AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

SWDEV-341212 - HIP header changes for supporting SPIR-V

فهرست منبع 
This removes almost all uses of the deprecated
__AMDGCN_WAVEFRONT_SIZE macro, which is unavailable
when targeting SPIR-V, and adds a SPIR-V compatible
formulation of warpSize (which should end up as the
sole definition of warpSize once we remove support
for treating it as a compile time constant). It
is incomplete in that the cooperative_groups
implementation will need additional surgery.

Squashed commit of the following:

commit 6840826c3fec8516857dc4f2092d84358550f588
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Fri Dec 6 23:36:32 2024 +0000

    Add deprecation warning for constexpr uses of `warpSize`.

commit a72307a7353034c2de53fd164e016967945fd0d1
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Fri Dec 6 23:12:14 2024 +0000

    Prepare HIP RT for SPIR-V.

commit 5e40dd746ac4f8c93b521ef048ff9d494905ba95
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Fri Dec 6 22:46:05 2024 +0000

    Revert stale change.

commit 231fe91c53dba4cabd832fc84eaa6ddb402271a0
Merge: a48905ec9 12dc02b4f
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Fri Dec 6 22:37:24 2024 +0000

    Merge branch 'amd-staging' of https://github.com/ROCm/clr into amd-staging

commit a48905ec9cfe0e017cc64943195be82b530117d7
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Tue Sep 17 03:14:56 2024 +0100

    Add scaffolding for SPIR-V support.

Change-Id: I2e84bbe90df58a5f9a8709b619905f04fa5b96dc


[ROCm/clr commit: dd4378611a]
This commit is contained in:
Konstantin Zhuravlyov
2024-12-10 08:08:59 -05:00
کامیت شده توسط Konstantin Zhuravlyov
والد 0795f00a14
کامیت 20b9b5a08c
5فایلهای تغییر یافته به همراه54 افزوده شده و 39 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_hip_bf16.h
+8 -8
مشاهده پرونده
@@ -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;
projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
+17 -17
مشاهده پرونده
@@ -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<unsigned int>(size()) - base_offset, tile_size);
lane_mask member_mask = static_cast<lane_mask>(-1) >> (__AMDGCN_WAVEFRONT_SIZE - masklength);
lane_mask member_mask = static_cast<lane_mask>(-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<int>(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
projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h
+13 -2
مشاهده پرونده
@@ -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));
}
projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h
+4 -4
مشاهده پرونده
@@ -230,7 +230,7 @@ unsigned long long __match_all_sync(MaskT mask, T value, int* pred) {
template <typename MaskT, typename T>
__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<MaskT>::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 <typename MaskT, typename T>
__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<MaskT>::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 <typename MaskT, typename T>
__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<MaskT>::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 <typename MaskT, typename T>
__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<MaskT>::value && sizeof(MaskT) == 8,
"The mask must be a 64-bit integer. "
projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+12 -8
مشاهده پرونده
@@ -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<lane_mask>(x), add);
counter = __builtin_amdgcn_mbcnt_hi(static_cast<lane_mask>(x >> 32), counter);
#endif
} else {
counter = __builtin_amdgcn_mbcnt_lo(static_cast<lane_mask>(x), add);
counter = __builtin_amdgcn_mbcnt_hi(static_cast<lane_mask>(x >> 32), counter);
}
return counter;
return counter;
}
} // namespace coalesced_group