[Fix] clang lint

.github/workflows/lint.yml

@@ -22,7 +22,7 @@ jobs:
       - name: Linting
         run: pre-commit run --all-files
       - name: Format c/cuda codes with clang-format
-        uses: DoozyX/clang-format-lint-action@v0.13
+        uses: DoozyX/clang-format-lint-action@v0.18
         with:
           source: mmcv/ops/csrc
           extensions: h,c,cpp,hpp,cu,cuh

mmcv/ops/csrc/common/cuda/bezier_align_cuda_kernel.cuh

@@ -222,9 +222,9 @@ __global__ void bezier_align_backward_cuda_kernel(
           atomicAdd(offset_bottom_diff + y_high * width + x_high,
                     static_cast<T>(g4));
         }  // if
-      }    // ix
-    }      // iy
-  }        // CUDA_1D_KERNEL_LOOP
+      }  // ix
+    }  // iy
+  }  // CUDA_1D_KERNEL_LOOP
 }  // BezierAlignBackward
 
 #endif  // BEZIER_ALIGN_CUDA_KERNEL_CUH

mmcv/ops/csrc/common/cuda/riroi_align_rotated_cuda_kernel.cuh

@@ -234,9 +234,9 @@ __global__ void riroi_align_rotated_backward_cuda_kernel(
                     g4 * l_var);
 
         }  // if
-      }    // ix
-    }      // iy
-  }        // CUDA_1D_KERNEL_LOOP
+      }  // ix
+    }  // iy
+  }  // CUDA_1D_KERNEL_LOOP
 }  // RiRoIAlignBackward
 
 #endif  // RIROI_ALIGN_ROTATED_CUDA_KERNEL_CUH

mmcv/ops/csrc/common/cuda/roi_align_rotated_cuda_kernel.cuh

@@ -194,9 +194,9 @@ __global__ void roi_align_rotated_backward_cuda_kernel(
           atomicAdd(offset_bottom_diff + y_high * width + x_low, g3);
           atomicAdd(offset_bottom_diff + y_high * width + x_high, g4);
         }  // if
-      }    // ix
-    }      // iy
-  }        // CUDA_1D_KERNEL_LOOP
+      }  // ix
+    }  // iy
+  }  // CUDA_1D_KERNEL_LOOP
 }  // RoIAlignBackward
 
 #endif  // ROI_ALIGN_ROTATED_CUDA_KERNEL_CUH

mmcv/ops/csrc/common/mlu/common_mlu_helper.hpp

@@ -33,7 +33,7 @@
 #define PAD_DOWN(x, y) (((x) / (y)) * (y))
 #endif
 
