From 9264e97cbbc2c6f1c8f261f997f44ce2430e9815 Mon Sep 17 00:00:00 2001 From: "Hernandez, Gerardo" Date: Fri, 11 Apr 2025 12:30:59 +0100 Subject: [PATCH] SWDEV-521920 - Fix compilation issues introduced by the reduce sync operations - 2 (#167) Fix pytorch 2.5 issues, by defining reduce sync operations for type __half in amd_hip_fp16.h and not in amd_warp_sync_functions.h which is problematic in case __half does not get included before that header. Only define types not supported by cuda if HIP_ENABLE_EXTRA_WARP_SYNC_TYPES is defined, to avoid portability issues [ROCm/clr commit: 66496258b4fcc5a2f1f8e51900c11dd27008fdf3] --- .../include/hip/amd_detail/amd_hip_bf16.h | 6 +- .../include/hip/amd_detail/amd_hip_fp16.h | 35 +++ .../hip/amd_detail/amd_warp_sync_functions.h | 238 ++++++++---------- projects/clr/hipamd/src/hiprtc/CMakeLists.txt | 8 +- 4 files changed, 147 insertions(+), 140 deletions(-) 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 096ce3f8aa..64f80e2173 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 @@ -696,7 +696,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_sync(const unsigned long long mask __hip_bfloat162 bf162; unsigned int ui; } u{in}; - u.ui = __shfl_sync(mask, u.ui, delta, width); + u.ui = __shfl_sync(mask, u.ui, delta, width); return u.bf162; } @@ -724,7 +724,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_up_sync(const unsigned long long m __hip_bfloat162 bf162; unsigned int ui; } u{in}; - u.ui = __shfl_up_sync(mask, u.ui, delta, width); + u.ui = __shfl_up_sync(mask, u.ui, delta, width); return u.bf162; } @@ -750,7 +750,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_xor_sync(const unsigned long long __hip_bfloat162 bf162; unsigned int ui; } u{in}; - u.ui = __shfl_xor_sync(mask, u.ui, delta, width); + u.ui = __shfl_xor_sync(mask, u.ui, delta, width); return u.bf162; } #endif // HIP_DISABLE_WARP_SYNC_BUILTINS diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 8dbc7a997c..36a40be4f6 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1870,6 +1870,41 @@ THE SOFTWARE. tmp.i = __shfl_xor(tmp.i, lane_mask, width); return tmp.h; } + + #if defined(HIP_ENABLE_EXTRA_WARP_SYNC_TYPES) && !defined(__HIP_NO_HALF_OPERATORS__) + extern "C" __device__ __attribute__((const)) __half __ockl_wfred_add_f16(__half); + extern "C" __device__ __attribute__((const)) __half __ockl_wfred_min_f16(__half); + extern "C" __device__ __attribute__((const)) __half __ockl_wfred_max_f16(__half); + + template + __device__ inline __half __reduce_add_sync(MaskT mask, __half val) + { + auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f16(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); + } + + template + __device__ inline __half __reduce_min_sync(MaskT mask, __half val) + { + auto op = [](decltype(val) lhs, decltype(val) rhs) { return rhs < lhs? rhs : lhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_min_f16(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); + } + + template + __device__ inline __half __reduce_max_sync(MaskT mask, __half val) + { + auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_f16(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); + } + + #endif // __HIP_NO_HALF_OPERATORS__ + #endif // defined(__cplusplus) #elif defined(__GNUC__) || defined(_MSC_VER) #if !defined(__HIPCC_RTC__) 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 80484b884e..9ba74fd3e1 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 @@ -28,55 +28,55 @@ THE SOFTWARE. // "HIP_ENABLE_WARP_SYNC_BUILTINS". This arrangement also applies to the // __activemask() builtin defined in amd_warp_functions.h. #ifdef HIP_ENABLE_WARP_SYNC_BUILTINS - #if !defined(__HIPCC_RTC__) #include "amd_warp_functions.h" +#include "amd_device_functions.h" #include "hip_assert.h" -#include "amd_hip_fp16.h" #include #include -#include #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); +extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_min_u32(unsigned int); +extern "C" __device__ __attribute__((const)) int __ockl_wfred_max_i32(int); +extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_max_u32(unsigned int); +extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_and_u32(unsigned int); +extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_or_u32(unsigned int); +extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_xor_u32(unsigned int); + +#ifdef HIP_ENABLE_EXTRA_WARP_SYNC_TYPES +// this macro enable types that are not in CUDA extern "C" __device__ __attribute__((const)) long long __ockl_wfred_add_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_add_u64(unsigned long long); -extern "C" __device__ __attribute__((const)) __half __ockl_wfred_add_f16(__half); extern "C" __device__ __attribute__((const)) float __ockl_wfred_add_f32(float); extern "C" __device__ __attribute__((const)) double __ockl_wfred_add_f64(double); -extern "C" __device__ __attribute__((const)) int __ockl_wfred_min_i32(int); -extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_min_u32(unsigned int); extern "C" __device__ __attribute__((const)) long long __ockl_wfred_min_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_min_u64(unsigned long long); -extern "C" __device__ __attribute__((const)) __half __ockl_wfred_min_f16(__half); extern "C" __device__ __attribute__((const)) float __ockl_wfred_min_f32(float); extern "C" __device__ __attribute__((const)) double __ockl_wfred_min_f64(double); -extern "C" __device__ __attribute__((const)) int __ockl_wfred_max_i32(int); -extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_max_u32(unsigned int); extern "C" __device__ __attribute__((const)) long long __ockl_wfred_max_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_max_u64(unsigned long long); -extern "C" __device__ __attribute__((const)) __half __ockl_wfred_max_f16(__half); extern "C" __device__ __attribute__((const)) float __ockl_wfred_max_f32(float); extern "C" __device__ __attribute__((const)) double __ockl_wfred_max_f64(double); extern "C" __device__ __attribute__((const)) int __ockl_wfred_and_i32(int); -extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_and_u32(unsigned int); extern "C" __device__ __attribute__((const)) long long __ockl_wfred_and_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_and_u64(unsigned long long); extern "C" __device__ __attribute__((const)) int __ockl_wfred_or_i32(int); -extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_or_u32(unsigned int); extern "C" __device__ __attribute__((const)) long long __ockl_wfred_or_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_or_u64(unsigned long long); extern "C" __device__ __attribute__((const)) int __ockl_wfred_xor_i32(int); -extern "C" __device__ __attribute__((const)) unsigned int __ockl_wfred_xor_u32(unsigned int); extern "C" __device__ __attribute__((const)) long long __ockl_wfred_xor_i64(long long); extern "C" __device__ __attribute__((const)) unsigned long long __ockl_wfred_xor_u64(unsigned long long); +#endif + template __device__ inline T __hip_readfirstlane(T val) { @@ -325,6 +325,7 @@ T __shfl_xor_sync(MaskT mask, T var, int laneMask, template __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wfReduce) { + using permuteType = __hip_internal::conditional::type; static constexpr auto kMaskNumBits = sizeof(MaskT) * 8; static_assert( __hip_internal::is_integral::value && sizeof(MaskT) == 8, @@ -343,9 +344,9 @@ __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wf int lastLane = kMaskNumBits - leadingZeroes - 1; int maskNumBits; int numIterations; - // int[2] is used when T is 64-bit wide - typename __hip_internal::conditional::type result, permuteResult; - auto backwardPermute = [](auto index, auto val) { + // unsigned int[2] is used when T is 64-bit wide + typename __hip_internal::conditional::type result, permuteResult; + auto backwardPermute = [](int index, permuteType val) { if constexpr (std::is_integral::value || std::is_same::value) return __hip_ds_bpermute(index, val); else @@ -402,8 +403,8 @@ __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wf } } - if constexpr (std::is_same::value) { - union { int i; __half f; } tmp; + if constexpr (sizeof(T) == 2) { + union { int i; T f; } tmp; tmp.f = result; tmp.i = __hip_ds_bpermute(nextBit << 2, tmp.i); @@ -434,8 +435,8 @@ __device__ inline T __reduce_op_sync(MaskT mask, T val, BinaryOp op, WfReduce wf numIterations--; } - if constexpr (std::is_same::value) { - union { int i; __half f; } tmp; + if constexpr (sizeof(T) == 2) { + union { int i; T f; } tmp; tmp.f = result; tmp.i = __hip_ds_bpermute(firstLane << 2, tmp.i); return tmp.f; @@ -452,7 +453,7 @@ template __device__ inline int __reduce_add_sync(MaskT mask, int val) { // although C++ has std::plus and other functors, we do not use them because - // the they are header and they were causing problem with hipRTC + // they are in the header and they were causing problem with hipRTC // at this time auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_i32(v); }; @@ -469,51 +470,6 @@ __device__ inline unsigned int __reduce_add_sync(MaskT mask, unsigned int val) return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline long long __reduce_add_sync(MaskT mask, long long val) -{ - auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_i64(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - -template -__device__ inline unsigned long long __reduce_add_sync(MaskT mask, unsigned long long val) -{ - auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_u64(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - -template -__device__ inline __half __reduce_add_sync(MaskT mask, __half val) -{ - auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f16(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - -template -__device__ inline float __reduce_add_sync(MaskT mask, float val) -{ - auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - -template -__device__ inline double __reduce_add_sync(MaskT mask, double val) -{ - auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f64(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline int __reduce_min_sync(MaskT mask, int val) { @@ -532,6 +488,88 @@ __device__ inline unsigned int __reduce_min_sync(MaskT mask, unsigned int val) return __reduce_op_sync(mask, val, op, wfReduce); } +template +__device__ inline int __reduce_max_sync(MaskT mask, int val) +{ + auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_i32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline unsigned int __reduce_max_sync(MaskT mask, unsigned int val) +{ + auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_u32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline unsigned int __reduce_or_sync(MaskT mask, unsigned int val) +{ + auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs || rhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_or_u32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline unsigned int __reduce_and_sync(MaskT mask, unsigned int val) +{ + auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs && rhs; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_and_u32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline unsigned int __reduce_xor_sync(MaskT mask, unsigned int val) +{ + auto op = [](decltype(val) lhs, decltype(val) rhs) { return (!lhs) != (!rhs) == 1; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_xor_u32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +#ifdef HIP_ENABLE_EXTRA_WARP_SYNC_TYPES +template +__device__ inline long long __reduce_add_sync(MaskT mask, long long val) +{ + auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_i64(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline unsigned long long __reduce_add_sync(MaskT mask, unsigned long long val) +{ + auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_u64(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline float __reduce_add_sync(MaskT mask, float val) +{ + auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f32(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + +template +__device__ inline double __reduce_add_sync(MaskT mask, double val) +{ + auto op = [](decltype(val)& a, decltype(val)& b) { return a + b; }; + auto wfReduce = [](decltype(val) v) { return __ockl_wfred_add_f64(v); }; + + return __reduce_op_sync(mask, val, op, wfReduce); +} + template __device__ inline long long __reduce_min_sync(MaskT mask, long long val) { @@ -550,15 +588,6 @@ __device__ inline unsigned long long __reduce_min_sync(MaskT mask, unsigned long return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline __half __reduce_min_sync(MaskT mask, __half val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return rhs < lhs? rhs : lhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_min_f16(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline float __reduce_min_sync(MaskT mask, float val) { @@ -577,24 +606,6 @@ __device__ inline double __reduce_min_sync(MaskT mask, double val) return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline int __reduce_max_sync(MaskT mask, int val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_i32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - -template -__device__ inline unsigned int __reduce_max_sync(MaskT mask, unsigned int val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_u32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline long long __reduce_max_sync(MaskT mask, long long val) { @@ -613,15 +624,6 @@ __device__ inline unsigned long long __reduce_max_sync(MaskT mask, unsigned long return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline __half __reduce_max_sync(MaskT mask, __half val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs < rhs? rhs : lhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_max_f16(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline float __reduce_max_sync(MaskT mask, float val) { @@ -649,15 +651,6 @@ __device__ inline int __reduce_and_sync(MaskT mask, int val) return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline unsigned int __reduce_and_sync(MaskT mask, unsigned int val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs && rhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_and_u32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline long long __reduce_and_sync(MaskT mask, long long val) { @@ -685,15 +678,6 @@ __device__ inline int __reduce_or_sync(MaskT mask, int val) return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline unsigned int __reduce_or_sync(MaskT mask, unsigned int val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return lhs || rhs; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_or_u32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline long long __reduce_or_sync(MaskT mask, long long val) { @@ -721,15 +705,6 @@ __device__ inline int __reduce_xor_sync(MaskT mask, int val) return __reduce_op_sync(mask, val, op, wfReduce); } -template -__device__ inline unsigned int __reduce_xor_sync(MaskT mask, unsigned int val) -{ - auto op = [](decltype(val) lhs, decltype(val) rhs) { return (!lhs) != (!rhs) == 1; }; - auto wfReduce = [](decltype(val) v) { return __ockl_wfred_xor_u32(v); }; - - return __reduce_op_sync(mask, val, op, wfReduce); -} - template __device__ inline long long __reduce_xor_sync(MaskT mask, long long val) { @@ -748,8 +723,5 @@ __device__ inline unsigned long long __reduce_xor_sync(MaskT mask, unsigned long return __reduce_op_sync(mask, val, op, wfReduce); } -#undef __hip_do_sync -#undef __hip_check_mask -#undef __hip_adjust_mask_for_wave32 - +#endif // HIP_ENABLE_EXTRA_WARP_SYNC_TYPES #endif // HIP_ENABLE_WARP_SYNC_BUILTINS diff --git a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt index f41573d66a..ef473f6237 100644 --- a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt +++ b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt @@ -175,15 +175,15 @@ ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_math_constants.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_device_functions.h -${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h -${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h -${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h -${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_math_functions.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_sync_functions.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_cooperative_groups_helper.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_cooperative_groups.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_unsafe_atomics.h ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_atomic.h +${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h +${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h +${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h +${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_math_functions.h ) # Generate required HIPRTC files.