Skip to content

Commit

Permalink
interpd.cu : added compilation TODO for Nvidia architectures < CC 6.0
Browse files Browse the repository at this point in the history
  • Loading branch information
thorstone25 committed Jul 12, 2023
1 parent 9521992 commit 60776b5
Showing 1 changed file with 9 additions and 0 deletions.
9 changes: 9 additions & 0 deletions src/interpd.cu
Original file line number Diff line number Diff line change
Expand Up @@ -190,10 +190,19 @@ __device__ void interpd_temp(T2 * __restrict__ y,
}
}

// atomicAdd supported @ CC 6.x+
#if (__CUDA_ARCH__ >= 600)
// nothing to do
#else
// TODO: implement atomicAdd for half/double via atomicCAS
#endif

// half type supported @ CC 5.3+
#if (__CUDA_ARCH__ >= 530)
inline __device__ void atomicAddStore(half2 * y, const half2 val){
atomicAdd(y, val);
}
#endif
inline __device__ void atomicAddStore(float2 * y, const float2 val){
atomicAdd(&y[0].x, val.x);
atomicAdd(&y[0].y, val.y);
Expand Down

0 comments on commit 60776b5

Please sign in to comment.