-#define CEIL_ALIGN(x, y) (((x) + (y)-1) / (y) * (y))
+#define CEIL_ALIGN(x, y) (((x) + (y) - 1) / (y) * (y))
 
 template <typename scalar_t>
 __mlu_func__ inline scalar_t min(scalar_t a, scalar_t b) {

mmcv/ops/csrc/common/pytorch_mlu_helper.hpp

@@ -21,9 +21,9 @@
 
 #define PAD_DOWN(x, y) (((x) / (y)) * (y))
 
-#define CEIL_DIV(x, y) (((x) + (y)-1) / (y))
+#define CEIL_DIV(x, y) (((x) + (y) - 1) / (y))
 
-#define CEIL_ALIGN(x, y) (((x) + (y)-1) / (y) * (y))
+#define CEIL_ALIGN(x, y) (((x) + (y) - 1) / (y) * (y))
 
 inline int32_t getJobLimitCapability() {
   CNcontext drv_ctx;

mmcv/ops/csrc/common/pytorch_npu_util.hpp

@@ -468,7 +468,7 @@ void ReleaseConvertTypes(Tuple &t) {
 }
 
 template <typename... Ts>
-constexpr auto ConvertTypes(Ts &... args) {
+constexpr auto ConvertTypes(Ts &...args) {
   return std::make_tuple(ConvertType(args)...);
 }
 
@@ -506,7 +506,7 @@ void AddParamToBuf(const string &);
 void AddParamToBuf();
 
 template <typename T, typename... Args>
-void AddParamToBuf(const T &arg, Args &... args) {
+void AddParamToBuf(const T &arg, Args &...args) {
   AddParamToBuf(arg);
   AddParamToBuf(args...);
 }

mmcv/ops/csrc/common/utils/spconv/paramsgrid.h

@@ -40,7 +40,7 @@ void assigner(TT &src, std::vector<int> counter, std::vector<scalar_t> &arg) {
 
 template <int Idx, class TT, class scalar_t, class... TArgs>
 void assigner(TT &src, std::vector<int> counter, std::vector<scalar_t> &arg,
-              std::vector<TArgs> &... args) {
+              std::vector<TArgs> &...args) {
   std::get<Idx>(src) = arg[counter[Idx]];
   assigner<Idx + 1>(src, counter, args...);
 }

mmcv/ops/csrc/common/utils/spconv/tensorview/tensorview.h

@@ -108,7 +108,7 @@ struct CPU {};
 template <typename scalar_t, size_t MaxDim = TV_MAX_DIM>
 struct SimpleVector {
  public:
-  TV_HOST_DEVICE_INLINE SimpleVector(){};
+  TV_HOST_DEVICE_INLINE SimpleVector() {};
   TV_HOST_DEVICE_INLINE SimpleVector(std::initializer_list<scalar_t> q) {
     TV_ASSERT(q.size() <= MaxDim);
     mSize = 0;
@@ -315,7 +315,7 @@ struct Slice {
 
 template <size_t MaxDim = TV_MAX_DIM>
 struct ShapeBase : public SimpleVector<int, MaxDim> {
-  TV_HOST_DEVICE_INLINE ShapeBase() : SimpleVector<int, MaxDim>(){};
+  TV_HOST_DEVICE_INLINE ShapeBase() : SimpleVector<int, MaxDim>() {};
   TV_HOST_DEVICE_INLINE ShapeBase(std::initializer_list<int> shape)
       : SimpleVector<int, MaxDim>(shape) {}
 

mmcv/ops/csrc/pytorch/cpu/bezier_align.cpp

@@ -220,9 +220,9 @@ void BezierAlignForward(const int nthreads, const T *input, const T *rois,
 
           output[index] = output_val;
         }  // for pw
-      }    // for ph
-    }      // for c
-  }        // for n
+      }  // for ph
+    }  // for c
+  }  // for n
 }
 
 template <typename T>
@@ -381,9 +381,9 @@ void BezierAlignBackward(const int nthreads, const T *grad_output,
           add(offset_grad_input + y_high * width + x_low, static_cast<T>(g3));
           add(offset_grad_input + y_high * width + x_high, static_cast<T>(g4));
         }  // if
-      }    // ix
-    }      // iy
-  }        // for
+      }  // ix
+    }  // iy
+  }  // for
 }  // BezierAlignBackward
 
 void BezierAlignForwardCPULauncher(Tensor input, Tensor rois, Tensor output,

mmcv/ops/csrc/pytorch/cpu/roi_align.cpp

@@ -207,10 +207,10 @@ void ROIAlignForward(const int nthreads, const T* input, const T* rois,
             // We do average (integral) pooling inside a bin
             output[index] = output_val / count;
           }  // if
-        }    // for pw
-      }      // for ph
-    }        // for c
-  }          // for n
+        }  // for pw
+      }  // for ph
+    }  // for c
+  }  // for n
 }
 
 template <typename T>
@@ -334,7 +334,7 @@ void ROIAlignBackward(const int nthreads, const T* grad_output, const T* rois,
           add(offset_grad_input + y_high * width + x_low, static_cast<T>(g3));
           add(offset_grad_input + y_high * width + x_high, static_cast<T>(g4));
         }  // if
