clang-format

diff_rast/cuda/csrc/CMakeLists.txt

@@ -54,4 +54,4 @@ target_include_directories(check_serial_backward PRIVATE
 )
 target_include_directories(check_serial_forward PRIVATE
     ${PROJECT_SOURCE_DIR}/third_party/glm
-)
+)

diff_rast/cuda/csrc/backward.cu

@@ -4,7 +4,7 @@
 
 namespace cg = cooperative_groups;
 
-template<int CHANNELS>
+template <int CHANNELS>
 __global__ void rasterize_backward_kernel(
     const dim3 tile_bounds,
     const dim3 img_size,
@@ -49,7 +49,8 @@ __global__ void rasterize_backward_kernel(
     float T_final = final_Ts[pix_id];
     float T = T_final;
     // the contribution from gaussians behind the current one
-    float S[CHANNELS] = {0.f}; // TODO: this currently doesn't match the channel count input.
+    float S[CHANNELS] = {
+        0.f}; // TODO: this currently doesn't match the channel count input.
     // S[0] = 0.0;
     // S[1] = 0.0;
     // S[2] = 0.0;
@@ -96,7 +97,6 @@ __global__ void rasterize_backward_kernel(
             S[c] += rgbs[CHANNELS * g + c] * fac;
         }
 
-
         // v_alpha = (rgb.x * T - S.x * ra) * v_out.x
         //     + (rgb.y * T - S.y * ra) * v_out.y
         //     + (rgb.z * T - S.z * ra) * v_out.z;
@@ -146,7 +146,7 @@ void rasterize_backward_impl(
     float *v_opacity
 
 ) {
-    rasterize_backward_kernel<3> <<<tile_bounds, block>>>(
+    rasterize_backward_kernel<3><<<tile_bounds, block>>>(
         tile_bounds,
         img_size,
         gaussians_ids_sorted,

diff_rast/cuda/csrc/bindings.cu

@@ -31,8 +31,9 @@ __global__ void compute_cov2d_bounds_forward_kernel(
     float3 conic;
     float radius;
     float3 cov2d{
-        (float)covs2d[index], (float)covs2d[index + 1], (float)covs2d[index + 2]
-    };
+        (float)covs2d[index],
+        (float)covs2d[index + 1],
+        (float)covs2d[index + 2]};
     compute_cov2d_bounds(cov2d, conic, radius);
     conics[index] = conic.x;
     conics[index + 1] = conic.y;
@@ -256,21 +257,18 @@ project_gaussians_backward_tensor(
     return std::make_tuple(v_cov2d, v_cov3d, v_mean3d, v_scale, v_quat);
 }
 
-std::tuple<
-    torch::Tensor,
-    torch::Tensor>
-compute_cumulative_intersects_tensor(
-    const int num_points,
-    torch::Tensor &num_tiles_hit
+std::tuple<torch::Tensor, torch::Tensor> compute_cumulative_intersects_tensor(
+    const int num_points, torch::Tensor &num_tiles_hit
 ) {
     // ref:
     // https://nvlabs.github.io/cub/structcub_1_1_device_scan.html#a9416ac1ea26f9fde669d83ddc883795a
     // allocate sum workspace
     CHECK_INPUT(num_tiles_hit);
-
-    torch::Tensor cum_tiles_hit =
-        torch::zeros({num_points}, num_tiles_hit.options().dtype(torch::kInt32));
-
+
+    torch::Tensor cum_tiles_hit = torch::zeros(
+        {num_points}, num_tiles_hit.options().dtype(torch::kInt32)
+    );
+
     int32_t num_intersects;
     compute_cumulative_intersects(
         num_points,
@@ -279,14 +277,15 @@ compute_cumulative_intersects_tensor(
         cum_tiles_hit.contiguous().data_ptr<int32_t>()
     );
 
-    return std::make_tuple(torch::tensor(num_intersects, num_tiles_hit.options().dtype(torch::kInt32)),
-                            cum_tiles_hit);
+    return std::make_tuple(
+        torch::tensor(
+            num_intersects, num_tiles_hit.options().dtype(torch::kInt32)
+        ),
+        cum_tiles_hit
+    );
 }
 
-std::tuple<
-    torch::Tensor,
-    torch::Tensor>
-map_gaussian_to_intersects_tensor(
+std::tuple<torch::Tensor, torch::Tensor> map_gaussian_to_intersects_tensor(
     const int num_points,
     torch::Tensor &xys,
     torch::Tensor &depths,
@@ -310,7 +309,7 @@ map_gaussian_to_intersects_tensor(
         torch::zeros({num_intersects}, xys.options().dtype(torch::kInt32));
     torch::Tensor isect_ids_unsorted =
         torch::zeros({num_intersects}, xys.options().dtype(torch::kInt64));
-    
+
     map_gaussian_to_intersects<<<
         (num_points + N_THREADS - 1) / N_THREADS,
         N_THREADS>>>(

diff_rast/cuda/csrc/bindings.h

@@ -80,18 +80,11 @@ project_gaussians_backward_tensor(
     torch::Tensor &v_conic
 );
 
-std::tuple<
-    torch::Tensor,
-    torch::Tensor>
-compute_cumulative_intersects_tensor(
-    const int num_points,
-    torch::Tensor &num_tiles_hit
+std::tuple<torch::Tensor, torch::Tensor> compute_cumulative_intersects_tensor(
+    const int num_points, torch::Tensor &num_tiles_hit
 );
 
-std::tuple<
-    torch::Tensor,
-    torch::Tensor>
-map_gaussian_to_intersects_tensor(
+std::tuple<torch::Tensor, torch::Tensor> map_gaussian_to_intersects_tensor(
     const int num_points,
     torch::Tensor &xys,
     torch::Tensor &depths,

diff_rast/cuda/csrc/helpers.cuh

@@ -31,7 +31,8 @@ inline __host__ __device__ void get_tile_bbox(
     uint2 &tile_min,
     uint2 &tile_max
 ) {
-    // gets gaussian dimensions in tile space, i.e. the span of a gaussian in tile_grid (image divided into tiles)
+    // gets gaussian dimensions in tile space, i.e. the span of a gaussian in
+    // tile_grid (image divided into tiles)
     float2 tile_center = {
         pix_center.x / (float)BLOCK_X, pix_center.y / (float)BLOCK_Y};
     float2 tile_radius = {
@@ -44,7 +45,8 @@ compute_cov2d_bounds(const float3 cov2d, float3 &conic, float &radius) {
     // find eigenvalues of 2d covariance matrix
     // expects upper triangular values of cov matrix as float3
     // then compute the radius and conic dimensions
-    // the conic is the inverse cov2d matrix, represented here with upper triangular values.  
+    // the conic is the inverse cov2d matrix, represented here with upper
+    // triangular values.
     float det = cov2d.x * cov2d.z - cov2d.y * cov2d.y;
     if (det == 0.f)
         return false;
@@ -64,8 +66,9 @@ compute_cov2d_bounds(const float3 cov2d, float3 &conic, float &radius) {
 }
 
 // compute vjp from df/d_conic to df/c_cov2d
-inline __host__ __device__ void
-cov2d_to_conic_vjp(const float3 &conic, const float3 &v_conic, float3 &v_cov2d) {
+inline __host__ __device__ void cov2d_to_conic_vjp(
+    const float3 &conic, const float3 &v_conic, float3 &v_cov2d
+) {
     // conic = inverse cov2d
     // df/d_cov2d = -conic * df/d_conic * conic
     glm::mat2 X = glm::mat2(conic.x, conic.y, conic.y, conic.z);
@@ -77,7 +80,8 @@ cov2d_to_conic_vjp(const float3 &conic, const float3 &v_conic, float3 &v_cov2d)
 }
 
 // helper for applying R * p + T, expect mat to be ROW MAJOR
-inline __host__ __device__ float3 transform_4x3(const float *mat, const float3 p) {
+inline __host__ __device__ float3
+transform_4x3(const float *mat, const float3 p) {
     float3 out = {
         mat[0] * p.x + mat[1] * p.y + mat[2] * p.z + mat[3],
         mat[4] * p.x + mat[5] * p.y + mat[6] * p.z + mat[7],
@@ -88,7 +92,8 @@ inline __host__ __device__ float3 transform_4x3(const float *mat, const float3 p
 
 // helper to apply 4x4 transform to 3d vector, return homo coords
 // expects mat to be ROW MAJOR
-inline __host__ __device__ float4 transform_4x4(const float *mat, const float3 p) {
+inline __host__ __device__ float4
+transform_4x4(const float *mat, const float3 p) {
     float4 out = {
         mat[0] * p.x + mat[1] * p.y + mat[2] * p.z + mat[3],
         mat[4] * p.x + mat[5] * p.y + mat[6] * p.z + mat[7],
@@ -117,8 +122,7 @@ inline __host__ __device__ float3 project_pix_vjp(
 
     float3 v_ndc = {0.5f * img_size.x * v_xy.x, 0.5f * img_size.y * v_xy.y};
     float4 v_proj = {
-        v_ndc.x * rw, v_ndc.y * rw, 0., -(v_ndc.x + v_ndc.y) * rw * rw
-    };
+        v_ndc.x * rw, v_ndc.y * rw, 0., -(v_ndc.x + v_ndc.y) * rw * rw};
     // df / d_world = df / d_cam * d_cam / d_world
     // = v_proj * P[:3, :3]
     return {
@@ -164,36 +168,36 @@ quat_to_rotmat_vjp(const float4 quat, const glm::mat3 v_R) {
     float4 v_quat;
     // v_R is COLUMN MAJOR
     // w element stored in x field
-    v_quat.x = 2.f * (
-    // v_quat.w = 2.f * (
-        x * (v_R[1][2] - v_R[2][1])
-        + y * (v_R[2][0] - v_R[0][2])
-        + z * (v_R[0][1] - v_R[1][0])
-    );
+    v_quat.x =
+        2.f * (
+                  // v_quat.w = 2.f * (
+                  x * (v_R[1][2] - v_R[2][1]) + y * (v_R[2][0] - v_R[0][2]) +
+                  z * (v_R[0][1] - v_R[1][0])
+              );
     // x element in y field
-    v_quat.y = 2.f * (
-    // v_quat.x = 2.f * (
-        -2.f * x * (v_R[1][1] + v_R[2][2])
-        + y * (v_R[0][1] + v_R[1][0])
-        + z * (v_R[0][2] + v_R[2][0])
-        + w * (v_R[1][2] - v_R[2][1])
-    );
+    v_quat.y =
+        2.f *
+        (
+            // v_quat.x = 2.f * (
+            -2.f * x * (v_R[1][1] + v_R[2][2]) + y * (v_R[0][1] + v_R[1][0]) +
+            z * (v_R[0][2] + v_R[2][0]) + w * (v_R[1][2] - v_R[2][1])
+        );
     // y element in z field
-    v_quat.z = 2.f * (
-    // v_quat.y = 2.f * (
-        x * (v_R[0][1] + v_R[1][0])
-        - 2.f * y * (v_R[0][0] + v_R[2][2])
-        + z * (v_R[1][2] + v_R[2][1])
-        + w * (v_R[2][0] - v_R[0][2])
-    );
+    v_quat.z =
+        2.f *
+        (
+            // v_quat.y = 2.f * (
+            x * (v_R[0][1] + v_R[1][0]) - 2.f * y * (v_R[0][0] + v_R[2][2]) +
+            z * (v_R[1][2] + v_R[2][1]) + w * (v_R[2][0] - v_R[0][2])
+        );
     // z element in w field
-    v_quat.w = 2.f * (
-    // v_quat.z = 2.f * (
-        x * (v_R[0][2] + v_R[2][0])
-        + y * (v_R[1][2] + v_R[2][1])
-        - 2.f * z * (v_R[0][0] + v_R[1][1])
-        + w * (v_R[0][1] - v_R[1][0])
-    );
+    v_quat.w =
+        2.f *
+        (
+            // v_quat.z = 2.f * (
+            x * (v_R[0][2] + v_R[2][0]) + y * (v_R[1][2] + v_R[2][1]) -
+            2.f * z * (v_R[0][0] + v_R[1][1]) + w * (v_R[0][1] - v_R[1][0])
+        );
     return v_quat;
 }
 
@@ -207,8 +211,9 @@ scale_to_mat(const float3 scale, const float glob_scale) {
 }
 
 // device helper for culling near points
-inline __host__ __device__ bool
-clip_near_plane(const float3 p, const float *viewmat, float3 &p_view, float thresh) {
+inline __host__ __device__ bool clip_near_plane(
+    const float3 p, const float *viewmat, float3 &p_view, float thresh
+) {
     p_view = transform_4x3(viewmat, p);
     if (p_view.z <= thresh) {
         return true;

diff_rast/cuda/csrc/serial_backward.cu

@@ -164,8 +164,8 @@ computeConicBackward(const float3 &cov2D, const float3 &dL_dconic) {
     float denom2inv = 1.0f / ((denom * denom) + 0.0000001f);
 
     if (denom2inv != 0) {
-        // This is slightly different from the original implementation, but we include this line to make
-        // equality checks easier.
+        // This is slightly different from the original implementation, but we
+        // include this line to make equality checks easier.
         float denom2inv = 1.0f / (denom * denom);
         // Gradients of loss w.r.t. entries of 2D covariance matrix,
         // given gradients of loss w.r.t. conic matrix (inverse covariance

diff_rast/cuda/csrc/serial_backward.cuh

@@ -1,31 +1,28 @@
 #include "cuda_runtime.h"
 
-
 __host__ __device__ float3 projectMean2DBackward(
-    const float3 m, const float* proj, const float2 dL_dmean2D
+    const float3 m, const float *proj, const float2 dL_dmean2D
 );
 
 __host__ __device__ void computeCov3DBackward(
     const float3 scale,
     const float mod,
     const float4 rot,
-    const float* dL_dcov3D,
+    const float *dL_dcov3D,
     float3 &dL_dscale,
     float4 &dL_dq
 );
 
-__host__ __device__ float3 computeConicBackward(
-    const float3 &cov2D,
-    const float3 &dL_dconic
-);
+__host__ __device__ float3
+computeConicBackward(const float3 &cov2D, const float3 &dL_dconic);
 
 __host__ __device__ void computeCov2DBackward(
     const float3 &mean,
     const float *cov3D,
     const float *view_matrix,
     const float h_x,
     const float h_y,
-	const float tan_fovx,
+    const float tan_fovx,
     const float tan_fovy,
     const float3 &dL_dcov2d,
     float3 &dL_dmean,

diff_rast/cuda/csrc/sh.cuh

@@ -10,17 +10,15 @@ __host__ __device__ const float SH_C2[] = {
     -1.0925484305920792f,
     0.31539156525252005f,
     -1.0925484305920792f,
-    0.5462742152960396f
-};
+    0.5462742152960396f};
 __host__ __device__ const float SH_C3[] = {
     -0.5900435899266435f,
     2.890611442640554f,
     -0.4570457994644658f,
     0.3731763325901154f,
     -0.4570457994644658f,
     1.445305721320277f,
-    -0.5900435899266435f
-};
+    -0.5900435899266435f};
 __host__ __device__ const float SH_C4[] = {
     2.5033429417967046f,
     -1.7701307697799304,
@@ -30,8 +28,7 @@ __host__ __device__ const float SH_C4[] = {
     -0.6690465435572892f,
     0.47308734787878004f,
     -1.7701307697799304f,
-    0.6258357354491761f
-};
+    0.6258357354491761f};
 
 __host__ __device__ unsigned num_sh_bases(const unsigned degree) {
     if (degree == 0)