From bcdbbc4864bd249d81efba8f8797729ce54c26ce Mon Sep 17 00:00:00 2001 From: maturk Date: Tue, 3 Oct 2023 11:36:35 +0300 Subject: [PATCH] clang-format --- diff_rast/cuda/csrc/CMakeLists.txt | 2 +- diff_rast/cuda/csrc/backward.cu | 8 +-- diff_rast/cuda/csrc/bindings.cu | 37 ++++++------ diff_rast/cuda/csrc/bindings.h | 13 +--- diff_rast/cuda/csrc/helpers.cuh | 79 +++++++++++++------------ diff_rast/cuda/csrc/serial_backward.cu | 4 +- diff_rast/cuda/csrc/serial_backward.cuh | 13 ++-- diff_rast/cuda/csrc/sh.cuh | 9 +-- 8 files changed, 78 insertions(+), 87 deletions(-) diff --git a/diff_rast/cuda/csrc/CMakeLists.txt b/diff_rast/cuda/csrc/CMakeLists.txt index 24cbc8364..12a4ba8e3 100644 --- a/diff_rast/cuda/csrc/CMakeLists.txt +++ b/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 -) +) \ No newline at end of file diff --git a/diff_rast/cuda/csrc/backward.cu b/diff_rast/cuda/csrc/backward.cu index a94a265cb..fedc475b1 100644 --- a/diff_rast/cuda/csrc/backward.cu +++ b/diff_rast/cuda/csrc/backward.cu @@ -4,7 +4,7 @@ namespace cg = cooperative_groups; -template +template __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> <<>>( + rasterize_backward_kernel<3><<>>( tile_bounds, img_size, gaussians_ids_sorted, diff --git a/diff_rast/cuda/csrc/bindings.cu b/diff_rast/cuda/csrc/bindings.cu index 9b5cc8753..672384b25 100644 --- a/diff_rast/cuda/csrc/bindings.cu +++ b/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 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() ); - 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 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 --git a/diff_rast/cuda/csrc/bindings.h b/diff_rast/cuda/csrc/bindings.h index 852004824..574c121de 100644 --- a/diff_rast/cuda/csrc/bindings.h +++ b/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 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 map_gaussian_to_intersects_tensor( const int num_points, torch::Tensor &xys, torch::Tensor &depths, diff --git a/diff_rast/cuda/csrc/helpers.cuh b/diff_rast/cuda/csrc/helpers.cuh index 884857d94..962c29dfc 100644 --- a/diff_rast/cuda/csrc/helpers.cuh +++ b/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 --git a/diff_rast/cuda/csrc/serial_backward.cu b/diff_rast/cuda/csrc/serial_backward.cu index 14a28d3d8..a902b3234 100644 --- a/diff_rast/cuda/csrc/serial_backward.cu +++ b/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 --git a/diff_rast/cuda/csrc/serial_backward.cuh b/diff_rast/cuda/csrc/serial_backward.cuh index d26d38c60..58b5930f3 100644 --- a/diff_rast/cuda/csrc/serial_backward.cuh +++ b/diff_rast/cuda/csrc/serial_backward.cuh @@ -1,23 +1,20 @@ #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, @@ -25,7 +22,7 @@ __host__ __device__ void computeCov2DBackward( 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 --git a/diff_rast/cuda/csrc/sh.cuh b/diff_rast/cuda/csrc/sh.cuh index eecfb654d..f709e155d 100644 --- a/diff_rast/cuda/csrc/sh.cuh +++ b/diff_rast/cuda/csrc/sh.cuh @@ -10,8 +10,7 @@ __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, @@ -19,8 +18,7 @@ __host__ __device__ const float SH_C3[] = { 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)