This is an archive of the discontinued LLVM Phabricator instance.

Differential D41521

[CUDA] fixes for __shfl_* intrinsics.
ClosedPublic

Authored by tra on Dec 21 2017, 2:37 PM.

Details

Reviewers
jlebar
Commits
rG3cebc738b6fd: [CUDA] More fixes for __shfl_* intrinsics.
rL321326: [CUDA] More fixes for __shfl_* intrinsics.
rC321326: [CUDA] More fixes for __shfl_* intrinsics.
Summary
  • __shfl_{up,down}* uses unsigned int for the third parameter.
  • added [unsigned] long overloads for non-sync shuffles. Augments r319908 which added long overload for sync shuffles.

Diff Detail

Repository
rC Clang

Event Timeline

tra created this revision.Dec 21 2017, 2:37 PM
Herald added a subscriber: sanjoy. · View Herald TranscriptDec 21 2017, 2:37 PM
jlebar accepted this revision.Dec 21 2017, 3:13 PM

Since this is tricky and we've seen it affecting user code, do you think it's a bad idea to add tests to the test-suite?

This revision is now accepted and ready to land.Dec 21 2017, 3:13 PM
tra added a comment.Dec 21 2017, 3:29 PM

Added to my todo list. There are few more gaps that I want to test in order to make sure we don't regress on compatibility with older CUDA versions while changing these wrappers.

Closed by commit rC321326: [CUDA] More fixes for __shfl_* intrinsics. (authored by tra). · Explain WhyDec 21 2017, 3:53 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
lib/
Headers/
77 lines

Diff 127962

lib/Headers/__clang_cuda_intrinsics.h

