| 524 | // Helper function for atomic accumulation implemented as CAS. |
| 525 | template <typename T, typename F> |
| 526 | __device__ T GpuAtomicCasHelper(T* ptr, F accumulate) { |
| 527 | T old = *ptr; |
| 528 | T assumed; |
| 529 | do { |
| 530 | assumed = old; |
| 531 | old = atomicCAS(ptr, assumed, accumulate(assumed)); |
| 532 | } while (assumed != old); |
| 533 | return old; |
| 534 | } |
| 535 | CREATE_CUDA_DEVICE_FUNCTION_ALIAS(GpuAtomicCasHelper, CudaAtomicCasHelper); |
| 536 | |
| 537 | // Overload for int64 |
no outgoing calls
no test coverage detected