Remove unused get_ray_directions and update test to use PixelCoordRange

Author: horizon-blue

genmetaballs/src/cuda/core/camera.cu

@@ -30,10 +30,10 @@ CUDA_CALLABLE bool PixelCoordRange::Sentinel::operator==(const Iterator& it) con
     return it.py >= py_end;
 }
 
-CUDA_CALLABLE constexpr PixelCoordRange::Iterator PixelCoordRange::begin() const {
+CUDA_CALLABLE PixelCoordRange::Iterator PixelCoordRange::begin() const {
     return Iterator{px_start, px_end, py_start, px_start, py_start};
 }
 
-CUDA_CALLABLE constexpr PixelCoordRange::Sentinel PixelCoordRange::end() const {
+CUDA_CALLABLE PixelCoordRange::Sentinel PixelCoordRange::end() const {
     return Sentinel{py_end};
 }

genmetaballs/src/cuda/core/camera.cuh

@@ -1,7 +1,6 @@
 #pragma once
 
 #include <cstdint>
-#include <cuda/std/ranges>
 #include <cuda/std/utility>
 #include <cuda_runtime.h>
 
@@ -19,25 +18,9 @@ struct Intrinsics {
     // Returns the direction of the ray going through pixel (px, py) in camera frame.
     // For efficiency, this function does not check if the pixel is within bounds.
     CUDA_CALLABLE Vec3D get_ray_direction(uint32_t px, uint32_t py) const;
-
-    // Returns a 2D array of ray directions in camera frame in the specified pixel range
-    // and store them in the provided buffer. By default, the full image is used
-    template <MemoryLocation location>
-    CUDA_CALLABLE Array2D<Vec3D, location>& get_ray_directions(Array2D<Vec3D, location>& buffer,
-                                                               uint32_t px_start = 0,
-                                                               uint32_t px_end = UINT32_MAX,
-                                                               uint32_t py_start = 0,
-                                                               uint32_t py_end = UINT32_MAX) const {
-        for (auto i = max(0, px_start); i < min(height, px_end); ++i) {
-            for (auto j = max(0, py_start); j < min(width, py_end); ++j) {
-                buffer[i][j] = get_ray_direction(j, i);
-            }
-        }
-        return buffer;
-    }
 };
 
-struct PixelCoordRange : public cuda::std::ranges::view_interface<PixelCoordRange> {
+struct PixelCoordRange {
     uint32_t px_start;
     uint32_t px_end;
     uint32_t py_start;
@@ -70,6 +53,6 @@ struct PixelCoordRange : public cuda::std::ranges::view_interface<PixelCoordRang
     };
 
     // range methods
-    CUDA_CALLABLE constexpr Iterator begin() const;
-    CUDA_CALLABLE constexpr Sentinel end() const;
+    CUDA_CALLABLE Iterator begin() const;
+    CUDA_CALLABLE Sentinel end() const;
 };

tests/cpp_tests/test_camera.cu

@@ -11,15 +11,19 @@
 
 namespace test_camera_gpu {
 
-// CUDA kernel to call get_ray_directions on device with multiple threads
-// Each thread processes one row of the image
+// CUDA kernel to call get_ray_direction on device with multiple threads
+// Each thread processes one row of the image via PixelCoordRange
 __global__ void get_ray_directions_kernel(Intrinsics intrinsics,
                                           Array2D<Vec3D, MemoryLocation::DEVICE> ray_buffer) {
     uint32_t row_start = threadIdx.x * 2;
     uint32_t row_end = max(row_start + 2, intrinsics.height);
     uint32_t col_start = threadIdx.y * 2;
     uint32_t col_end = max(col_start + 2, intrinsics.width);
-    intrinsics.get_ray_directions(ray_buffer, row_start, row_end, col_start, col_end);
+    auto pixel_coords = PixelCoordRange{row_start, row_end, col_start, col_end};
+
+    for (auto [px, py] : pixel_coords) {
+        ray_buffer[px][py] = intrinsics.get_ray_direction(px, py);
+    }
 }
 
 } // namespace test_camera_gpu