diff --git a/mmcv/ops/csrc/common/cuda/scatter_points_cuda_kernel.cuh b/mmcv/ops/csrc/common/cuda/scatter_points_cuda_kernel.cuh
index ffac86244..1bafa9361 100644
--- a/mmcv/ops/csrc/common/cuda/scatter_points_cuda_kernel.cuh
+++ b/mmcv/ops/csrc/common/cuda/scatter_points_cuda_kernel.cuh
@@ -37,8 +37,13 @@ __device__ __forceinline__ static void reduceMax(double *address, double val) {
 #ifdef __CUDA_ARCH__
 __device__ __forceinline__ static void reduceAdd(float *address, float val) {
 #if (__CUDA_ARCH__ < 200)
+#ifdef _MSC_VER
+#pragma message( \
+    "compute capability lower than 2.x. fall back to use CAS version of atomicAdd for float32")
+#else
 #warning \
     "compute capability lower than 2.x. fall back to use CAS version of atomicAdd for float32"
+#endif
   int *address_as_i = reinterpret_cast<int *>(address);
   int old = *address_as_i, assumed;
   do {
@@ -53,8 +58,13 @@ __device__ __forceinline__ static void reduceAdd(float *address, float val) {
 
 __device__ __forceinline__ static void reduceAdd(double *address, double val) {
 #if (__CUDA_ARCH__ < 600)
+#ifdef _MSC_VER
+#pragma message( \
+    "compute capability lower than 6.x. fall back to use CAS version of atomicAdd for float64")
+#else
 #warning \
     "compute capability lower than 6.x. fall back to use CAS version of atomicAdd for float64"
+#endif
   unsigned long long *address_as_ull =
       reinterpret_cast<unsigned long long *>(address);
   unsigned long long old = *address_as_ull, assumed;