mirrors/mmcv
1
0
Fork 0
mirror of https://github.com/open-mmlab/mmcv.git synced 2025-06-03 21:54:52 +08:00
Code Issues Projects Releases Wiki Activity

[Fix] Fix compiled error on windows (#1510)

Browse Source
This commit is contained in:
Zaida Zhou 2021-11-23 18:41:25 +08:00 committed by GitHub
parent 71341e9a0a
commit 9cad97bcf8
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 10 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
mmcv/ops/csrc/common/cuda/scatter_points_cuda_kernel.cuh
View File

@ -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;