From 17b708c38a51394c604fe4ea64124c014fb0c7a0 Mon Sep 17 00:00:00 2001 From: Arijit Dasgupta Date: Thu, 20 Nov 2025 19:58:23 +0000 Subject: [PATCH 1/4] convert test_utils.cu from a GT test to a bounds-only test --- genmetaballs/src/cuda/bindings.cu | 1 - genmetaballs/src/cuda/core/confidence.cuh | 1 - tests/cpp_tests/test_utils.cu | 27 +++++------------------ 3 files changed, 5 insertions(+), 24 deletions(-) diff --git a/genmetaballs/src/cuda/bindings.cu b/genmetaballs/src/cuda/bindings.cu index 6a530e0..b60f230 100644 --- a/genmetaballs/src/cuda/bindings.cu +++ b/genmetaballs/src/cuda/bindings.cu @@ -2,7 +2,6 @@ #include #include #include -#include #include "core/confidence.cuh" #include "core/geometry.cuh" diff --git a/genmetaballs/src/cuda/core/confidence.cuh b/genmetaballs/src/cuda/core/confidence.cuh index 5320b0e..d605c96 100644 --- a/genmetaballs/src/cuda/core/confidence.cuh +++ b/genmetaballs/src/cuda/core/confidence.cuh @@ -2,7 +2,6 @@ #include #include -#include #include "utils.cuh" diff --git a/tests/cpp_tests/test_utils.cu b/tests/cpp_tests/test_utils.cu index 575502a..8c26120 100644 --- a/tests/cpp_tests/test_utils.cu +++ b/tests/cpp_tests/test_utils.cu @@ -13,7 +13,7 @@ namespace test_utils_gpu { -// CUDA kernel for computing sigmoid element-wise (relies on __device__ sigmoid in utils.cuh) +// CUDA kernel for computing sigmoid element-wise __global__ void sigmoid_kernel(const float* x, float* result, uint32_t n) { uint32_t i = threadIdx.x + blockIdx.x * blockDim.x; if (i < n) { @@ -21,7 +21,7 @@ __global__ void sigmoid_kernel(const float* x, float* result, uint32_t n) { } } -// GPU function to compute sigmoid for a vector (float only) +// GPU function to compute sigmoid for a vector template std::vector gpu_sigmoid(const std::vector& x_vec) { uint32_t n = x_vec.size(); @@ -45,11 +45,6 @@ std::vector gpu_sigmoid(const std::vector& x_vec) { return result; } -// Host sigmoid for reference -inline float host_sigmoid(float x) { - return 1.0f / (1.0f + std::exp(-x)); -} - } // namespace test_utils_gpu // Parameters matching the removed Python test @@ -65,7 +60,7 @@ static std::vector sigmoid_test_sizes() { return sizes; } -TEST(GpuSigmoidTest, SigmoidVectorCorrectness) { +TEST(GpuSigmoidTest, SigmoidGPUWithinBounds) { // Generate seeds std::mt19937 master_gen(SEED_MASTER); std::uniform_int_distribution seed_dist(0, std::numeric_limits::max()); @@ -85,22 +80,10 @@ TEST(GpuSigmoidTest, SigmoidVectorCorrectness) { for (int i = 0; i < N; ++i) x_vec[i] = dist(rng); - // Compute expected (host) - std::vector expected(N); - for (int i = 0; i < N; ++i) - expected[i] = test_utils_gpu::host_sigmoid(x_vec[i]); - - // Compute actual (GPU) + // Run on GPU constexpr uint32_t block_dim = 256; uint32_t grid_dim = (N + block_dim - 1) / block_dim; - std::vector actual = test_utils_gpu::gpu_sigmoid<1024, block_dim>(x_vec); - - // Compare - ASSERT_EQ(actual.size(), expected.size()); - for (int i = 0; i < N; ++i) { - ASSERT_NEAR(actual[i], expected[i], 1e-5) - << "at idx=" << i << " for N=" << N << " seed=" << seed; - } + std::vector actual = test_utils_gpu::gpu_sigmoid(x_vec); // Check [0, 1] bounds ASSERT_TRUE(std::all_of(actual.begin(), actual.end(), From 81acfc4db88494641f15dc1c7310fb2ad78d4e8d Mon Sep 17 00:00:00 2001 From: Arijit Dasgupta Date: Thu, 20 Nov 2025 20:05:11 +0000 Subject: [PATCH 2/4] fixed compilation issue with test_utils --- genmetaballs/src/cuda/bindings.cu | 3 --- tests/cpp_tests/test_utils.cu | 3 +-- 2 files changed, 1 insertion(+), 5 deletions(-) diff --git a/genmetaballs/src/cuda/bindings.cu b/genmetaballs/src/cuda/bindings.cu index b60f230..5a48a8a 100644 --- a/genmetaballs/src/cuda/bindings.cu +++ b/genmetaballs/src/cuda/bindings.cu @@ -7,9 +7,6 @@ #include "core/geometry.cuh" #include "core/utils.cuh" -constexpr uint32_t GRID_DIM = 4096; -constexpr uint32_t BLOCK_DIM = 1024; - namespace nb = nanobind; NB_MODULE(_genmetaballs_bindings, m) { diff --git a/tests/cpp_tests/test_utils.cu b/tests/cpp_tests/test_utils.cu index 8c26120..b4cc098 100644 --- a/tests/cpp_tests/test_utils.cu +++ b/tests/cpp_tests/test_utils.cu @@ -82,8 +82,7 @@ TEST(GpuSigmoidTest, SigmoidGPUWithinBounds) { // Run on GPU constexpr uint32_t block_dim = 256; - uint32_t grid_dim = (N + block_dim - 1) / block_dim; - std::vector actual = test_utils_gpu::gpu_sigmoid(x_vec); + std::vector actual = test_utils_gpu::gpu_sigmoid<1024, block_dim>(x_vec); // Check [0, 1] bounds ASSERT_TRUE(std::all_of(actual.begin(), actual.end(), From 54d84a0a4429142c667cf84f2b6acf0a0a718e43 Mon Sep 17 00:00:00 2001 From: Arijit Dasgupta Date: Thu, 20 Nov 2025 20:08:58 +0000 Subject: [PATCH 3/4] removed isnan check for sigmoid --- genmetaballs/src/cuda/core/utils.cuh | 3 --- 1 file changed, 3 deletions(-) diff --git a/genmetaballs/src/cuda/core/utils.cuh b/genmetaballs/src/cuda/core/utils.cuh index 1dcdf36..d9b8f28 100644 --- a/genmetaballs/src/cuda/core/utils.cuh +++ b/genmetaballs/src/cuda/core/utils.cuh @@ -15,9 +15,6 @@ void cuda_check(cudaError_t code, const char* file, int line); CUDA_CALLABLE __forceinline__ float sigmoid(float x) { - if (isnan(x)) { - return x; - } return 1.0f / (1.0f + expf(-x)); } From 49f19d33e7b81b4d24d0fd72833542ecf671650e Mon Sep 17 00:00:00 2001 From: Arijit Dasgupta Date: Thu, 20 Nov 2025 20:14:21 +0000 Subject: [PATCH 4/4] removed GT testing from confidence cpp testing, GT testing remains in python --- tests/cpp_tests/test_confidence.cu | 28 ---------------------------- 1 file changed, 28 deletions(-) diff --git a/tests/cpp_tests/test_confidence.cu b/tests/cpp_tests/test_confidence.cu index cda499a..2b28b05 100644 --- a/tests/cpp_tests/test_confidence.cu +++ b/tests/cpp_tests/test_confidence.cu @@ -10,17 +10,6 @@ #include "core/confidence.cuh" -// Helper: Python ground truth, as in test_confidence.py -inline float ground_truth_expit(float x) { - return 1.0F / (1.0F + std::exp(-x)); -} -float ground_truth_two_parameter_confidence(float beta4, float beta5, float sumexpd) { - return ground_truth_expit((beta4 * sumexpd) + beta5); -} -float ground_truth_zero_parameter_confidence(float sumexpd) { - return 1.0F - std::exp(-sumexpd); -} - template __global__ void confidence_kernel(const float* sumexpd, float* confidences, uint32_t n, Confidence confidence) { @@ -69,7 +58,6 @@ static std::vector confidence_test_sizes() { return sizes; } -// Define simple struct to match python CONFIDENCE_TEST_CASES struct ConfidenceCase { std::string name; float beta4 = 0.0F; @@ -112,16 +100,6 @@ TEST(GpuConfidenceTest, ConfidenceMultipleValuesGPU_AllTypes) { for (int i = 0; i < N; ++i) sumexpd_vec[i] = dist(rng); - std::vector expected(N); - if (conf_case.is_two_param) { - for (int i = 0; i < N; ++i) - expected[i] = ground_truth_two_parameter_confidence( - conf_case.beta4, conf_case.beta5, sumexpd_vec[i]); - } else { - for (int i = 0; i < N; ++i) - expected[i] = ground_truth_zero_parameter_confidence(sumexpd_vec[i]); - } - std::vector actual; if (conf_case.is_two_param) { TwoParameterConfidence conf(conf_case.beta4, conf_case.beta5); @@ -131,12 +109,6 @@ TEST(GpuConfidenceTest, ConfidenceMultipleValuesGPU_AllTypes) { actual = gpu_get_confidence(sumexpd_vec, conf); } - ASSERT_EQ(actual.size(), expected.size()); - for (int i = 0; i < N; ++i) { - ASSERT_NEAR(actual[i], expected[i], 1e-6F) - << "at idx=" << i << " N=" << N << " conf_type=" << conf_case.name - << " exp=" << expected[i] << " act=" << actual[i]; - } // Ensure all actual values are in [0, 1] ASSERT_TRUE(std::all_of(actual.begin(), actual.end(), [](float v) { return v >= 0.0F && v <= 1.0F; }))