From 4a186e17e4aa096bb59c263eb917a5479410b906 Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 00:15:21 -0500 Subject: [PATCH 1/7] Compute ray direction in camera frame & world frame --- CMakeLists.txt | 2 ++ genmetaballs/src/cuda/core/camera.cu | 17 +++++++++++++++++ genmetaballs/src/cuda/core/camera.cuh | 17 +++++++++++++++++ genmetaballs/src/cuda/core/geometry.cuh | 4 ++-- 4 files changed, 38 insertions(+), 2 deletions(-) create mode 100644 genmetaballs/src/cuda/core/camera.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 664c437..989b9ce 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -14,6 +14,8 @@ set(CMAKE_EXPORT_COMPILE_COMMANDS ON) ################ add_library(genmetaballs_core + genmetaballs/src/cuda/core/camera.cu + genmetaballs/src/cuda/core/camera.cuh genmetaballs/src/cuda/core/utils.cu genmetaballs/src/cuda/core/utils.cuh genmetaballs/src/cuda/core/geometry.cuh diff --git a/genmetaballs/src/cuda/core/camera.cu b/genmetaballs/src/cuda/core/camera.cu new file mode 100644 index 0000000..385a261 --- /dev/null +++ b/genmetaballs/src/cuda/core/camera.cu @@ -0,0 +1,17 @@ +#include +#include + +#include "camera.cuh" +#include "geometry.cuh" +#include "utils.cuh" + +CUDA_CALLABLE Vec3D Intrinsics::get_ray_direction(uint32_t px, uint32_t py) const { + auto x = (static_cast(px) - cx) / fx; + auto y = (static_cast(py) - cy) / fy; + return Vec3D{x, y, -1.0f}; +} + +CUDA_CALLABLE Vec3D Camera::get_ray_direction(uint32_t px, uint32_t py) const { + auto ray_cam = intrinsics.get_ray_direction(px, py); + return extrinsics.get_rot().apply(ray_cam); +} diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index 4b2d2b2..fa887d6 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -1,6 +1,10 @@ #pragma once #include +#include + +#include "geometry.cuh" +#include "utils.cuh" struct Intrinsics { uint32_t height; @@ -11,4 +15,17 @@ struct Intrinsics { float cy; float near; float far; + + // 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; +}; + +struct Camera { + Intrinsics intrinsics; + Pose extrinsics; + + // returns the direction of the ray going through pixel (px, py) in world 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; }; diff --git a/genmetaballs/src/cuda/core/geometry.cuh b/genmetaballs/src/cuda/core/geometry.cuh index f64dbd2..dd67adb 100644 --- a/genmetaballs/src/cuda/core/geometry.cuh +++ b/genmetaballs/src/cuda/core/geometry.cuh @@ -72,11 +72,11 @@ public: return {rot, tran}; } - CUDA_CALLABLE Rotation get_rot() const { + CUDA_CALLABLE const Rotation& get_rot() const { return rot_; } - CUDA_CALLABLE Vec3D get_tran() const { + CUDA_CALLABLE const Vec3D& get_tran() const { return tran_; } From 8167ce2cf28e18238a910fe77b9970759fc93462 Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 00:15:22 -0500 Subject: [PATCH 2/7] Refactor get_camera_rays to assist testing --- examples/run_fmb.py | 6 ++---- genmetaballs/src/genmetaballs/fmb/utils.py | 9 +++++++++ 2 files changed, 11 insertions(+), 4 deletions(-) diff --git a/examples/run_fmb.py b/examples/run_fmb.py index d8d27b3..a3ba4a4 100644 --- a/examples/run_fmb.py +++ b/examples/run_fmb.py @@ -33,7 +33,7 @@ from tqdm import tqdm # Import local utilities -from genmetaballs.fmb.utils import DegradeLR, image_grid +from genmetaballs.fmb.utils import DegradeLR, get_camera_rays, image_grid CURRENT_DIR = Path(__file__).parent PROJECT_ROOT = CURRENT_DIR.parent @@ -475,14 +475,12 @@ def main(): # Setup camera rays height, width = image_size - K = np.array([[focal_length, 0, cx], [0, focal_length, cy], [0, 0, 1]]) pixel_list = ( (np.array(np.meshgrid(np.arange(width), height - np.arange(height) - 1, [0]))[:, :, :, 0]) .reshape((3, -1)) .T ) - camera_rays = (pixel_list - K[:, 2]) / np.diag(K) - camera_rays[:, -1] = -1 + camera_rays = get_camera_rays(focal_length, focal_length, cx, cy, pixel_list) cameras_list = [] for tran, quat in zip(trans, rand_quats, strict=False): R = transforms3d.quaternions.quat2mat(quat) diff --git a/genmetaballs/src/genmetaballs/fmb/utils.py b/genmetaballs/src/genmetaballs/fmb/utils.py index 6418c31..4ec8ea0 100644 --- a/genmetaballs/src/genmetaballs/fmb/utils.py +++ b/genmetaballs/src/genmetaballs/fmb/utils.py @@ -141,3 +141,12 @@ def compute_normals(camera_rays, depth_py_px, eps=1e-20): norms = nan_ddiff / (eps + jnp.linalg.norm(nan_ddiff, axis=1, keepdims=True)) return norms + + +def get_camera_rays( + fx: float, fy: float, cx: float, cy: float, pixel_list: np.ndarray[int] +) -> np.ndarray[np.float32]: + K = np.array([[fx, 0, cx], [0, fy, cy], [0, 0, 1]], dtype=np.float32) + camera_rays = (pixel_list - K[:, 2]) / np.diag(K) + camera_rays[:, -1] = -1 + return camera_rays From b59aaf2270fbaed9927c96cb18aa6511d1cb450f Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 00:15:22 -0500 Subject: [PATCH 3/7] Binding & testing for intrinsics.get_ray_direction --- genmetaballs/src/cuda/bindings.cu | 28 +++++++++++- genmetaballs/src/cuda/core/camera.cuh | 2 - .../src/genmetaballs/core/__init__.py | 3 ++ tests/python_tests/test_camera.py | 43 +++++++++++++++++++ 4 files changed, 73 insertions(+), 3 deletions(-) create mode 100644 tests/python_tests/test_camera.py diff --git a/genmetaballs/src/cuda/bindings.cu b/genmetaballs/src/cuda/bindings.cu index dff316f..d11e897 100644 --- a/genmetaballs/src/cuda/bindings.cu +++ b/genmetaballs/src/cuda/bindings.cu @@ -5,6 +5,7 @@ #include #include "core/blender.cuh" +#include "core/camera.cuh" #include "core/confidence.cuh" #include "core/geometry.cuh" #include "core/utils.cuh" @@ -65,7 +66,32 @@ NB_MODULE(_genmetaballs_bindings, m) { .def_rw("direction", &Ray::direction); /* - * Confidence submodule bindings + * Camera module bindings + */ + nb::module_ camera = m.def_submodule("camera", "Camera intrinsics and extrinsics"); + nb::class_(camera, "Intrinsics") + .def(nb::init(), nb::arg("height"), + nb::arg("width"), nb::arg("fx"), nb::arg("fy"), nb::arg("cx"), nb::arg("cy")) + .def_ro("height", &Intrinsics::height) + .def_ro("width", &Intrinsics::width) + .def_ro("fx", &Intrinsics::fx) + .def_ro("fy", &Intrinsics::fy) + .def_ro("cx", &Intrinsics::cx) + .def_ro("cy", &Intrinsics::cy) + .def("get_ray_direction", &Intrinsics::get_ray_direction, + "Get the direction of the ray going through pixel (px, py) in camera frame", + nb::arg("px"), nb::arg("py")); + + nb::class_(camera, "Camera") + .def(nb::init()) + .def_ro("intrinsics", &Camera::intrinsics) + .def_ro("extrinsics", &Camera::extrinsics) + .def("get_ray_direction", &Camera::get_ray_direction, + "Get the direction of the ray going through pixel (px, py) in world frame", + nb::arg("px"), nb::arg("py")); + + /* + * Confidence module bindings */ nb::module_ confidence = m.def_submodule("confidence"); diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index fa887d6..a48252b 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -13,8 +13,6 @@ struct Intrinsics { float fy; float cx; float cy; - float near; - float far; // 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 diff --git a/genmetaballs/src/genmetaballs/core/__init__.py b/genmetaballs/src/genmetaballs/core/__init__.py index 89c845b..2f9f86d 100644 --- a/genmetaballs/src/genmetaballs/core/__init__.py +++ b/genmetaballs/src/genmetaballs/core/__init__.py @@ -3,6 +3,7 @@ FourParameterBlender, ThreeParameterBlender, ) +from genmetaballs._genmetaballs_bindings.camera import Camera, Intrinsics from genmetaballs._genmetaballs_bindings.confidence import ( TwoParameterConfidence, ZeroParameterConfidence, @@ -30,6 +31,8 @@ def array2d_float(data, device) -> CPUFloatArray2D | GPUFloatArray2D: "ZeroParameterConfidence", "TwoParameterConfidence", "geometry", + "Camera", + "Intrinsics", "sigmoid", "FourParameterBlender", "ThreeParameterBlender", diff --git a/tests/python_tests/test_camera.py b/tests/python_tests/test_camera.py new file mode 100644 index 0000000..2bd0bfb --- /dev/null +++ b/tests/python_tests/test_camera.py @@ -0,0 +1,43 @@ +import numpy as np +import pytest + +from genmetaballs.core import Intrinsics +from genmetaballs.fmb.utils import get_camera_rays + + +@pytest.fixture +def rng() -> np.random.Generator: + return np.random.default_rng(0) + + +@pytest.fixture +def intrinsics() -> Intrinsics: + return Intrinsics(height=480, width=640, fx=500.0, fy=520.0, cx=320.0, cy=240.0) + + +def test_get_ray_direction_in_camera_frame(intrinsics: Intrinsics): + # selet a few random pixels to test + pixel_list = np.array( + [ + [0, 0], + [intrinsics.width - 1, 0], + [0, intrinsics.height - 1], + [intrinsics.width - 1, intrinsics.height - 1], + [intrinsics.width // 2, intrinsics.height // 2], + ] + ) + + reference_rays = get_camera_rays( + intrinsics.fx, + intrinsics.fy, + intrinsics.cx, + intrinsics.cy, + np.concatenate([pixel_list, np.zeros((pixel_list.shape[0], 1))], axis=1), + ) + + for (pixel_x, pixel_y), reference_ray in zip(pixel_list, reference_rays, strict=True): + ray_direction = intrinsics.get_ray_direction(pixel_x, pixel_y) + assert np.allclose( + [ray_direction.x, ray_direction.y, ray_direction.z], + reference_ray, + ) From bf69733c998aa94199920f35655c2128cecf0495 Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 00:15:22 -0500 Subject: [PATCH 4/7] Delete extrinsics-related code from this branch --- genmetaballs/src/cuda/bindings.cu | 8 -------- genmetaballs/src/cuda/core/camera.cu | 5 ----- genmetaballs/src/cuda/core/camera.cuh | 9 --------- genmetaballs/src/genmetaballs/core/__init__.py | 2 +- 4 files changed, 1 insertion(+), 23 deletions(-) diff --git a/genmetaballs/src/cuda/bindings.cu b/genmetaballs/src/cuda/bindings.cu index d11e897..d51c30c 100644 --- a/genmetaballs/src/cuda/bindings.cu +++ b/genmetaballs/src/cuda/bindings.cu @@ -82,14 +82,6 @@ NB_MODULE(_genmetaballs_bindings, m) { "Get the direction of the ray going through pixel (px, py) in camera frame", nb::arg("px"), nb::arg("py")); - nb::class_(camera, "Camera") - .def(nb::init()) - .def_ro("intrinsics", &Camera::intrinsics) - .def_ro("extrinsics", &Camera::extrinsics) - .def("get_ray_direction", &Camera::get_ray_direction, - "Get the direction of the ray going through pixel (px, py) in world frame", - nb::arg("px"), nb::arg("py")); - /* * Confidence module bindings */ diff --git a/genmetaballs/src/cuda/core/camera.cu b/genmetaballs/src/cuda/core/camera.cu index 385a261..0f4d124 100644 --- a/genmetaballs/src/cuda/core/camera.cu +++ b/genmetaballs/src/cuda/core/camera.cu @@ -10,8 +10,3 @@ CUDA_CALLABLE Vec3D Intrinsics::get_ray_direction(uint32_t px, uint32_t py) cons auto y = (static_cast(py) - cy) / fy; return Vec3D{x, y, -1.0f}; } - -CUDA_CALLABLE Vec3D Camera::get_ray_direction(uint32_t px, uint32_t py) const { - auto ray_cam = intrinsics.get_ray_direction(px, py); - return extrinsics.get_rot().apply(ray_cam); -} diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index a48252b..c1ffcad 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -18,12 +18,3 @@ struct Intrinsics { // 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; }; - -struct Camera { - Intrinsics intrinsics; - Pose extrinsics; - - // returns the direction of the ray going through pixel (px, py) in world 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; -}; diff --git a/genmetaballs/src/genmetaballs/core/__init__.py b/genmetaballs/src/genmetaballs/core/__init__.py index 2f9f86d..17c4ba5 100644 --- a/genmetaballs/src/genmetaballs/core/__init__.py +++ b/genmetaballs/src/genmetaballs/core/__init__.py @@ -3,7 +3,7 @@ FourParameterBlender, ThreeParameterBlender, ) -from genmetaballs._genmetaballs_bindings.camera import Camera, Intrinsics +from genmetaballs._genmetaballs_bindings.camera import Intrinsics from genmetaballs._genmetaballs_bindings.confidence import ( TwoParameterConfidence, ZeroParameterConfidence, From f4982ed3053f8e5095b10b55cd040cfa79947dd5 Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 00:15:23 -0500 Subject: [PATCH 5/7] Compute directions for a bunch of rays --- genmetaballs/src/cuda/core/camera.cu | 15 +++++++++++++++ genmetaballs/src/cuda/core/camera.cuh | 14 ++++++++++++-- 2 files changed, 27 insertions(+), 2 deletions(-) diff --git a/genmetaballs/src/cuda/core/camera.cu b/genmetaballs/src/cuda/core/camera.cu index 0f4d124..d923abc 100644 --- a/genmetaballs/src/cuda/core/camera.cu +++ b/genmetaballs/src/cuda/core/camera.cu @@ -1,4 +1,5 @@ #include +#include #include #include "camera.cuh" @@ -10,3 +11,17 @@ CUDA_CALLABLE Vec3D Intrinsics::get_ray_direction(uint32_t px, uint32_t py) cons auto y = (static_cast(py) - cy) / fy; return Vec3D{x, y, -1.0f}; } + +template +CUDA_CALLABLE Array2D& Intrinsics::get_ray_directions(Array2D buffer, + uint32_t px_start, + uint32_t px_end, + uint32_t py_start, + uint32_t py_end) const { + for (auto i = max(0, py_start); i < min(height, py_end); ++i) { + for (auto j = max(0, px_start); j < min(width, px_end); ++j) { + buffer[i][j] = get_ray_direction(j, i); + } + } + return buffer; +} diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index c1ffcad..f4d0d95 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -14,7 +14,17 @@ struct Intrinsics { float cx; float cy; - // 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 + // 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 + // For efficiency, this function does not check if Array2D buffer has correct size + template + CUDA_CALLABLE Array2D& get_ray_directions(Array2D buffer, + uint32_t px_start = 0, + uint32_t px_end = UINT32_MAX, + uint32_t py_start = 0, + uint32_t py_end = UINT32_MAX) const; }; From cc09b1b59597a9dad47a83fc2968c753200af0bb Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 01:40:45 -0500 Subject: [PATCH 6/7] Add unit tests for camera rays on GPU --- genmetaballs/src/cuda/core/camera.cu | 14 ------- genmetaballs/src/cuda/core/camera.cuh | 14 +++++-- tests/cpp_tests/test_camera.cu | 53 +++++++++++++++++++++++++++ tests/cpp_tests/test_confidence.cu | 1 - 4 files changed, 63 insertions(+), 19 deletions(-) create mode 100644 tests/cpp_tests/test_camera.cu diff --git a/genmetaballs/src/cuda/core/camera.cu b/genmetaballs/src/cuda/core/camera.cu index d923abc..07c65c1 100644 --- a/genmetaballs/src/cuda/core/camera.cu +++ b/genmetaballs/src/cuda/core/camera.cu @@ -11,17 +11,3 @@ CUDA_CALLABLE Vec3D Intrinsics::get_ray_direction(uint32_t px, uint32_t py) cons auto y = (static_cast(py) - cy) / fy; return Vec3D{x, y, -1.0f}; } - -template -CUDA_CALLABLE Array2D& Intrinsics::get_ray_directions(Array2D buffer, - uint32_t px_start, - uint32_t px_end, - uint32_t py_start, - uint32_t py_end) const { - for (auto i = max(0, py_start); i < min(height, py_end); ++i) { - for (auto j = max(0, px_start); j < min(width, px_end); ++j) { - buffer[i][j] = get_ray_direction(j, i); - } - } - return buffer; -} diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index f4d0d95..d4a0084 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -7,8 +7,8 @@ #include "utils.cuh" struct Intrinsics { - uint32_t height; - uint32_t width; + uint32_t height; // in x direction + uint32_t width; // in y direction float fx; float fy; float cx; @@ -20,11 +20,17 @@ struct Intrinsics { // 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 - // For efficiency, this function does not check if Array2D buffer has correct size template CUDA_CALLABLE Array2D& get_ray_directions(Array2D buffer, uint32_t px_start = 0, uint32_t px_end = UINT32_MAX, uint32_t py_start = 0, - uint32_t py_end = UINT32_MAX) const; + 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; + } }; diff --git a/tests/cpp_tests/test_camera.cu b/tests/cpp_tests/test_camera.cu new file mode 100644 index 0000000..96a793b --- /dev/null +++ b/tests/cpp_tests/test_camera.cu @@ -0,0 +1,53 @@ +#include +#include +#include +#include +#include +#include + +#include "core/camera.cuh" +#include "core/geometry.cuh" +#include "core/utils.cuh" + +namespace test_camera_gpu { + +// CUDA kernel to call get_ray_directions on device with multiple threads +// Each thread processes one row of the image +__global__ void get_ray_directions_kernel(Intrinsics intrinsics, + Array2D 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); +} + +} // namespace test_camera_gpu + +// Test get_ray_directions on GPU (device) +TEST(CameraTest, GetRayDirectionsDevice) { + // Create a small camera intrinsics + Intrinsics intrinsics{4, 6, 100.0f, 100.0f, 3.0f, 2.0f}; + + // Create Array2D buffer on device + thrust::device_vector data(intrinsics.height * intrinsics.width); + Array2D ray_buffer(data.data(), intrinsics.height, intrinsics.width); + + // Launch kernel with multiple threads -- divide into 2x2 tiles + test_camera_gpu:: + get_ray_directions_kernel<<<1, dim3(intrinsics.height / 2, intrinsics.width / 2)>>>( + intrinsics, ray_buffer); + CUDA_CHECK(cudaGetLastError()); + CUDA_CHECK(cudaDeviceSynchronize()); + + // Copy data back to host for sanity check + thrust::host_vector ray_data = data; + + // Sanity check: adjacent rays should be different + constexpr float eps = 1e-6f; + for (auto i = 0; i < data.size() - 1; ++i) { + auto diff = ray_data[i + 1] - ray_data[i]; + float diff_mag = sqrtf(dot(diff, diff)); + EXPECT_GT(diff_mag, eps) << "Adjacent rays are too similar at index " << i; + } +} diff --git a/tests/cpp_tests/test_confidence.cu b/tests/cpp_tests/test_confidence.cu index 2b28b05..bc7b990 100644 --- a/tests/cpp_tests/test_confidence.cu +++ b/tests/cpp_tests/test_confidence.cu @@ -75,7 +75,6 @@ static std::vector confidence_cases() { TEST(GpuConfidenceTest, ConfidenceMultipleValuesGPU_AllTypes) { using test_float = float; - constexpr float rtol = 1e-6F; auto sizes = confidence_test_sizes(); std::mt19937 master_gen(MASTER_SEED); From d26d491f64572fce621f8ab552fb26a904c4f3d7 Mon Sep 17 00:00:00 2001 From: Xiaoyan Wang Date: Tue, 25 Nov 2025 15:54:51 -0500 Subject: [PATCH 7/7] Rename `DeviceType` -> `MemoryLocation` due to changes in `Array2D` --- genmetaballs/src/cuda/core/camera.cuh | 12 ++++++------ tests/cpp_tests/test_camera.cu | 5 +++-- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/genmetaballs/src/cuda/core/camera.cuh b/genmetaballs/src/cuda/core/camera.cuh index d4a0084..d43c490 100644 --- a/genmetaballs/src/cuda/core/camera.cuh +++ b/genmetaballs/src/cuda/core/camera.cuh @@ -20,12 +20,12 @@ struct Intrinsics { // 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 - CUDA_CALLABLE Array2D& get_ray_directions(Array2D buffer, - uint32_t px_start = 0, - uint32_t px_end = UINT32_MAX, - uint32_t py_start = 0, - uint32_t py_end = UINT32_MAX) const { + template + CUDA_CALLABLE Array2D& get_ray_directions(Array2D 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); diff --git a/tests/cpp_tests/test_camera.cu b/tests/cpp_tests/test_camera.cu index 96a793b..1448c8a 100644 --- a/tests/cpp_tests/test_camera.cu +++ b/tests/cpp_tests/test_camera.cu @@ -14,7 +14,7 @@ namespace test_camera_gpu { // CUDA kernel to call get_ray_directions on device with multiple threads // Each thread processes one row of the image __global__ void get_ray_directions_kernel(Intrinsics intrinsics, - Array2D ray_buffer) { + Array2D 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; @@ -31,7 +31,8 @@ TEST(CameraTest, GetRayDirectionsDevice) { // Create Array2D buffer on device thrust::device_vector data(intrinsics.height * intrinsics.width); - Array2D ray_buffer(data.data(), intrinsics.height, intrinsics.width); + Array2D ray_buffer(data.data(), intrinsics.height, + intrinsics.width); // Launch kernel with multiple threads -- divide into 2x2 tiles test_camera_gpu::