8 Commits

Author SHA1 Message Date
Bram Veenboer
2a10cad3dd Fix compilation errors 2025-09-02 16:59:19 +02:00
Bram Veenboer
2c2a59d6d6 Apply formatting 2025-09-02 16:26:59 +02:00
Bram Veenboer
a1f2dd6c4d Apply suggestions from code review
Co-authored-by: Wiebe van Breukelen <breukelen@astron.nl>
2025-09-02 16:26:27 +02:00
Bram Veenboer
3dcca92b79 Remove remaining init and std::memcpy 2025-09-02 13:19:50 +02:00
Bram Veenboer
8df4bbf54e Add allocate_memory and free_memory 2025-09-02 12:03:31 +02:00
Bram Veenboer
716f323b26 Update GPU memory management
- Move device memory allocation for output out of init
- Copy directly from device memory to host pointers
2025-09-02 09:33:36 +02:00
Wiebe van Breukelen
9d3af8c202 Fixed broken pybind11 target check (#28)
* Fixed broken pybind11 target check

* Update python/CMakeLists.txt

Co-authored-by: Bram Veenboer <bram.veenboer@gmail.com>

---------

Co-authored-by: Bram Veenboer <bram.veenboer@gmail.com>
2025-08-27 17:07:30 +02:00
Bram Veenboer
0f7fd06be8 Extend CI (#27)
* Add build with Intel compiler
* Switch to upload/download-artifacts that retain permissions
2025-08-21 15:15:25 +02:00
6 changed files with 100 additions and 68 deletions

View File

@@ -13,6 +13,11 @@ jobs:
gres: gpu:A4000 gres: gpu:A4000
cmake_flags: "-DTRIGDX_USE_MKL=1 -DTRIGDX_USE_GPU=1 -DTRIGDX_USE_MKL=1 -DTRIGDX_USE_XSIMD=1 -DCMAKE_CUDA_ARCHITECTURES=86" cmake_flags: "-DTRIGDX_USE_MKL=1 -DTRIGDX_USE_GPU=1 -DTRIGDX_USE_MKL=1 -DTRIGDX_USE_XSIMD=1 -DCMAKE_CUDA_ARCHITECTURES=86"
environment_modules: "spack/20250403 intel-oneapi-mkl cuda python" environment_modules: "spack/20250403 intel-oneapi-mkl cuda python"
- name: Intel, NVIDIA A4000
partition: defq
gres: gpu:A4000
cmake_flags: "-DTRIGDX_USE_MKL=1 -DTRIGDX_USE_GPU=1 -DTRIGDX_USE_MKL=1 -DTRIGDX_USE_XSIMD=1 -DCMAKE_CUDA_ARCHITECTURES=86"
environment_modules: "spack/20250403 intel-oneapi-compilers intel-oneapi-mkl cuda python"
- name: NVIDIA GH200 - name: NVIDIA GH200
partition: ghq partition: ghq
gres: gpu:GH200 gres: gpu:GH200
@@ -38,7 +43,7 @@ jobs:
cmake -S . -B build ${CMAKE_FLAGS} cmake -S . -B build ${CMAKE_FLAGS}
make -C build -j make -C build -j
- name: Upload build - name: Upload build
uses: actions/upload-artifact@v4 uses: pyTooling/upload-artifact@v4
with: with:
name: build-${{ matrix.name }} name: build-${{ matrix.name }}
path: build path: build
@@ -53,7 +58,7 @@ jobs:
PARTITION_NAME: ${{ matrix.partition }} PARTITION_NAME: ${{ matrix.partition }}
steps: steps:
- *cleanup - *cleanup
- uses: actions/download-artifact@v4 - uses: pyTooling/download-artifact@v4
with: with:
name: build-${{ matrix.name }} name: build-${{ matrix.name }}
- uses: astron-rd/slurm-action@v1.2 - uses: astron-rd/slurm-action@v1.2
@@ -61,7 +66,7 @@ jobs:
partition: ${{ matrix.partition }} partition: ${{ matrix.partition }}
gres: ${{ matrix.gres }} gres: ${{ matrix.gres }}
commands: | commands: |
find tests -type f -executable -exec {} \; find build/tests -type f -executable -exec {} \;
benchmark: benchmark:
runs-on: [slurm] runs-on: [slurm]
@@ -73,7 +78,7 @@ jobs:
PARTITION_NAME: ${{ matrix.partition }} PARTITION_NAME: ${{ matrix.partition }}
steps: steps:
- *cleanup - *cleanup
- uses: actions/download-artifact@v4 - uses: pyTooling/download-artifact@v4
with: with:
name: build-${{ matrix.name }} name: build-${{ matrix.name }}
- uses: astron-rd/slurm-action@v1.2 - uses: astron-rd/slurm-action@v1.2
@@ -81,4 +86,4 @@ jobs:
partition: ${{ matrix.partition }} partition: ${{ matrix.partition }}
gres: ${{ matrix.gres }} gres: ${{ matrix.gres }}
commands: | commands: |
find benchmarks -type f -executable -exec {} \; find build/benchmarks -type f -executable -exec {} \;

View File

@@ -2,13 +2,14 @@
#include <chrono> #include <chrono>
#include <cmath> #include <cmath>
#include <stdexcept>
#include <string> #include <string>
#include <vector> #include <vector>
#include <benchmark/benchmark.h> #include <benchmark/benchmark.h>
void init_x(std::vector<float> &x) { void init_x(float *x, size_t n) {
for (size_t i = 0; i < x.size(); ++i) { for (size_t i = 0; i < n; ++i) {
x[i] = (i % 360) * 0.0174533f; // degrees to radians x[i] = (i % 360) * 0.0174533f; // degrees to radians
} }
} }
@@ -16,24 +17,31 @@ void init_x(std::vector<float> &x) {
template <typename Backend> template <typename Backend>
static void benchmark_sinf(benchmark::State &state) { static void benchmark_sinf(benchmark::State &state) {
const size_t N = static_cast<size_t>(state.range(0)); const size_t N = static_cast<size_t>(state.range(0));
std::vector<float> x(N), s(N);
init_x(x);
Backend backend; Backend backend;
auto start = std::chrono::high_resolution_clock::now(); auto start = std::chrono::high_resolution_clock::now();
backend.init(N); backend.init(N);
float *x =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
float *s =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
auto end = std::chrono::high_resolution_clock::now(); auto end = std::chrono::high_resolution_clock::now();
state.counters["init_ms"] = state.counters["init_ms"] =
std::chrono::duration_cast<std::chrono::microseconds>(end - start) std::chrono::duration_cast<std::chrono::microseconds>(end - start)
.count() / .count() /
1.e3; 1.e3;
init_x(x, N);
for (auto _ : state) { for (auto _ : state) {
backend.compute_sinf(N, x.data(), s.data()); backend.compute_sinf(N, x, s);
benchmark::DoNotOptimize(s); benchmark::DoNotOptimize(s);
} }
backend.free_memory(x);
backend.free_memory(s);
state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) * state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
static_cast<int64_t>(N)); static_cast<int64_t>(N));
} }
@@ -41,24 +49,35 @@ static void benchmark_sinf(benchmark::State &state) {
template <typename Backend> template <typename Backend>
static void benchmark_cosf(benchmark::State &state) { static void benchmark_cosf(benchmark::State &state) {
const size_t N = static_cast<size_t>(state.range(0)); const size_t N = static_cast<size_t>(state.range(0));
std::vector<float> x(N), c(N);
init_x(x);
Backend backend; Backend backend;
auto start = std::chrono::high_resolution_clock::now(); auto start = std::chrono::high_resolution_clock::now();
backend.init(N); backend.init(N);
float *x =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
float *c =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
if (!x || !c) {
throw std::runtime_error("Buffer allocation failed");
}
auto end = std::chrono::high_resolution_clock::now(); auto end = std::chrono::high_resolution_clock::now();
state.counters["init_ms"] = state.counters["init_ms"] =
std::chrono::duration_cast<std::chrono::microseconds>(end - start) std::chrono::duration_cast<std::chrono::microseconds>(end - start)
.count() / .count() /
1.e3; 1.e3;
init_x(x, N);
for (auto _ : state) { for (auto _ : state) {
backend.compute_cosf(N, x.data(), c.data()); backend.compute_cosf(N, x, c);
benchmark::DoNotOptimize(c); benchmark::DoNotOptimize(c);
} }
backend.free_memory(x);
backend.free_memory(c);
state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) * state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
static_cast<int64_t>(N)); static_cast<int64_t>(N));
} }
@@ -66,25 +85,38 @@ static void benchmark_cosf(benchmark::State &state) {
template <typename Backend> template <typename Backend>
static void benchmark_sincosf(benchmark::State &state) { static void benchmark_sincosf(benchmark::State &state) {
const size_t N = static_cast<size_t>(state.range(0)); const size_t N = static_cast<size_t>(state.range(0));
std::vector<float> x(N), s(N), c(N);
init_x(x);
Backend backend; Backend backend;
auto start = std::chrono::high_resolution_clock::now(); auto start = std::chrono::high_resolution_clock::now();
backend.init(N); backend.init(N);
float *x =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
float *s =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
float *c =
reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
if (!x || !s || !c) {
throw std::runtime_error("Buffer allocation failed");
}
auto end = std::chrono::high_resolution_clock::now(); auto end = std::chrono::high_resolution_clock::now();
state.counters["init_ms"] = state.counters["init_ms"] =
std::chrono::duration_cast<std::chrono::microseconds>(end - start) std::chrono::duration_cast<std::chrono::microseconds>(end - start)
.count() / .count() /
1.e3; 1.e3;
init_x(x, N);
for (auto _ : state) { for (auto _ : state) {
backend.compute_sincosf(N, x.data(), s.data(), c.data()); backend.compute_sincosf(N, x, s, c);
benchmark::DoNotOptimize(s); benchmark::DoNotOptimize(s);
benchmark::DoNotOptimize(c); benchmark::DoNotOptimize(c);
} }
backend.free_memory(x);
backend.free_memory(s);
backend.free_memory(c);
state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) * state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
static_cast<int64_t>(N)); static_cast<int64_t>(N));
} }

