6 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
12 changed files with 88 additions and 138 deletions

View File

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

View File

@@ -11,12 +11,12 @@ public:
GPUBackend();
~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_cosf(size_t n, const float *x, float *c) const override;
void compute_sincosf(size_t n, const float *x, float *s,
float *c) const override;
void compute_expf(size_t n, const float *x, float *e) const override;
private:
struct Impl;

View File

@@ -1,6 +1,8 @@
#pragma once
#include <cstddef>
#include <cstdint>
#include <cstdlib>
// Base interface for all math backends
class Backend {
@@ -10,6 +12,12 @@ public:
// Optional initialization
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
virtual void compute_sinf(size_t n, const float *x, float *s) const = 0;
@@ -19,7 +27,4 @@ public:
// Compute sine and cosine for n elements
virtual void compute_sincosf(size_t n, const float *x, float *s,
float *c) const = 0;
// Compute exponent for n elements
virtual void compute_expf(size_t n, const float *x, float *e) const = 0;
};

View File

@@ -10,6 +10,4 @@ public:
void compute_sincosf(size_t n, const float *x, float *s,
float *c) const override;
void compute_expf(size_t n, const float *x, float *e) const override;
};

View File

@@ -10,6 +10,4 @@ public:
void compute_sincosf(size_t n, const float *x, float *s,
float *c) const override;
void compute_expf(size_t n, const float *x, float *e) const override;
};

View File

