diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h index 0bed52578f..1042d459ba 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h @@ -685,7 +685,7 @@ __BF16_HOST_DEVICE_STATIC__ __hip_bfloat16 __ushort_as_bfloat16(const unsigned s __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_down_sync(const unsigned long long mask, const __hip_bfloat16 in, const unsigned int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { return __ushort_as_bfloat16(__shfl_down_sync(mask, __bfloat16_as_ushort(in), delta, width)); } @@ -696,7 +696,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_down_sync(const unsigned long long __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_down_sync(const unsigned long long mask, const __hip_bfloat162 in, const unsigned int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { static_assert(sizeof(__hip_bfloat162) == sizeof(unsigned int)); union { __hip_bfloat162 bf162; @@ -712,7 +712,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_down_sync(const unsigned long long */ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_sync(const unsigned long long mask, const __hip_bfloat16 in, const int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { return __ushort_as_bfloat16(__shfl_sync(mask, __bfloat16_as_ushort(in), delta, width)); } @@ -722,7 +722,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_sync(const unsigned long long mask, */ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_sync(const unsigned long long mask, const __hip_bfloat162 in, const int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { static_assert(sizeof(__hip_bfloat162) == sizeof(unsigned int)); union { __hip_bfloat162 bf162; @@ -739,7 +739,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_sync(const unsigned long long mask __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_up_sync(const unsigned long long mask, const __hip_bfloat16 in, const unsigned int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { return __ushort_as_bfloat16(__shfl_up_sync(mask, __bfloat16_as_ushort(in), delta, width)); } @@ -750,7 +750,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_up_sync(const unsigned long long ma __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_up_sync(const unsigned long long mask, const __hip_bfloat162 in, const unsigned int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { static_assert(sizeof(__hip_bfloat162) == sizeof(unsigned int)); union { __hip_bfloat162 bf162; @@ -766,7 +766,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_up_sync(const unsigned long long m */ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_xor_sync(const unsigned long long mask, const __hip_bfloat16 in, const int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { return __ushort_as_bfloat16(__shfl_xor_sync(mask, __bfloat16_as_ushort(in), delta, width)); } @@ -776,7 +776,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 __shfl_xor_sync(const unsigned long long m */ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_xor_sync(const unsigned long long mask, const __hip_bfloat162 in, const int delta, - const int width = __AMDGCN_WAVEFRONT_SIZE) { + const int width = warpSize) { static_assert(sizeof(__hip_bfloat162) == sizeof(unsigned int)); union { __hip_bfloat162 bf162; diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 1901b49aa4..3a66c3c00c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -242,7 +242,7 @@ class thread_block : public thread_group { __CG_QUALIFIER__ thread_group new_tiled_group(unsigned int tile_size) const { const bool pow2 = ((tile_size & (tile_size - 1)) == 0); // Invalid tile size, assert - if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) { + if (!tile_size || (tile_size > warpSize) || !pow2) { __hip_assert(false && "invalid tile size"); } @@ -308,7 +308,7 @@ class tiled_group : public thread_group { __CG_QUALIFIER__ tiled_group new_tiled_group(unsigned int tile_size) const { const bool pow2 = ((tile_size & (tile_size - 1)) == 0); - if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) { + if (!tile_size || (tile_size > warpSize) || !pow2) { __hip_assert(false && "invalid tile size"); } @@ -374,7 +374,7 @@ class coalesced_group : public thread_group { if (coalesced_info.tiled_info.is_tiled) { unsigned int base_offset = (thread_rank() & (~(tile_size - 1))); unsigned int masklength = min(static_cast(size()) - base_offset, tile_size); - lane_mask member_mask = static_cast(-1) >> (__AMDGCN_WAVEFRONT_SIZE - masklength); + lane_mask member_mask = static_cast(-1) >> (warpSize - masklength); member_mask <<= (__lane_id() & ~(tile_size - 1)); coalesced_group coalesced_tile = coalesced_group(member_mask); @@ -389,7 +389,7 @@ class coalesced_group : public thread_group { unsigned int tile_rank = 0; int lanes_to_skip = ((thread_rank()) / tile_size) * tile_size; - for (unsigned int i = 0; i < __AMDGCN_WAVEFRONT_SIZE; i++) { + for (unsigned int i = 0; i < warpSize; i++) { lane_mask active = coalesced_info.member_mask & (1 << i); // Make sure the lane is active if (active) { @@ -466,11 +466,11 @@ class coalesced_group : public thread_group { srcRank = srcRank % static_cast(size()); - int lane = (size() == __AMDGCN_WAVEFRONT_SIZE) ? srcRank - : (__AMDGCN_WAVEFRONT_SIZE == 64) ? __fns64(coalesced_info.member_mask, 0, (srcRank + 1)) - : __fns32(coalesced_info.member_mask, 0, (srcRank + 1)); + int lane = (size() == warpSize) ? srcRank + : (warpSize == 64) ? __fns64(coalesced_info.member_mask, 0, (srcRank + 1)) + : __fns32(coalesced_info.member_mask, 0, (srcRank + 1)); - return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); + return __shfl(var, lane, warpSize); } /** \brief Shuffle down operation on group level. @@ -495,12 +495,12 @@ class coalesced_group : public thread_group { // and WARP_SIZE as the shift value rather than lane_delta itself. // This is not described in the documentation and is not done here. - if (size() == __AMDGCN_WAVEFRONT_SIZE) { - return __shfl_down(var, lane_delta, __AMDGCN_WAVEFRONT_SIZE); + if (size() == warpSize) { + return __shfl_down(var, lane_delta, warpSize); } int lane; - if (__AMDGCN_WAVEFRONT_SIZE == 64) { + if (warpSize == 64) { lane = __fns64(coalesced_info.member_mask, __lane_id(), lane_delta + 1); } else { @@ -511,7 +511,7 @@ class coalesced_group : public thread_group { lane = __lane_id(); } - return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); + return __shfl(var, lane, warpSize); } /** \brief Shuffle up operation on group level. @@ -536,15 +536,15 @@ class coalesced_group : public thread_group { // and WARP_SIZE as the shift value rather than lane_delta itself. // This is not described in the documentation and is not done here. - if (size() == __AMDGCN_WAVEFRONT_SIZE) { - return __shfl_up(var, lane_delta, __AMDGCN_WAVEFRONT_SIZE); + if (size() == warpSize) { + return __shfl_up(var, lane_delta, warpSize); } int lane; - if (__AMDGCN_WAVEFRONT_SIZE == 64) { + if (warpSize == 64) { lane = __fns64(coalesced_info.member_mask, __lane_id(), -(lane_delta + 1)); } - else if (__AMDGCN_WAVEFRONT_SIZE == 32) { + else if (warpSize == 32) { lane = __fns32(coalesced_info.member_mask, __lane_id(), -(lane_delta + 1)); } @@ -552,7 +552,7 @@ class coalesced_group : public thread_group { lane = __lane_id(); } - return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); + return __shfl(var, lane, warpSize); } #ifdef HIP_ENABLE_WARP_SYNC_BUILTINS diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h index a42440daa4..463fc06c4b 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -83,8 +83,18 @@ __device__ static inline int __hip_move_dpp_N(int src) { bound_ctrl); } -__device__ -static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; +#if defined(__SPIRV__) + inline __device__ const struct final { + __device__ + __attribute__((always_inline, const)) + operator int() const noexcept { + return __builtin_amdgcn_wavefrontsize(); + } + } warpSize{}; +#else + __device__ + static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; +#endif // warp vote function __all __any __ballot __device__ @@ -123,6 +133,7 @@ unsigned long long __activemask() { #endif // HIP_ENABLE_WARP_SYNC_BUILTINS __device__ static inline unsigned int __lane_id() { + if (warpSize == 32) return __builtin_amdgcn_mbcnt_lo(-1, 0); return __builtin_amdgcn_mbcnt_hi( -1, __builtin_amdgcn_mbcnt_lo(-1, 0)); } diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h index a84dada0ad..bd78b47aa6 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h @@ -230,7 +230,7 @@ unsigned long long __match_all_sync(MaskT mask, T value, int* pred) { template __device__ inline T __shfl_sync(MaskT mask, T var, int srcLane, - int width = __AMDGCN_WAVEFRONT_SIZE) { + int width = warpSize) { static_assert( __hip_internal::is_integral::value && sizeof(MaskT) == 8, "The mask must be a 64-bit integer. " @@ -243,7 +243,7 @@ T __shfl_sync(MaskT mask, T var, int srcLane, template __device__ inline T __shfl_up_sync(MaskT mask, T var, unsigned int delta, - int width = __AMDGCN_WAVEFRONT_SIZE) { + int width = warpSize) { static_assert( __hip_internal::is_integral::value && sizeof(MaskT) == 8, "The mask must be a 64-bit integer. " @@ -256,7 +256,7 @@ T __shfl_up_sync(MaskT mask, T var, unsigned int delta, template __device__ inline T __shfl_down_sync(MaskT mask, T var, unsigned int delta, - int width = __AMDGCN_WAVEFRONT_SIZE) { + int width = warpSize) { static_assert( __hip_internal::is_integral::value && sizeof(MaskT) == 8, "The mask must be a 64-bit integer. " @@ -269,7 +269,7 @@ T __shfl_down_sync(MaskT mask, T var, unsigned int delta, template __device__ inline T __shfl_xor_sync(MaskT mask, T var, int laneMask, - int width = __AMDGCN_WAVEFRONT_SIZE) { + int width = warpSize) { static_assert( __hip_internal::is_integral::value && sizeof(MaskT) == 8, "The mask must be a 64-bit integer. " diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index 95379bcaef..995c8e5949 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -52,6 +52,10 @@ THE SOFTWARE. #define _CG_STATIC_CONST_DECL_ static constexpr #endif +#if defined(__SPIRV__) && !defined(__AMDGCN_WAVEFRONT_SIZE) +#error "TEMPORARY LIMITATION: when targeting AMDGCN SPIR-V" + "__AMDGCN_WAVEFRONT_SIZE is not defined, and must be defined by the user" +#endif #if __AMDGCN_WAVEFRONT_SIZE == 32 using lane_mask = unsigned int; #else @@ -94,7 +98,7 @@ typedef enum { * @ingroup CooperativeG * @{ * This section describes the cooperative groups functions of HIP runtime API. - * + * * The cooperative groups provides flexible thread parallel programming algorithms, threads * cooperate and share data to perform collective computations. * @@ -118,7 +122,7 @@ namespace helper { __CG_STATIC_QUALIFIER__ unsigned long long adjust_mask( unsigned long long base_mask, unsigned long long input_mask) { unsigned long long out = 0; - for (unsigned int i = 0, index = 0; i < __AMDGCN_WAVEFRONT_SIZE; i++) { + for (unsigned int i = 0, index = 0; i < warpSize; i++) { auto lane_active = base_mask & (1ull << i); if (lane_active) { auto result = input_mask & (1ull << i); @@ -245,14 +249,14 @@ __CG_STATIC_QUALIFIER__ void sync() { __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, " // have i-th bit of x set and come before the current thread. __CG_STATIC_QUALIFIER__ unsigned int masked_bit_count(lane_mask x, unsigned int add = 0) { unsigned int counter=0; - #if __AMDGCN_WAVEFRONT_SIZE == 32 + if (warpSize == 32) { counter = __builtin_amdgcn_mbcnt_lo(x, add); - #else - counter = __builtin_amdgcn_mbcnt_lo(static_cast(x), add); - counter = __builtin_amdgcn_mbcnt_hi(static_cast(x >> 32), counter); - #endif + } else { + counter = __builtin_amdgcn_mbcnt_lo(static_cast(x), add); + counter = __builtin_amdgcn_mbcnt_hi(static_cast(x >> 32), counter); + } - return counter; + return counter; } } // namespace coalesced_group