From 9cad97bcf8f94051330ac62ef90eb0ed7a470bcd Mon Sep 17 00:00:00 2001
From: Zaida Zhou <58739961+zhouzaida@users.noreply.github.com>
Date: Tue, 23 Nov 2021 18:41:25 +0800
Subject: [PATCH] [Fix] Fix compiled error on windows (#1510)

---
 .../csrc/common/cuda/scatter_points_cuda_kernel.cuh    | 10 ++++++++++
 1 file changed, 10 insertions(+)

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;