-      }    // mode
+      }  // mode
     } else if (pool_mode == 1) {
       // We do average (integral) pooling inside a bin
       // We use roi_bin_grid to sample the grid and mimic integral
@@ -375,10 +375,10 @@ void ROIAlignBackward(const int nthreads, const T* grad_output, const T* rois,
             add(offset_grad_input + y_high * width + x_high,
                 static_cast<T>(g4));
           }  // if
-        }    // ix
-      }      // iy
-    }        // mode
-  }          // for
+        }  // ix
+      }  // iy
+    }  // mode
+  }  // for
 }  // ROIAlignBackward
 
 void ROIAlignForwardCPULauncher(Tensor input, Tensor rois, Tensor output,

mmcv/ops/csrc/pytorch/cpu/roi_align_rotated.cpp

@@ -206,9 +206,9 @@ void ROIAlignRotatedForward(const int nthreads, const T* input,
 
           output[index] = output_val;
         }  // for pw
-      }    // for ph
-    }      // for c
-  }        // for n
+      }  // for ph
+    }  // for c
+  }  // for n
 }
 
 template <typename T>
@@ -366,9 +366,9 @@ void ROIAlignRotatedBackward(
           add(offset_grad_input + y_high * width + x_low, static_cast<T>(g3));
           add(offset_grad_input + y_high * width + x_high, static_cast<T>(g4));
         }  // if