@@ -10,98 +10,63 @@
struct GPUBackend::Impl {
~Impl() {
if (h_x) {
cudaFreeHost(h_x);
}
if (h_s) {
cudaFreeHost(h_s);
}
if (h_c) {
cudaFreeHost(h_c);
}
if (h_e) {
cudaFreeHost(h_e);
}
if (d_x) {
cudaFree(d_x);
}
if (d_s) {
cudaFree(d_s);
}
if (d_c) {
cudaFree(d_c);
}
if (d_e) {
cudaFree(d_e);
}
void *allocate_memory(size_t bytes) const {
void *ptr;
cudaMallocHost(&ptr, bytes);
return ptr;
}
void init(size_t n) {
const size_t bytes = n * sizeof(float);
cudaMallocHost(&h_x, bytes);
cudaMallocHost(&h_s, bytes);
cudaMallocHost(&h_c, bytes);
cudaMallocHost(&h_e, bytes);
cudaMalloc(&d_x, bytes);
cudaMalloc(&d_s, bytes);
cudaMalloc(&d_c, bytes);
cudaMalloc(&d_e, bytes);
}
void free_memory(void *ptr) const { cudaFreeHost(ptr); }
void compute_sinf(size_t n, const float *x, float *s) const {
const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes);
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
float *d_x, *d_s;
cudaMalloc(&d_x, bytes);
cudaMalloc(&d_s, bytes);
cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
launch_sinf_kernel(d_x, d_s, n);
cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost);
std::memcpy(s, h_s, bytes);
cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
cudaFree(d_x);
cudaFree(d_s);
}
void compute_cosf(size_t n, const float *x, float *c) const {
const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes);
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
float *d_x, *d_c;
cudaMalloc(&d_x, bytes);
cudaMalloc(&d_c, bytes);
cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
launch_cosf_kernel(d_x, d_c, n);
cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
std::memcpy(c, h_c, bytes);
cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
cudaFree(d_x);
cudaFree(d_c);
}
void compute_sincosf(size_t n, const float *x, float *s, float *c) const {
const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes);
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
float *d_x, *d_s, *d_c;
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);
cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost);
cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
std::memcpy(s, h_s, bytes);
std::memcpy(c, h_c, bytes);
cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
cudaFree(d_x);
cudaFree(d_s);
cudaFree(d_c);
}
void compute_expf(size_t n, const float *x, float *e) const {
const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes);
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
launch_expf_kernel(d_x, d_e, n);
cudaMemcpy(h_e, d_e, bytes, cudaMemcpyDeviceToHost);
std::memcpy(e, h_e, bytes);
}
float *h_x = nullptr;
float *h_s = nullptr;
float *h_c = nullptr;
float *h_e = nullptr;
float *d_x = nullptr;
float *d_s = nullptr;
float *d_c = nullptr;
float *d_e = nullptr;
};
GPUBackend::GPUBackend() : impl(std::make_unique<Impl>()) {}
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 {
impl->compute_sinf(n, x, s);
@@ -114,8 +79,4 @@ void GPUBackend::compute_cosf(size_t n, const float *x, float *c) const {
void GPUBackend::compute_sincosf(size_t n, const float *x, float *s,
float *c) const {
impl->compute_sincosf(n, x, s, c);
}
void GPUBackend::compute_expf(size_t n, const float *x, float *e) const {
impl->compute_expf(n, x, e);
}
}

View File

@@ -31,15 +31,6 @@ __global__ void kernel_sincosf(const float *__restrict__ x,
}
}
__global__ void kernel_expf(const float *__restrict__ x, float *__restrict__ e,
size_t n) {
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
// e[idx] = __expf(x[idx]);
e[idx] = expf(x[idx]);
}
}
namespace {
inline dim3 make_grid(size_t n, size_t threadsPerBlock = 256) {
return dim3((n + threadsPerBlock - 1) / threadsPerBlock);
@@ -63,9 +54,3 @@ void launch_sincosf_kernel(const float *d_x, float *d_s, float *d_c, size_t n) {
dim3 grid = make_grid(n, blocks.x);
kernel_sincosf<<<grid, blocks>>>(d_x, d_s, d_c, n);
}
void launch_expf_kernel(const float *d_x, float *d_e, size_t n) {
dim3 blocks(256);
dim3 grid = make_grid(n, blocks.x);
kernel_expf<<<grid, blocks>>>(d_x, d_e, n);
}

View File

@@ -6,4 +6,3 @@ void launch_sinf_kernel(const float *d_x, float *d_s, size_t n);
void launch_cosf_kernel(const float *d_x, float *d_c, size_t n);
void launch_sincosf_kernel(const float *d_x, float *d_s, float *d_c,
std::size_t n);
void launch_expf_kernel(const float *d_x, float *d_e, size_t n);

View File

@@ -14,7 +14,3 @@ void MKLBackend::compute_sincosf(size_t n, const float *x, float *s,
float *c) const {
vmsSinCos(static_cast<MKL_INT>(n), x, s, c, VML_HA);
}
void MKLBackend::compute_expf(size_t n, const float *x, float *e) const {
vmsExp(static_cast<MKL_INT>(n), x, e, VML_HA);
}

View File

@@ -21,9 +21,3 @@ void ReferenceBackend::compute_sincosf(size_t n, const float *x, float *s,
c[i] = cosf(x[i]);
}
}
void ReferenceBackend::compute_expf(size_t n, const float *x, float *e) const {
for (size_t i = 0; i < n; ++i) {
e[i] = expf(x[i]);
}
}

View File

@@ -8,5 +8,3 @@ TEST_CASE("sinf") { test_sinf<MKLBackend>(1e-6f); }
TEST_CASE("cosf") { test_cosf<MKLBackend>(1e-6f); }
TEST_CASE("sincosf") { test_sincosf<MKLBackend>(1e-6f); }
TEST_CASE("expf") { test_expf<MKLBackend>(1e-6f); }

View File

@@ -63,19 +63,3 @@ template <typename Backend> inline void test_sincosf(float tol) {
REQUIRE_THAT(c[i], Catch::Matchers::WithinAbs(c_ref[i], tol));
}
}
template <typename Backend> inline void test_expf(float tol) {
std::vector<float> x(N), e_ref(N), e(N);
init_x(x);
ReferenceBackend ref;
Backend backend;
backend.init(N);
ref.compute_expf(N, x.data(), e_ref.data());
backend.compute_expf(N, x.data(), e.data());
for (size_t i = 0; i < N; ++i) {
REQUIRE_THAT(e[i], Catch::Matchers::WithinAbs(e_ref[i], tol));
}
}