Skip to content

SWDEV-573004 - fix shfl_sync for compiler init value #2533

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
shadidashmiz wants to merge 1 commit into
base: develop
Choose a base branch
from
+14 −4
Open

SWDEV-573004 - fix shfl_sync for compiler init value #2533

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 14 additions & 4 deletions projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,13 @@ THE SOFTWARE.
#include <algorithm>
#endif

#pragma push_macro("MAYBE_UNDEF")
#if defined(__has_attribute) && __has_attribute(maybe_undef)
#define MAYBE_UNDEF __attribute__((maybe_undef))
#else
#define MAYBE_UNDEF
#endif

extern "C" __device__ __attribute__((const)) int __ockl_wfred_add_i32(int);
extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_add_u32(unsigned int);
extern "C" __device__ __attribute__((const)) int __ockl_wfred_min_i32(int);
Expand Down Expand Up @@ -273,7 +280,7 @@ __device__ inline unsigned long long __match_all_sync(MaskT mask, T value, int*
// various variants of shfl

template <typename MaskT, typename T>
__device__ inline T __shfl_sync(MaskT mask, T var, int srcLane, int width = warpSize) {
__device__ inline T __shfl_sync(MaskT mask, MAYBE_UNDEF T var, int srcLane, int width = warpSize) {
static_assert(__hip_internal::is_integral<MaskT>::value && sizeof(MaskT) == 8,
"The mask must be a 64-bit integer. "
"Implicitly promoting a smaller integer is almost always an error.");
Expand All @@ -283,7 +290,7 @@ __device__ inline T __shfl_sync(MaskT mask, T var, int srcLane, int width = warp
}

template <typename MaskT, typename T>
__device__ inline T __shfl_up_sync(MaskT mask, T var, unsigned int delta, int width = warpSize) {
__device__ inline T __shfl_up_sync(MaskT mask, MAYBE_UNDEF T var, unsigned int delta, int width = warpSize) {
static_assert(__hip_internal::is_integral<MaskT>::value && sizeof(MaskT) == 8,
"The mask must be a 64-bit integer. "
"Implicitly promoting a smaller integer is almost always an error.");
Expand All @@ -293,7 +300,7 @@ __device__ inline T __shfl_up_sync(MaskT mask, T var, unsigned int delta, int wi
}

template <typename MaskT, typename T>
__device__ inline T __shfl_down_sync(MaskT mask, T var, unsigned int delta, int width = warpSize) {
__device__ inline T __shfl_down_sync(MaskT mask, MAYBE_UNDEF T var, unsigned int delta, int width = warpSize) {
static_assert(__hip_internal::is_integral<MaskT>::value && sizeof(MaskT) == 8,
"The mask must be a 64-bit integer. "
"Implicitly promoting a smaller integer is almost always an error.");
Expand All @@ -303,7 +310,7 @@ __device__ inline T __shfl_down_sync(MaskT mask, T var, unsigned int delta, int
}

template <typename MaskT, typename T>
__device__ inline T __shfl_xor_sync(MaskT mask, T var, int laneMask, int width = warpSize) {
__device__ inline T __shfl_xor_sync(MaskT mask, MAYBE_UNDEF T var, int laneMask, int width = warpSize) {
static_assert(__hip_internal::is_integral<MaskT>::value && sizeof(MaskT) == 8,
"The mask must be a 64-bit integer. "
"Implicitly promoting a smaller integer is almost always an error.");
Expand Down Expand Up @@ -679,4 +686,7 @@ __device__ inline unsigned long long __reduce_xor_sync(MaskT mask, unsigned long
#undef __hip_adjust_mask_for_wave32

#endif // HIP_ENABLE_EXTRA_WARP_SYNC_TYPES

#pragma pop_macro("MAYBE_UNDEF")

#endif // HIP_DISABLE_WARP_SYNC_BUILTINS