View File

@@ -11,7 +11,8 @@ public:
GPUBackend(); GPUBackend();
~GPUBackend() override; ~GPUBackend() override;
void init(size_t n = 0) override; void *allocate_memory(size_t bytes) const override;
void free_memory(void *ptr) const override;
void compute_sinf(size_t n, const float *x, float *s) const override; void compute_sinf(size_t n, const float *x, float *s) const override;
void compute_cosf(size_t n, const float *x, float *c) const override; void compute_cosf(size_t n, const float *x, float *c) const override;
void compute_sincosf(size_t n, const float *x, float *s, void compute_sincosf(size_t n, const float *x, float *s,

View File

@@ -1,6 +1,8 @@
#pragma once #pragma once
#include <cstddef> #include <cstddef>
#include <cstdint>
#include <cstdlib>
// Base interface for all math backends // Base interface for all math backends
class Backend { class Backend {
@@ -10,6 +12,12 @@ public:
// Optional initialization // Optional initialization
virtual void init(size_t n = 0) {} virtual void init(size_t n = 0) {}
virtual void *allocate_memory(size_t bytes) const {
return static_cast<void *>(new uint8_t[bytes]);
};
virtual void free_memory(void *ptr) const { std::free(ptr); };
// Compute sine for n elements // Compute sine for n elements
virtual void compute_sinf(size_t n, const float *x, float *s) const = 0; virtual void compute_sinf(size_t n, const float *x, float *s) const = 0;

View File

@@ -1,4 +1,6 @@
if(NOT TARGET pybind11) find_package(pybind11 CONFIG QUIET)
if(NOT pybind11_FOUND)
FetchContent_Declare( FetchContent_Declare(
pybind11 pybind11
GIT_REPOSITORY https://github.com/pybind/pybind11.git GIT_REPOSITORY https://github.com/pybind/pybind11.git

View File

@@ -10,79 +10,63 @@
struct GPUBackend::Impl { struct GPUBackend::Impl {
~Impl() { void *allocate_memory(size_t bytes) const {
if (h_x) { void *ptr;
cudaFreeHost(h_x); cudaMallocHost(&ptr, bytes);
} return ptr;
if (h_s) {
cudaFreeHost(h_s);
}
if (h_c) {
cudaFreeHost(h_c);
}
if (d_x) {
cudaFree(d_x);
}
if (d_s) {
cudaFree(d_s);
}
if (d_c) {
cudaFree(d_c);
}
} }
void init(size_t n) { void free_memory(void *ptr) const { cudaFreeHost(ptr); }
const size_t bytes = n * sizeof(float);
cudaMallocHost(&h_x, bytes);
cudaMallocHost(&h_s, bytes);
cudaMallocHost(&h_c, bytes);
cudaMalloc(&d_x, bytes);
cudaMalloc(&d_s, bytes);
cudaMalloc(&d_c, bytes);
}
void compute_sinf(size_t n, const float *x, float *s) const { void compute_sinf(size_t n, const float *x, float *s) const {
const size_t bytes = n * sizeof(float); const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes); float *d_x, *d_s;
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice); cudaMalloc(&d_x, bytes);
cudaMalloc(&d_s, bytes);
cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
launch_sinf_kernel(d_x, d_s, n); launch_sinf_kernel(d_x, d_s, n);
cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost); cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
std::memcpy(s, h_s, bytes); cudaFree(d_x);
cudaFree(d_s);
} }
void compute_cosf(size_t n, const float *x, float *c) const { void compute_cosf(size_t n, const float *x, float *c) const {
const size_t bytes = n * sizeof(float); const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes); float *d_x, *d_c;
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice); cudaMalloc(&d_x, bytes);
cudaMalloc(&d_c, bytes);
cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
launch_cosf_kernel(d_x, d_c, n); launch_cosf_kernel(d_x, d_c, n);
cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
std::memcpy(c, h_c, bytes); cudaFree(d_x);
cudaFree(d_c);
} }
void compute_sincosf(size_t n, const float *x, float *s, float *c) const { void compute_sincosf(size_t n, const float *x, float *s, float *c) const {
const size_t bytes = n * sizeof(float); const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes); float *d_x, *d_s, *d_c;
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice); cudaMalloc(&d_x, bytes);
cudaMalloc(&d_s, bytes);
cudaMalloc(&d_c, bytes);
cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
launch_sincosf_kernel(d_x, d_s, d_c, n); launch_sincosf_kernel(d_x, d_s, d_c, n);
cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost); cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
std::memcpy(s, h_s, bytes); cudaFree(d_x);
std::memcpy(c, h_c, bytes); cudaFree(d_s);
cudaFree(d_c);
} }
float *h_x = nullptr;
float *h_s = nullptr;
float *h_c = nullptr;
float *d_x = nullptr;
float *d_s = nullptr;
float *d_c = nullptr;
}; };
GPUBackend::GPUBackend() : impl(std::make_unique<Impl>()) {} GPUBackend::GPUBackend() : impl(std::make_unique<Impl>()) {}
GPUBackend::~GPUBackend() = default; GPUBackend::~GPUBackend() = default;
void GPUBackend::init(size_t n) { impl->init(n); } void *GPUBackend::allocate_memory(size_t bytes) const {
return impl->allocate_memory(bytes);
}
void GPUBackend::free_memory(void *ptr) const { impl->free_memory(ptr); }
void GPUBackend::compute_sinf(size_t n, const float *x, float *s) const { void GPUBackend::compute_sinf(size_t n, const float *x, float *s) const {
impl->compute_sinf(n, x, s); impl->compute_sinf(n, x, s);