Show All 28 Lines
// sm_30 intrinsics: __shfl_{up,down,xor}. // sm_30 intrinsics: __shfl_{up,down,xor}.
#define __SM_30_INTRINSICS_H__ #define __SM_30_INTRINSICS_H__
#define __SM_30_INTRINSICS_HPP__ #define __SM_30_INTRINSICS_HPP__
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300
#pragma push_macro("__MAKE_SHUFFLES") #pragma push_macro("__MAKE_SHUFFLES")
#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask) \ #define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask, \
inline __device__ int __FnName(int __val, int __offset, \ __Type) \
inline __device__ int __FnName(int __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return __IntIntrinsic(__val, __offset, \ return __IntIntrinsic(__val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((warpSize - __width) << 8) | (__Mask)); \
} \ } \
inline __device__ float __FnName(float __val, int __offset, \ inline __device__ float __FnName(float __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return __FloatIntrinsic(__val, __offset, \ return __FloatIntrinsic(__val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((warpSize - __width) << 8) | (__Mask)); \
} \ } \
inline __device__ unsigned int __FnName(unsigned int __val, int __offset, \ inline __device__ unsigned int __FnName(unsigned int __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return static_cast<unsigned int>( \ return static_cast<unsigned int>( \
::__FnName(static_cast<int>(__val), __offset, __width)); \ ::__FnName(static_cast<int>(__val), __offset, __width)); \
} \ } \
inline __device__ long long __FnName(long long __val, int __offset, \ inline __device__ long long __FnName(long long __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
struct __Bits { \ struct __Bits { \
int __a, __b; \ int __a, __b; \
}; \ }; \
_Static_assert(sizeof(__val) == sizeof(__Bits)); \ _Static_assert(sizeof(__val) == sizeof(__Bits)); \
_Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \ _Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \
__Bits __tmp; \ __Bits __tmp; \
memcpy(&__val, &__tmp, sizeof(__val)); \ memcpy(&__val, &__tmp, sizeof(__val)); \
__tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \ __tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \
__tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \ __tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \
long long __ret; \ long long __ret; \
memcpy(&__ret, &__tmp, sizeof(__tmp)); \ memcpy(&__ret, &__tmp, sizeof(__tmp)); \
return __ret; \ return __ret; \
} \ } \
inline __device__ long __FnName(long __val, __Type __offset, \
int __width = warpSize) { \
_Static_assert(sizeof(long) == sizeof(long long) || \
sizeof(long) == sizeof(int)); \
if (sizeof(long) == sizeof(long long)) { \
return static_cast<long>( \
::__FnName(static_cast<long long>(__val), __offset, __width)); \
} else if (sizeof(long) == sizeof(int)) { \
return static_cast<long>( \
::__FnName(static_cast<int>(__val), __offset, __width)); \
} \
} \
inline __device__ unsigned long __FnName( \
unsigned long __val, __Type __offset, int __width = warpSize) { \
return static_cast<unsigned long>( \
::__FnName(static_cast<long>(__val), __offset, __width)); \
} \
inline __device__ unsigned long long __FnName( \ inline __device__ unsigned long long __FnName( \
unsigned long long __val, int __offset, int __width = warpSize) { \ unsigned long long __val, __Type __offset, int __width = warpSize) { \
return static_cast<unsigned long long>(::__FnName( \ return static_cast<unsigned long long>(::__FnName( \
static_cast<unsigned long long>(__val), __offset, __width)); \ static_cast<unsigned long long>(__val), __offset, __width)); \
} \ } \
inline __device__ double __FnName(double __val, int __offset, \ inline __device__ double __FnName(double __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
long long __tmp; \ long long __tmp; \
_Static_assert(sizeof(__tmp) == sizeof(__val)); \ _Static_assert(sizeof(__tmp) == sizeof(__val)); \
memcpy(&__tmp, &__val, sizeof(__val)); \ memcpy(&__tmp, &__val, sizeof(__val)); \
__tmp = ::__FnName(__tmp, __offset, __width); \ __tmp = ::__FnName(__tmp, __offset, __width); \
double __ret; \ double __ret; \
memcpy(&__ret, &__tmp, sizeof(__ret)); \ memcpy(&__ret, &__tmp, sizeof(__ret)); \
return __ret; \ return __ret; \
} }
__MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f); __MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f, int);
// We use 0 rather than 31 as our mask, because shfl.up applies to lanes >= // We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=
// maxLane. // maxLane.
__MAKE_SHUFFLES(__shfl_up, __nvvm_shfl_up_i32, __nvvm_shfl_up_f32, 0); __MAKE_SHUFFLES(__shfl_up, __nvvm_shfl_up_i32, __nvvm_shfl_up_f32, 0,
__MAKE_SHUFFLES(__shfl_down, __nvvm_shfl_down_i32, __nvvm_shfl_down_f32, 0x1f); unsigned int);
__MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f); __MAKE_SHUFFLES(__shfl_down, __nvvm_shfl_down_i32, __nvvm_shfl_down_f32, 0x1f,
unsigned int);
__MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f,
int);
#pragma pop_macro("__MAKE_SHUFFLES") #pragma pop_macro("__MAKE_SHUFFLES")
#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300 #endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300
#if CUDA_VERSION >= 9000 #if CUDA_VERSION >= 9000
#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300) #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300)
// __shfl_sync_* variants available in CUDA-9 // __shfl_sync_* variants available in CUDA-9
#pragma push_macro("__MAKE_SYNC_SHUFFLES") #pragma push_macro("__MAKE_SYNC_SHUFFLES")
#define __MAKE_SYNC_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, \ #define __MAKE_SYNC_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, \
__Mask) \ __Mask, __Type) \
inline __device__ int __FnName(unsigned int __mask, int __val, int __offset, \ inline __device__ int __FnName(unsigned int __mask, int __val, \
int __width = warpSize) { \ __Type __offset, int __width = warpSize) { \
return __IntIntrinsic(__mask, __val, __offset, \ return __IntIntrinsic(__mask, __val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((warpSize - __width) << 8) | (__Mask)); \
} \ } \
inline __device__ float __FnName(unsigned int __mask, float __val, \ inline __device__ float __FnName(unsigned int __mask, float __val, \
int __offset, int __width = warpSize) { \ __Type __offset, int __width = warpSize) { \
return __FloatIntrinsic(__mask, __val, __offset, \ return __FloatIntrinsic(__mask, __val, __offset, \
((warpSize - __width) << 8) | (__Mask)); \ ((warpSize - __width) << 8) | (__Mask)); \
} \ } \
inline __device__ unsigned int __FnName(unsigned int __mask, \ inline __device__ unsigned int __FnName(unsigned int __mask, \
unsigned int __val, int __offset, \ unsigned int __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return static_cast<unsigned int>( \ return static_cast<unsigned int>( \
::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \ ::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \
} \ } \
inline __device__ long long __FnName(unsigned int __mask, long long __val, \ inline __device__ long long __FnName(unsigned int __mask, long long __val, \
int __offset, int __width = warpSize) { \ __Type __offset, \
int __width = warpSize) { \
struct __Bits { \ struct __Bits { \
int __a, __b; \ int __a, __b; \
}; \ }; \
_Static_assert(sizeof(__val) == sizeof(__Bits)); \ _Static_assert(sizeof(__val) == sizeof(__Bits)); \
_Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \ _Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \
__Bits __tmp; \ __Bits __tmp; \
memcpy(&__val, &__tmp, sizeof(__val)); \ memcpy(&__val, &__tmp, sizeof(__val)); \
__tmp.__a = ::__FnName(__mask, __tmp.__a, __offset, __width); \ __tmp.__a = ::__FnName(__mask, __tmp.__a, __offset, __width); \
__tmp.__b = ::__FnName(__mask, __tmp.__b, __offset, __width); \ __tmp.__b = ::__FnName(__mask, __tmp.__b, __offset, __width); \
long long __ret; \ long long __ret; \
memcpy(&__ret, &__tmp, sizeof(__tmp)); \ memcpy(&__ret, &__tmp, sizeof(__tmp)); \
return __ret; \ return __ret; \
} \ } \
inline __device__ unsigned long long __FnName( \ inline __device__ unsigned long long __FnName( \
unsigned int __mask, unsigned long long __val, int __offset, \ unsigned int __mask, unsigned long long __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return static_cast<unsigned long long>(::__FnName( \ return static_cast<unsigned long long>(::__FnName( \
__mask, static_cast<unsigned long long>(__val), __offset, __width)); \ __mask, static_cast<unsigned long long>(__val), __offset, __width)); \
} \ } \
inline __device__ long __FnName(unsigned int __mask, long __val, \ inline __device__ long __FnName(unsigned int __mask, long __val, \
int __offset, int __width = warpSize) { \ __Type __offset, int __width = warpSize) { \
_Static_assert(sizeof(long) == sizeof(long long) || \ _Static_assert(sizeof(long) == sizeof(long long) || \
sizeof(long) == sizeof(int)); \ sizeof(long) == sizeof(int)); \
if (sizeof(long) == sizeof(long long)) { \ if (sizeof(long) == sizeof(long long)) { \
return static_cast<long>(::__FnName( \ return static_cast<long>(::__FnName( \
__mask, static_cast<long long>(__val), __offset, __width)); \ __mask, static_cast<long long>(__val), __offset, __width)); \
} else if (sizeof(long) == sizeof(int)) { \ } else if (sizeof(long) == sizeof(int)) { \
return static_cast<long>( \ return static_cast<long>( \
::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \ ::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \
} \ } \
} \ } \
inline __device__ unsigned long __FnName(unsigned int __mask, \ inline __device__ unsigned long __FnName( \
unsigned long __val, int __offset, \ unsigned int __mask, unsigned long __val, __Type __offset, \
int __width = warpSize) { \ int __width = warpSize) { \
return static_cast<unsigned long>( \ return static_cast<unsigned long>( \
::__FnName(__mask, static_cast<long>(__val), __offset, __width)); \ ::__FnName(__mask, static_cast<long>(__val), __offset, __width)); \
} \ } \
inline __device__ double __FnName(unsigned int __mask, double __val, \ inline __device__ double __FnName(unsigned int __mask, double __val, \
int __offset, int __width = warpSize) { \ __Type __offset, int __width = warpSize) { \
long long __tmp; \ long long __tmp; \
_Static_assert(sizeof(__tmp) == sizeof(__val)); \ _Static_assert(sizeof(__tmp) == sizeof(__val)); \
memcpy(&__tmp, &__val, sizeof(__val)); \ memcpy(&__tmp, &__val, sizeof(__val)); \
__tmp = ::__FnName(__mask, __tmp, __offset, __width); \ __tmp = ::__FnName(__mask, __tmp, __offset, __width); \
double __ret; \ double __ret; \
memcpy(&__ret, &__tmp, sizeof(__ret)); \ memcpy(&__ret, &__tmp, sizeof(__ret)); \
return __ret; \ return __ret; \
} }
__MAKE_SYNC_SHUFFLES(__shfl_sync, __nvvm_shfl_sync_idx_i32, __MAKE_SYNC_SHUFFLES(__shfl_sync, __nvvm_shfl_sync_idx_i32,
__nvvm_shfl_sync_idx_f32, 0x1f); __nvvm_shfl_sync_idx_f32, 0x1f, int);
// We use 0 rather than 31 as our mask, because shfl.up applies to lanes >= // We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=
// maxLane. // maxLane.
__MAKE_SYNC_SHUFFLES(__shfl_up_sync, __nvvm_shfl_sync_up_i32, __MAKE_SYNC_SHUFFLES(__shfl_up_sync, __nvvm_shfl_sync_up_i32,
__nvvm_shfl_sync_up_f32, 0); __nvvm_shfl_sync_up_f32, 0, unsigned int);
__MAKE_SYNC_SHUFFLES(__shfl_down_sync, __nvvm_shfl_sync_down_i32, __MAKE_SYNC_SHUFFLES(__shfl_down_sync, __nvvm_shfl_sync_down_i32,
__nvvm_shfl_sync_down_f32, 0x1f); __nvvm_shfl_sync_down_f32, 0x1f, unsigned int);
__MAKE_SYNC_SHUFFLES(__shfl_xor_sync, __nvvm_shfl_sync_bfly_i32, __MAKE_SYNC_SHUFFLES(__shfl_xor_sync, __nvvm_shfl_sync_bfly_i32,
__nvvm_shfl_sync_bfly_f32, 0x1f); __nvvm_shfl_sync_bfly_f32, 0x1f, int);
#pragma pop_macro("__MAKE_SYNC_SHUFFLES") #pragma pop_macro("__MAKE_SYNC_SHUFFLES")
inline __device__ void __syncwarp(unsigned int mask = 0xffffffff) { inline __device__ void __syncwarp(unsigned int mask = 0xffffffff) {
return __nvvm_bar_warp_sync(mask); return __nvvm_bar_warp_sync(mask);
} }
inline __device__ void __barrier_sync(unsigned int id) { inline __device__ void __barrier_sync(unsigned int id) {
__nvvm_barrier_sync(id); __nvvm_barrier_sync(id);
▲ Show 20 Lines • Show All 285 Lines • Show Last 20 Lines