Shuenhoy/vtk-m2
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge topic 'cuda_array_handles_on_cuda8_reprise'

Browse Source 
f6da09214 Use CUDA_ARCH instead of CUDACC to guard device-only code.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1361
This commit is contained in:
Allison Vacanti 2018-08-09 16:48:06 +00:00 committed by Kitware Robot
commit a891e6d90a
1 changed files with 4 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
vtkm/cont/cuda/internal/DeviceAdapterAtomicArrayImplementationCuda.h
View File

@ -59,11 +59,9 @@ public:
// __host__ __device__, and nvcc 8.0.61 errors when calling the __device__
// function vtkmAtomicAdd. VTKM_SUPPRESS_EXEC_WARNINGS does not fix this.
// We work around this by calling the __device__ function inside of a
// __CUDACC__ guard, as nvcc is smart enough to recognize that this is a
// __CUDA_ARCH__ guard, as nvcc is smart enough to recognize that this is a
// safe usage of a __device__ function in a __host__ __device__ context.
// However, it is not smart enough to recognize that VTKM_CUDA is equivalent
// to __CUDACC__, so we must use the lower-level define here.
#ifdef __CUDACC__
#ifdef __CUDA_ARCH__
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
return this->vtkmAtomicAdd(lockedValue, value);
#else
@ -81,11 +79,9 @@ public:
// __host__ __device__, and nvcc 8.0.61 errors when calling the __device__
// function vtkmAtomicAdd. VTKM_SUPPRESS_EXEC_WARNINGS does not fix this.
// We work around this by calling the __device__ function inside of a
// __CUDACC__ guard, as nvcc is smart enough to recognize that this is a
// __CUDA_ARCH__ guard, as nvcc is smart enough to recognize that this is a
// safe usage of a __device__ function in a __host__ __device__ context.
// However, it is not smart enough to recognize that VTKM_CUDA is equivalent
// to __CUDACC__, so we must use the lower-level define here.
#ifdef __CUDACC__
#ifdef __CUDA_ARCH__
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
return this->vtkmCompareAndSwap(lockedValue, newValue, oldValue);
#else