-      }    // ix
-    }      // iy
-  }        // for
+      }  // ix
+    }  // iy
+  }  // for
 }  // ROIAlignRotatedBackward
 
 void ROIAlignRotatedForwardCPULauncher(Tensor input, Tensor rois, Tensor output,

mmcv/ops/csrc/pytorch/cuda/filtered_lrelu.cu

@@ -100,8 +100,9 @@ void *choose_filtered_lrelu_act_kernel(void);
 //------------------------------------------------------------------------
 // Helpers.
 
-enum              // Filter modes.
-{ MODE_SUSD = 0,  // Separable upsampling, separable downsampling.
+enum  // Filter modes.
+{
+  MODE_SUSD = 0,  // Separable upsampling, separable downsampling.
   MODE_FUSD = 1,  // Full upsampling, separable downsampling.
   MODE_SUFD = 2,  // Separable upsampling, full downsampling.
   MODE_FUFD = 3,  // Full upsampling, full downsampling.
@@ -157,12 +158,11 @@ struct InternalType<c10::Half> {
 
 #define MIN(A, B) ((A) < (B) ? (A) : (B))
 #define MAX(A, B) ((A) > (B) ? (A) : (B))
-#define CEIL_DIV(A, B)                                   \
-  (((B) == 1)                                            \
-       ? (A)                                             \
-       : ((B) == 2) ? ((int)((A) + 1) >> 1)              \
-                    : ((B) == 4) ? ((int)((A) + 3) >> 2) \
-                                 : (((A) + ((A) > 0 ? (B)-1 : 0)) / (B)))
+#define CEIL_DIV(A, B)                  \
+  (((B) == 1)   ? (A)                   \
+   : ((B) == 2) ? ((int)((A) + 1) >> 1) \
+   : ((B) == 4) ? ((int)((A) + 3) >> 2) \
+                : (((A) + ((A) > 0 ? (B) - 1 : 0)) / (B)))
 
 // This works only up to blocks of size 256 x 256 and for all N that are powers
 // of two.
@@ -333,22 +333,16 @@ static __global__ void filtered_lrelu_kernel(filtered_lrelu_kernel_params p) {
   const int szDownX = tileUpH * tileOutW;
 
   // Sizes for shared memory arrays.
-  const int s_buf0_size_base =
-      (filterMode == MODE_SUSD)
-          ? MAX(szIn, szUpXY)
-          : (filterMode == MODE_FUSD)
-                ? MAX(szIn, szDownX)
-                : (filterMode == MODE_SUFD)
-                      ? MAX(szIn, szUpXY)
-                      : (filterMode == MODE_FUFD) ? szIn : -1;
-  const int s_buf1_size_base =
-      (filterMode == MODE_SUSD)
-          ? MAX(szUpX, szDownX)
-          : (filterMode == MODE_FUSD)
-                ? szUpXY
-                : (filterMode == MODE_SUFD)
-                      ? szUpX
-                      : (filterMode == MODE_FUFD) ? szUpXY : -1;
+  const int s_buf0_size_base = (filterMode == MODE_SUSD)   ? MAX(szIn, szUpXY)
+                               : (filterMode == MODE_FUSD) ? MAX(szIn, szDownX)
+                               : (filterMode == MODE_SUFD) ? MAX(szIn, szUpXY)
+                               : (filterMode == MODE_FUFD) ? szIn
+                                                           : -1;
+  const int s_buf1_size_base = (filterMode == MODE_SUSD)   ? MAX(szUpX, szDownX)
+                               : (filterMode == MODE_FUSD) ? szUpXY
+                               : (filterMode == MODE_SUFD) ? szUpX
+                               : (filterMode == MODE_FUFD) ? szUpXY
+                                                           : -1;
 
   // Ensure U128 alignment.
   const int s_buf0_size = (s_buf0_size_base + 3) & ~3;
@@ -980,17 +974,17 @@ static __global__ void filtered_lrelu_kernel(filtered_lrelu_kernel_params p) {
 
 #define X_LOOP(TAPY, PX)                                             \
   for (int sx = 0; sx < fuSize / up; sx++) {                         \
-    v.x += a * (scalar_t)c_fu[(sx * up + (((PX)-0) & (up - 1))) +    \
+    v.x += a * (scalar_t)c_fu[(sx * up + (((PX) - 0) & (up - 1))) +  \
                               (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \
-    v.z += b * (scalar_t)c_fu[(sx * up + (((PX)-0) & (up - 1))) +    \
+    v.z += b * (scalar_t)c_fu[(sx * up + (((PX) - 0) & (up - 1))) +  \
                               (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \
     if ((PX) == 0) {                                                 \
       a = b;                                                         \
       b = s_tileIn[src0 + 2 + sx + sy * tileInW];                    \
     }                                                                \
-    v.y += a * (scalar_t)c_fu[(sx * up + (((PX)-1) & (up - 1))) +    \
+    v.y += a * (scalar_t)c_fu[(sx * up + (((PX) - 1) & (up - 1))) +  \
                               (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \
-    v.w += b * (scalar_t)c_fu[(sx * up + (((PX)-1) & (up - 1))) +    \
+    v.w += b * (scalar_t)c_fu[(sx * up + (((PX) - 1) & (up - 1))) +  \
                               (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \
     if ((PX) == 1) {                                                 \
       a = b;                                                         \
@@ -1447,7 +1441,7 @@ static __global__ void filtered_lrelu_act_kernel(
         s |= __shfl_xor(s, 4);
         s |= __shfl_xor(s, 8);
 #else
-        s |= __shfl_xor_sync(m, s, 1);                  // Distribute.
+        s |= __shfl_xor_sync(m, s, 1);  // Distribute.
         s |= __shfl_xor_sync(m, s, 2);
         s |= __shfl_xor_sync(m, s, 4);
         s |= __shfl_xor_sync(m, s, 8);

mmcv/ops/csrc/pytorch/cuda/upfirdn2d_kernel.cu

@@ -225,9 +225,9 @@ static __global__ void upfirdn2d_kernel_small(upfirdn2d_kernel_params p) {
         scalar_t v = 0;
         if (inX >= 0 & inY >= 0 & inX < p.inSize.x & inY < p.inSize.y &
             c < p.inSize.z)
-          v = (scalar_t)(
-              (const T *)p.x)[inX * p.inStride.x + inY * p.inStride.y +
-                              c * p.inStride.z + n * p.inStride.w];
+          v = (scalar_t)((const T *)
+                             p.x)[inX * p.inStride.x + inY * p.inStride.y +
+                                  c * p.inStride.z + n * p.inStride.w];
         sx[relInY][relInX][relC] = v;
       }
 

mmcv/ops/csrc/pytorch/npu/bbox_overlaps_npu.cpp

@@ -26,7 +26,7 @@ void bbox_overlaps_npu(const Tensor bboxes1, const Tensor bboxes2, Tensor ious,
     gtboxesFP32 = gtboxesFP32.to(at::kFloat);
   }
   c10::SmallVector<int64_t, 8> iousSize = {gtboxesFP32.size(0),
-                                              bboxesFP32.size(0)};
+                                           bboxesFP32.size(0)};
   if (aligned) {
     iousSize = {gtboxesFP32.size(0), 1};
   }

mmcv/ops/csrc/pytorch/npu/box_iou_quadri_npu.cpp

@@ -4,11 +4,10 @@ using namespace NPU_NAME_SPACE;
 using namespace std;
 
 void box_iou_quadri_impl(const Tensor boxes1, const Tensor boxes2, Tensor ious,
-                          const int mode_flag, const bool aligned);
+                         const int mode_flag, const bool aligned);
 
 void box_iou_quadri_npu(const Tensor boxes1, const Tensor boxes2, Tensor ious,
-                         const int mode_flag, const bool aligned) {
-
+                        const int mode_flag, const bool aligned) {
   TORCH_CHECK(boxes1.size(1) == 8, "boxes1 must be 2D tensor (N, 8)");
   TORCH_CHECK(boxes1.size(1) == 8, "boxes1 must be 2D tensor (N, 8)");
 

mmcv/ops/csrc/pytorch/npu/box_iou_rotated_npu.cpp

@@ -8,14 +8,13 @@ void box_iou_rotated_impl(const Tensor boxes1, const Tensor boxes2, Tensor ious,
 
 void box_iou_rotated_npu(const Tensor boxes1, const Tensor boxes2, Tensor ious,
                          const int mode_flag, const bool aligned) {
-
   TORCH_CHECK(boxes1.size(1) == 5, "boxes1 must be 2D tensor (N, 5)");
   TORCH_CHECK(boxes1.size(1) == 5, "boxes1 must be 2D tensor (N, 5)");
 
   auto trans = false;
   auto is_clockwise = false;
   EXEC_NPU_CMD(aclnnBoxesOverlapBev, boxes1, boxes2, trans, is_clockwise,
-                aligned, mode_flag, ious);
+               aligned, mode_flag, ious);
   return;
 }
 

mmcv/ops/csrc/pytorch/npu/boxes_overlap_bev_npu.cpp

@@ -10,16 +10,17 @@ void iou3d_boxes_overlap_bev_forward_impl(const int num_a, const Tensor boxes_a,
 void iou3d_boxes_overlap_bev_forward_npu(const int num_a, const Tensor boxes_a,
                                          const int num_b, const Tensor boxes_b,
                                          Tensor ans_overlap) {
+  TORCH_CHECK(boxes_a.size(1) == 7, "boxes_a must be 2D tensor (N, 7)");
+  TORCH_CHECK(boxes_b.size(1) == 7, "boxes_b must be 2D tensor (N, 7)");
 
-    TORCH_CHECK(boxes_a.size(1) == 7, "boxes_a must be 2D tensor (N, 7)");
-    TORCH_CHECK(boxes_b.size(1) == 7, "boxes_b must be 2D tensor (N, 7)");
-
-    auto trans = false;
-    auto is_clockwise = false;
-    auto aligned = false;
-    auto mode_flag = 2;
-    EXEC_NPU_CMD(aclnnBoxesOverlapBev, boxes_a, boxes_b, trans, is_clockwise, aligned, mode_flag, ans_overlap);
-    return;
+  auto trans = false;
+  auto is_clockwise = false;
+  auto aligned = false;
+  auto mode_flag = 2;
+  EXEC_NPU_CMD(aclnnBoxesOverlapBev, boxes_a, boxes_b, trans, is_clockwise,
+               aligned, mode_flag, ans_overlap);
+  return;
 }
 
-REGISTER_NPU_IMPL(iou3d_boxes_overlap_bev_forward_impl, iou3d_boxes_overlap_bev_forward_npu);
+REGISTER_NPU_IMPL(iou3d_boxes_overlap_bev_forward_impl,
+                  iou3d_boxes_overlap_bev_forward_npu);