Spaces:
Runtime error
Runtime error
| /* | |
| * Functions to share code between CPU and GPU | |
| */ | |
| // CUDA versions | |
| // Recent compute capabilities have block-level atomicAdd for all data types, so we use that | |
| // Older architectures don't have block-level atomicAdd, nor atomicAdd for doubles, so we defer to atomicAdd for float | |
| // and use the known atomicCAS-based implementation for double | |
| template<typename data_t> | |
| __device__ inline data_t atomic_add(data_t *address, data_t val) { | |
| return atomicAdd(address, val); | |
| } | |
| template<> | |
| __device__ inline double atomic_add(double *address, double val) { | |
| unsigned long long int* address_as_ull = (unsigned long long int*)address; | |
| unsigned long long int old = *address_as_ull, assumed; | |
| do { | |
| assumed = old; | |
| old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); | |
| } while (assumed != old); | |
| return __longlong_as_double(old); | |
| } | |
| // CPU versions | |