Function shfl_up_sync

src/backend/cuda/kernel/shfl_intrinsics.hpp:81–87

//__shfl_up_sync wrapper
template<typename T>
__device__ T shfl_up_sync(T var, int delta) {
#if (CUDA_VERSION >= 9000)
    return __shfl_up_sync(FULL_MASK, var, delta);
#else
    return __shfl_up(var, delta);
#endif
}
// specialization for cfloat
template<>
inline __device__ cfloat shfl_up_sync(cfloat var, int delta) {