| 46 | //__shfl_down_sync wrapper |
| 47 | template<typename T> |
| 48 | __device__ T shfl_down_sync(T var, int delta) { |
| 49 | #if (CUDA_VERSION >= 9000) |
| 50 | return __shfl_down_sync(FULL_MASK, var, delta); |
| 51 | #else |
| 52 | return __shfl_down(var, delta); |
| 53 | #endif |
| 54 | } |
| 55 | // specialization for cfloat |
| 56 | template<> |
| 57 | inline __device__ cfloat shfl_down_sync(cfloat var, int delta) { |
no outgoing calls