mirror/trigdx
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #11 from astron-rd/fix-cuda-test

Browse Source 
Bugfixes in data copies in CUDA backend
This commit is contained in:
Bram Veenboer
2025-08-12 17:16:10 +02:00
committed by GitHub
parent eb3806442d 33f98abc48
commit cdc94ab2cb
4 changed files with 26 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files

22
benchmarks/benchmark_utils.hpp
View File

@@ -7,17 +7,17 @@
#include <benchmark/benchmark.h>
// Default values if not overridden by range multipliers
constexpr size_t DEFAULT_N = 10'000'000;
void init_x(std::vector<float> &x) {
for (size_t i = 0; i < x.size(); ++i) {
x[i] = (i % 360) * 0.0174533f; // degrees to radians
}
}
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);
for (size_t i = 0; i < N; ++i) {
x[i] = (i % 360) * 0.0174533f; // degrees to radians
}
init_x(x);
Backend backend;
@@ -42,10 +42,7 @@ 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);
for (size_t i = 0; i < N; ++i) {
x[i] = (i % 360) * 0.0174533f;
}
init_x(x);
Backend backend;
@@ -70,10 +67,7 @@ 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);
for (size_t i = 0; i < N; ++i) {
x[i] = (i % 360) * 0.0174533f;
}
init_x(x);
Backend backend;

8
src/gpu.cpp
View File

@@ -45,25 +45,29 @@ struct GPUBackend::Impl {
const size_t bytes = n * sizeof(float);
std::memcpy(h_x, x, bytes);
cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
launch_sincosf_kernel(d_x, d_s, d_c, n);
launch_sinf_kernel(d_x, d_s, n);
cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost);
std::memcpy(s, h_s, bytes);
}
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);
launch_sincosf_kernel(d_x, d_s, d_c, n);
launch_cosf_kernel(d_x, d_c, n);
cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
std::memcpy(c, h_c, bytes);
}
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);
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);
}
float *h_x = nullptr;

6
tests/test_gpu.cpp
View File

@@ -3,8 +3,8 @@
#include "test_utils.hpp"
TEST_CASE("sinf") { test_sinf<GPUBackend>(1e-1f); }
TEST_CASE("sinf") { test_sinf<GPUBackend>(1e-6f); }
TEST_CASE("cosf") { test_cosf<GPUBackend>(1e-1f); }
TEST_CASE("cosf") { test_cosf<GPUBackend>(1e-6f); }
TEST_CASE("sincosf") { test_sincosf<GPUBackend>(1e-1f); }
TEST_CASE("sincosf") { test_sincosf<GPUBackend>(1e-6f); }

21
tests/test_utils.hpp
View File

@@ -9,12 +9,15 @@
const size_t N = 1e7;
void init_x(std::vector<float> &x) {
for (size_t i = 0; i < x.size(); ++i) {
x[i] = (i % 360) * 0.0174533f; // degrees to radians
}
}
template <typename Backend> inline void test_sinf(float tol) {
std::vector<float> x(N), s_ref(N), s(N);
for (size_t i = 0; i < N; ++i) {
x[i] = float(i) * 0.01f;
}
init_x(x);
ReferenceBackend ref;
Backend backend;
@@ -30,10 +33,7 @@ template <typename Backend> inline void test_sinf(float tol) {
template <typename Backend> inline void test_cosf(float tol) {
std::vector<float> x(N), c_ref(N), c(N);
for (size_t i = 0; i < N; ++i) {
x[i] = float(i) * 0.01f;
}
init_x(x);
ReferenceBackend ref;
Backend backend;
@@ -49,10 +49,7 @@ template <typename Backend> inline void test_cosf(float tol) {
template <typename Backend> inline void test_sincosf(float tol) {
std::vector<float> x(N), s_ref(N), c_ref(N), s(N), c(N);
for (size_t i = 0; i < N; ++i) {
x[i] = float(i) * 0.01f;
}
init_x(x);
ReferenceBackend ref;
Backend backend;