MET-46: Cleanup Confidence Implementation

genmetaballs/src/cuda/bindings.cu

@@ -2,15 +2,11 @@
 #include <nanobind/nanobind.h>
 #include <nanobind/operators.h>
 #include <nanobind/stl/vector.h>
-#include <stdexcept>
 
 #include "core/confidence.cuh"
 #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) {

genmetaballs/src/cuda/core/confidence.cuh

@@ -2,7 +2,6 @@
 
 #include <cmath>
 #include <cuda_runtime.h>
-#include <vector>
 
 #include "utils.cuh"
 

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));
 }
 

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 <typename Confidence>
 __global__ void confidence_kernel(const float* sumexpd, float* confidences, uint32_t n,
                                   Confidence confidence) {
@@ -69,7 +58,6 @@ static std::vector<int> 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<float> 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<float> 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; }))

tests/cpp_tests/test_utils.cu

@@ -13,15 +13,15 @@
 
 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) {
         result[i] = sigmoid(x[i]);
     }
 }
 
-// GPU function to compute sigmoid for a vector (float only)
+// GPU function to compute sigmoid for a vector
 template <uint32_t grid_dim, uint32_t block_dim>
 std::vector<float> gpu_sigmoid(const std::vector<float>& x_vec) {
     uint32_t n = x_vec.size();
@@ -45,11 +45,6 @@ std::vector<float> gpu_sigmoid(const std::vector<float>& 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<int> sigmoid_test_sizes() {
     return sizes;
 }
 
-TEST(GpuSigmoidTest, SigmoidVectorCorrectness) {
+TEST(GpuSigmoidTest, SigmoidGPUWithinBounds) {
     // Generate seeds
     std::mt19937 master_gen(SEED_MASTER);
     std::uniform_int_distribution<uint32_t> seed_dist(0, std::numeric_limits<uint32_t>::max());
@@ -85,23 +80,10 @@ TEST(GpuSigmoidTest, SigmoidVectorCorrectness) {
             for (int i = 0; i < N; ++i)
                 x_vec[i] = dist(rng);
 
-            // Compute expected (host)
-            std::vector<float> 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<float> 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;
-            }
-
             // Check [0, 1] bounds
             ASSERT_TRUE(std::all_of(actual.begin(), actual.end(),
                                     [](float v) { return v >= 0.0f && v <= 1.0f; }))