Skip to content

Commit 7e85a0e

Browse files
fwilliamsfwilliams
committed
orthographic
Signed-off-by: Francis Williams <francis@fwilliams.info>
1 parent ba96359 commit 7e85a0e

File tree

3 files changed

+107
-32
lines changed

3 files changed

+107
-32
lines changed

‎src/fvdb/detail/ops/gsplat/GaussianProjectionUT.cu‎

Lines changed: 29 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,17 @@ template <typename T> class OpenCVCameraModel {
6060
: K(K_in) {
6161
deviceAssertOrTrap(K != nullptr);
6262

63+
// Orthographic is a camera model choice; it is only supported for the undistorted pinhole
64+
// math path (no OpenCV distortion).
65+
orthographic = (cameraModel == CameraModel::ORTHOGRAPHIC);
66+
67+
if (cameraModel == CameraModel::ORTHOGRAPHIC) {
68+
radial = tangential = thinPrism = nullptr;
69+
numRadial = numTangential = numThinPrism = 0;
70+
model = Model::NONE;
71+
return;
72+
}
73+
6374
if (cameraModel == CameraModel::PINHOLE) {
6475
radial = tangential = thinPrism = nullptr;
6576
numRadial = numTangential = numThinPrism = 0;
@@ -120,12 +131,17 @@ template <typename T> class OpenCVCameraModel {
120131
}
121132

122133
// Project a 3D point in camera coordinates to pixel coordinates using this camera model
123-
// (pinhole + distortion + intrinsics).
134+
// (pinhole/orthographic + distortion + intrinsics).
124135
__device__ Vec2
125136
project(const Vec3 &p_cam) const {
126-
// Normalize by depth.
127-
const T z_inv = T(1) / max(p_cam[2], T(1e-6));
128-
const Vec2 p_normalized(p_cam[0] * z_inv, p_cam[1] * z_inv);
137+
// Normalize to camera plane.
138+
Vec2 p_normalized;
139+
if (orthographic) {
140+
p_normalized = Vec2(p_cam[0], p_cam[1]);
141+
} else {
142+
const T z_inv = T(1) / max(p_cam[2], T(1e-6));
143+
p_normalized = Vec2(p_cam[0] * z_inv, p_cam[1] * z_inv);
144+
}
129145

130146
const Vec2 p_distorted = applyDistortion(p_normalized);
131147

@@ -150,7 +166,8 @@ template <typename T> class OpenCVCameraModel {
150166

151167
// Camera intrinsics pointer (typically points into shared memory).
152168
// This avoids copying a Mat3 into registers per-thread.
153-
const Mat3 *K = nullptr;
169+
const Mat3 *K = nullptr;
170+
bool orthographic = false;
154171

155172
// Coefficients for the distortion model.
156173
const T *radial = nullptr; // k1..k6 (but k4..k6 only used in rational model)
@@ -283,7 +300,8 @@ template <typename ScalarType> struct WorldToPixelTransform {
283300
const RigidTransform<ScalarType> &xf,
284301
Vec2 &out_pix) const {
285302
const Vec3 p_cam = xf.apply(p_world);
286-
// Perspective only (ortho is not meaningful for distorted camera models).
303+
// Note: `CameraModel::ORTHOGRAPHIC` still uses the same behind-camera check, but does not
304+
// divide by depth in the camera model projection step.
287305
if (p_cam[2] <= ScalarType(0)) {
288306
// Ensure deterministic output to avoid UB on callers that assign/read even on invalid
289307
// projections. This value is ignored when we treat BehindCamera as a hard reject.
@@ -803,8 +821,7 @@ dispatchGaussianProjectionForwardUT<torch::kCUDA>(
803821
const float nearPlane,
804822
const float farPlane,
805823
const float minRadius2d,
806-
const bool calcCompensations,
807-
const bool ortho) {
824+
const bool calcCompensations) {
808825
FVDB_FUNC_RANGE();
809826

810827
TORCH_CHECK_VALUE(means.is_cuda(), "means must be a CUDA tensor");
@@ -817,7 +834,7 @@ dispatchGaussianProjectionForwardUT<torch::kCUDA>(
817834
TORCH_CHECK_VALUE(projectionMatrices.is_cuda(), "projectionMatrices must be a CUDA tensor");
818835
TORCH_CHECK_VALUE(distortionCoeffs.is_cuda(), "distortionCoeffs must be a CUDA tensor");
819836
TORCH_CHECK_VALUE(distortionCoeffs.dim() == 2, "distortionCoeffs must be 2D");
820-
if (cameraModel == CameraModel::PINHOLE) {
837+
if (cameraModel == CameraModel::PINHOLE || cameraModel == CameraModel::ORTHOGRAPHIC) {
821838
// Accept any K (including 0); ignored.
822839
} else if (cameraModel == CameraModel::OPENCV_RADTAN_5 ||
823840
cameraModel == CameraModel::OPENCV_RATIONAL_8 ||
@@ -857,10 +874,7 @@ dispatchGaussianProjectionForwardUT<torch::kCUDA>(
857874
using scalar_t = float;
858875

859876
const size_t NUM_BLOCKS = GET_BLOCKS(C * N, 256);
860-
// This kernel currently implements the (distorted) perspective camera model.
861-
// Keep parity with the reference kernel: orthographic is not supported here.
862-
TORCH_CHECK_VALUE(!ortho,
863-
"GaussianProjectionForwardUT does not support orthographic projection");
877+
// Orthographic is supported only for CameraModel::ORTHOGRAPHIC (undistorted).
864878

865879
const size_t SHARED_MEM_SIZE = C * (3 * sizeof(nanovdb::math::Mat3<scalar_t>) +
866880
2 * sizeof(nanovdb::math::Vec3<scalar_t>)) +
@@ -915,8 +929,7 @@ dispatchGaussianProjectionForwardUT<torch::kCPU>(
915929
const float nearPlane,
916930
const float farPlane,
917931
const float minRadius2d,
918-
const bool calcCompensations,
919-
const bool ortho) {
932+
const bool calcCompensations) {
920933
TORCH_CHECK_NOT_IMPLEMENTED(false, "GaussianProjectionForwardUT not implemented on the CPU");
921934
}
922935

@@ -939,8 +952,7 @@ dispatchGaussianProjectionForwardUT<torch::kPrivateUse1>(
939952
const float nearPlane,
940953
const float farPlane,
941954
const float minRadius2d,
942-
const bool calcCompensations,
943-
const bool ortho) {
955+
const bool calcCompensations) {
944956
TORCH_CHECK_NOT_IMPLEMENTED(false,
945957
"GaussianProjectionForwardUT not implemented for this device type");
946958
}

‎src/fvdb/detail/ops/gsplat/GaussianProjectionUT.h‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,9 @@ enum class CameraModel : int32_t {
2525
// Pinhole intrinsics only (no distortion).
2626
PINHOLE = 0,
2727

28+
// Orthographic intrinsics (no distortion).
29+
ORTHOGRAPHIC = 5,
30+
2831
// OpenCV variants which are just pinhole intrinsics + optional distortion (all of them use the
2932
// same [C,12] distortion coefficients layout: [k1,k2,k3,k4,k5,k6,p1,p2,s1,s2,s3,s4]).
3033
OPENCV_RADTAN_5 = 1, // polynomial radial (k1,k2,k3) + tangential (p1,p2)).
@@ -83,6 +86,7 @@ struct UTParams {
8386
/// @param[in] cameraModel Camera model used to interpret `distortionCoeffs`.
8487
/// @param[in] distortionCoeffs Distortion coefficients for each camera.
8588
/// - CameraModel::PINHOLE: ignored (use [C,0] or [C,K] tensor).
89+
/// - CameraModel::ORTHOGRAPHIC: ignored (use [C,0] or [C,K] tensor).
8690
/// - CameraModel::OPENCV_*: expects [C,12] coefficients in the following order:
8791
/// [k1,k2,k3,k4,k5,k6,p1,p2,s1,s2,s3,s4]
8892
/// where k1..k6 are radial (rational), p1,p2 are tangential, and s1..s4 are thin-prism.
@@ -94,7 +98,6 @@ struct UTParams {
9498
/// @param[in] minRadius2d Minimum 2D radius threshold; Gaussians with projected radius <= this
9599
/// value are clipped/discarded
96100
/// @param[in] calcCompensations Whether to calculate view-dependent compensation factors
97-
/// @param[in] ortho Whether to use orthographic projection instead of perspective
98101
///
99102
/// @return std::tuple containing:
100103
/// - Radii of 2D Gaussians [C, N]
@@ -114,15 +117,14 @@ dispatchGaussianProjectionForwardUT(
114117
const RollingShutterType rollingShutterType,
115118
const UTParams &utParams,
116119
const CameraModel cameraModel,
117-
const torch::Tensor &distortionCoeffs, // [C, 12] for OPENCV_*, or [C, 0] for NONE
120+
const torch::Tensor &distortionCoeffs, // [C, 12] for OPENCV_*, or [C, 0] for PINHOLE/ORTHO
118121
const int64_t imageWidth,
119122
const int64_t imageHeight,
120123
const float eps2d,
121124
const float nearPlane,
122125
const float farPlane,
123126
const float minRadius2d,
124-
const bool calcCompensations,
125-
const bool ortho);
127+
const bool calcCompensations);
126128

127129
} // namespace ops
128130
} // namespace detail

‎src/tests/GaussianProjectionUTTest.cpp‎

Lines changed: 72 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -189,7 +189,6 @@ TEST_F(GaussianProjectionUTTestFixture, CenteredGaussian_NoDistortion_AnalyticMe
189189
nearPlane,
190190
farPlane,
191191
minRadius2d,
192-
false,
193192
false);
194193

195194
auto means2d_cpu = means2d.cpu();
@@ -219,6 +218,78 @@ TEST_F(GaussianProjectionUTTestFixture, CenteredGaussian_NoDistortion_AnalyticMe
219218
EXPECT_NEAR(conics_cpu[0][0][2].item<float>(), expected_c, 1e-3f);
220219
}
221220

221+
TEST_F(GaussianProjectionUTTestFixture, Orthographic_NoDistortion_AnalyticMeanAndDepth) {
222+
const int64_t C = 1;
223+
224+
const float x = 1.0f, y = -2.0f, z = 10.0f;
225+
means = torch::tensor({{x, y, z}}, torch::kFloat32);
226+
quats = torch::tensor({{1.0f, 0.0f, 0.0f, 0.0f}}, torch::kFloat32);
227+
logScales = torch::log(torch::tensor({{0.2f, 0.3f, 0.4f}}, torch::kFloat32));
228+
229+
worldToCamMatricesStart =
230+
torch::eye(4, torch::TensorOptions().dtype(torch::kFloat32)).unsqueeze(0).expand({C, 4, 4});
231+
worldToCamMatricesEnd = worldToCamMatricesStart.clone();
232+
233+
const float fx = 123.0f, fy = 77.0f, cx = 320.0f, cy = 240.0f;
234+
projectionMatrices = torch::zeros({C, 3, 3}, torch::TensorOptions().dtype(torch::kFloat32));
235+
auto projectionMatricesAcc = projectionMatrices.accessor<float, 3>();
236+
projectionMatricesAcc[0][0][0] = fx;
237+
projectionMatricesAcc[0][1][1] = fy;
238+
projectionMatricesAcc[0][0][2] = cx;
239+
projectionMatricesAcc[0][1][2] = cy;
240+
projectionMatricesAcc[0][2][2] = 1.0f;
241+
242+
cameraModel = CameraModel::ORTHOGRAPHIC;
243+
distortionCoeffs = torch::zeros({C, 0}, torch::kFloat32);
244+
245+
imageWidth = 640;
246+
imageHeight = 480;
247+
eps2d = 0.3f;
248+
nearPlane = 0.1f;
249+
farPlane = 100.0f;
250+
minRadius2d = 0.0f;
251+
252+
utParams = UTParams{};
253+
utParams.inImageMargin = 0.1f;
254+
utParams.requireAllSigmaPointsInImage = true;
255+
256+
means = means.cuda();
257+
quats = quats.cuda();
258+
logScales = logScales.cuda();
259+
worldToCamMatricesStart = worldToCamMatricesStart.cuda();
260+
worldToCamMatricesEnd = worldToCamMatricesEnd.cuda();
261+
projectionMatrices = projectionMatrices.cuda();
262+
distortionCoeffs = distortionCoeffs.cuda();
263+
264+
const auto [radii, means2d, depths, conics, compensations] =
265+
dispatchGaussianProjectionForwardUT<torch::kCUDA>(means,
266+
quats,
267+
logScales,
268+
worldToCamMatricesStart,
269+
worldToCamMatricesEnd,
270+
projectionMatrices,
271+
RollingShutterType::NONE,
272+
utParams,
273+
cameraModel,
274+
distortionCoeffs,
275+
imageWidth,
276+
imageHeight,
277+
eps2d,
278+
nearPlane,
279+
farPlane,
280+
minRadius2d,
281+
false);
282+
283+
auto means2d_cpu = means2d.cpu();
284+
auto depths_cpu = depths.cpu();
285+
auto radii_cpu = radii.cpu();
286+
287+
EXPECT_GT(radii_cpu[0][0].item<int32_t>(), 0);
288+
EXPECT_NEAR(depths_cpu[0][0].item<float>(), z, 1e-4f);
289+
EXPECT_NEAR(means2d_cpu[0][0][0].item<float>(), fx * x + cx, 1e-3f);
290+
EXPECT_NEAR(means2d_cpu[0][0][1].item<float>(), fy * y + cy, 1e-3f);
291+
}
292+
222293
TEST_F(GaussianProjectionUTTestFixture, OffAxisTinyGaussian_NoDistortion_MeanMatchesPinhole) {
223294
const int64_t C = 1;
224295

@@ -283,7 +354,6 @@ TEST_F(GaussianProjectionUTTestFixture, OffAxisTinyGaussian_NoDistortion_MeanMat
283354
nearPlane,
284355
farPlane,
285356
minRadius2d,
286-
false,
287357
false);
288358

289359
auto means2d_cpu = means2d.cpu();
@@ -379,7 +449,6 @@ TEST_F(GaussianProjectionUTTestFixture, MultiCamera_RadTanDistortion_PerCameraPa
379449
nearPlane,
380450
farPlane,
381451
minRadius2d,
382-
false,
383452
false);
384453

385454
auto radii_cpu = radii.cpu();
@@ -475,7 +544,6 @@ TEST_F(GaussianProjectionUTTestFixture,
475544
nearPlane,
476545
farPlane,
477546
minRadius2d,
478-
false,
479547
false);
480548

481549
auto radii_cpu = radii.cpu();
@@ -569,7 +637,6 @@ TEST_F(GaussianProjectionUTTestFixture,
569637
nearPlane,
570638
farPlane,
571639
minRadius2d,
572-
false,
573640
false);
574641

575642
auto radii_cpu = radii.cpu();
@@ -671,7 +738,6 @@ TEST_F(GaussianProjectionUTTestFixture,
671738
nearPlane,
672739
farPlane,
673740
minRadius2d,
674-
false,
675741
false);
676742

677743
auto radii_cpu = radii.cpu();
@@ -768,7 +834,6 @@ TEST_F(GaussianProjectionUTTestFixture,
768834
nearPlane,
769835
farPlane,
770836
minRadius2d,
771-
false,
772837
false);
773838

774839
auto radii_cpu = radii.cpu();
@@ -837,7 +902,6 @@ TEST_F(GaussianProjectionUTTestFixture, RadTanThinPrism_IgnoresK456EvenIfNonZero
837902
nearPlane,
838903
farPlane,
839904
minRadius2d,
840-
false,
841905
false);
842906

843907
auto radii_cpu = radii.cpu();
@@ -922,7 +986,6 @@ TEST_F(GaussianProjectionUTTestFixture,
922986
nearPlane,
923987
farPlane,
924988
minRadius2d,
925-
false,
926989
false);
927990

928991
// When the UT kernel discards a Gaussian, only radii are defined to be 0; other outputs are
@@ -997,7 +1060,6 @@ TEST_F(GaussianProjectionUTTestFixture,
9971060
nearPlane,
9981061
farPlane,
9991062
minRadius2d,
1000-
false,
10011063
false);
10021064

10031065
auto radii_cpu = radii.cpu();
@@ -1069,7 +1131,6 @@ TEST_F(GaussianProjectionUTTestFixture, RollingShutterNone_DepthUsesStartPoseNot
10691131
nearPlane,
10701132
farPlane,
10711133
minRadius2d,
1072-
false,
10731134
false);
10741135

10751136
auto depths_cpu = depths.cpu();

0 commit comments

Comments
 (0)