From 1db9a7d48b702c0442348201fabf7a92ab929fc7 Mon Sep 17 00:00:00 2001 From: "Arsenault, Matthew" Date: Mon, 5 May 2025 19:35:12 +0200 Subject: [PATCH] SWDEV-1 - Stop using ocml rounding functions (#228) Directly use the builtins. Use the elementwise versions since there's no implied errno, regardless of -f[no]-math-errno. I didn't change the cases unnecessarily casting. The bfloat and vector cases should work directly. --- .../hip/amd_detail/amd_device_functions.h | 56 ++++++++++--------- hipamd/include/hip/amd_detail/amd_hip_bf16.h | 8 +-- hipamd/include/hip/amd_detail/amd_hip_fp16.h | 20 +++---- .../hip/amd_detail/device_library_decls.h | 5 -- 4 files changed, 43 insertions(+), 46 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index 36353096f3..794efb8a41 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -404,41 +404,47 @@ __device__ static inline int __double2loint(double x) { return tmp[0]; } -__device__ static inline int __double2int_rd(double x) { return (int)__ocml_floor_f64(x); } -__device__ static inline int __double2int_rn(double x) { return (int)__ocml_rint_f64(x); } -__device__ static inline int __double2int_ru(double x) { return (int)__ocml_ceil_f64(x); } +__device__ static inline int __double2int_rd(double x) { + return (int)__builtin_elementwise_floor(x); +} +__device__ static inline int __double2int_rn(double x) { + return (int)__builtin_elementwise_rint(x); +} +__device__ static inline int __double2int_ru(double x) { + return (int)__builtin_elementwise_ceil(x); +} __device__ static inline int __double2int_rz(double x) { return (int)x; } __device__ static inline long long int __double2ll_rd(double x) { - return (long long)__ocml_floor_f64(x); + return (long long)__builtin_elementwise_floor(x); } __device__ static inline long long int __double2ll_rn(double x) { - return (long long)__ocml_rint_f64(x); + return (long long)__builtin_elementwise_rint(x); } __device__ static inline long long int __double2ll_ru(double x) { - return (long long)__ocml_ceil_f64(x); + return (long long)__builtin_elementwise_ceil(x); } __device__ static inline long long int __double2ll_rz(double x) { return (long long)x; } __device__ static inline unsigned int __double2uint_rd(double x) { - return (unsigned int)__ocml_floor_f64(x); + return (unsigned int)__builtin_elementwise_floor(x); } __device__ static inline unsigned int __double2uint_rn(double x) { - return (unsigned int)__ocml_rint_f64(x); + return (unsigned int)__builtin_elementwise_rint(x); } __device__ static inline unsigned int __double2uint_ru(double x) { - return (unsigned int)__ocml_ceil_f64(x); + return (unsigned int)__builtin_elementwise_ceil(x); } __device__ static inline unsigned int __double2uint_rz(double x) { return (unsigned int)x; } __device__ static inline unsigned long long int __double2ull_rd(double x) { - return (unsigned long long int)__ocml_floor_f64(x); + return (unsigned long long int)__builtin_elementwise_floor(x); } __device__ static inline unsigned long long int __double2ull_rn(double x) { - return (unsigned long long int)__ocml_rint_f64(x); + return (unsigned long long int)__builtin_elementwise_rint(x); } __device__ static inline unsigned long long int __double2ull_ru(double x) { - return (unsigned long long int)__ocml_ceil_f64(x); + return (unsigned long long int)__builtin_elementwise_ceil(x); } __device__ static inline unsigned long long int __double2ull_rz(double x) { return (unsigned long long int)x; @@ -466,41 +472,41 @@ CUDA implements half as unsigned short whereas, HIP doesn't. */ -__device__ static inline int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); } -__device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); } -__device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); } -__device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); } +__device__ static inline int __float2int_rd(float x) { return (int)__builtin_elementwise_floor(x); } +__device__ static inline int __float2int_rn(float x) { return (int)__builtin_elementwise_rint(x); } +__device__ static inline int __float2int_ru(float x) { return (int)__builtin_elementwise_ceil(x); } +__device__ static inline int __float2int_rz(float x) { return (int)__builtin_elementwise_trunc(x); } __device__ static inline long long int __float2ll_rd(float x) { - return (long long int)__ocml_floor_f32(x); + return (long long int)__builtin_elementwise_floor(x); } __device__ static inline long long int __float2ll_rn(float x) { - return (long long int)__ocml_rint_f32(x); + return (long long int)__builtin_elementwise_rint(x); } __device__ static inline long long int __float2ll_ru(float x) { - return (long long int)__ocml_ceil_f32(x); + return (long long int)__builtin_elementwise_ceil(x); } __device__ static inline long long int __float2ll_rz(float x) { return (long long int)x; } __device__ static inline unsigned int __float2uint_rd(float x) { - return (unsigned int)__ocml_floor_f32(x); + return (unsigned int)__builtin_elementwise_floor(x); } __device__ static inline unsigned int __float2uint_rn(float x) { - return (unsigned int)__ocml_rint_f32(x); + return (unsigned int)__builtin_elementwise_rint(x); } __device__ static inline unsigned int __float2uint_ru(float x) { - return (unsigned int)__ocml_ceil_f32(x); + return (unsigned int)__builtin_elementwise_ceil(x); } __device__ static inline unsigned int __float2uint_rz(float x) { return (unsigned int)x; } __device__ static inline unsigned long long int __float2ull_rd(float x) { - return (unsigned long long int)__ocml_floor_f32(x); + return (unsigned long long int)__builtin_elementwise_floor(x); } __device__ static inline unsigned long long int __float2ull_rn(float x) { - return (unsigned long long int)__ocml_rint_f32(x); + return (unsigned long long int)__builtin_elementwise_rint(x); } __device__ static inline unsigned long long int __float2ull_ru(float x) { - return (unsigned long long int)__ocml_ceil_f32(x); + return (unsigned long long int)__builtin_elementwise_ceil(x); } __device__ static inline unsigned long long int __float2ull_rz(float x) { return (unsigned long long int)x; diff --git a/hipamd/include/hip/amd_detail/amd_hip_bf16.h b/hipamd/include/hip/amd_detail/amd_hip_bf16.h index 85626fe860..405e320bb5 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_bf16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_bf16.h @@ -1548,7 +1548,7 @@ __BF16_HOST_DEVICE_STATIC__ bool operator>=(const __hip_bfloat162& l, const __hi * \brief Calculate ceil of bfloat16 */ __BF16_DEVICE_STATIC__ __hip_bfloat16 hceil(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_ceil_f32(__bfloat162float(h))); + return __float2bfloat16(__builtin_elementwise_ceil(__bfloat162float(h))); } /** @@ -1588,7 +1588,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hexp2(const __hip_bfloat16 h) { * \brief Calculate floor of bfloat16 */ __BF16_DEVICE_STATIC__ __hip_bfloat16 hfloor(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_floor_f32(__bfloat162float(h))); + return __float2bfloat16(__builtin_elementwise_floor(__bfloat162float(h))); } /** @@ -1628,7 +1628,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hrcp(const __hip_bfloat16 h) { * \brief Round to nearest int */ __BF16_DEVICE_STATIC__ __hip_bfloat16 hrint(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_rint_f32(__bfloat162float(h))); + return __float2bfloat16(__builtin_elementwise_rint(__bfloat162float(h))); } /** @@ -1660,7 +1660,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hsqrt(const __hip_bfloat16 h) { * \brief Calculate truncate of bfloat16 */ __BF16_DEVICE_STATIC__ __hip_bfloat16 htrunc(const __hip_bfloat16 h) { - return __float2bfloat16(__ocml_trunc_f32(__bfloat162float(h))); + return __float2bfloat16(__builtin_elementwise_trunc(__bfloat162float(h))); } /** diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 36a40be4f6..994478d23e 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1593,29 +1593,25 @@ THE SOFTWARE. __device__ __half htrunc(__half x) { - return __half_raw{ - __ocml_trunc_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_trunc(static_cast<__half_raw>(x).data)}; } inline __device__ __half hceil(__half x) { - return __half_raw{ - __ocml_ceil_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_ceil(static_cast<__half_raw>(x).data)}; } inline __device__ __half hfloor(__half x) { - return __half_raw{ - __ocml_floor_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_floor(static_cast<__half_raw>(x).data)}; } inline __device__ __half hrint(__half x) { - return __half_raw{ - __ocml_rint_f16(static_cast<__half_raw>(x).data)}; + return __half_raw{__builtin_elementwise_rint(static_cast<__half_raw>(x).data)}; } inline __device__ @@ -1713,25 +1709,25 @@ THE SOFTWARE. __device__ __half2 h2trunc(__half2 x) { - return __half2{__ocml_trunc_2f16(x)}; + return __half2{__builtin_elementwise_trunc(static_cast<__half2_raw>(x).data)}; } inline __device__ __half2 h2ceil(__half2 x) { - return __half2{__ocml_ceil_2f16(x)}; + return __half2{__builtin_elementwise_ceil(static_cast<__half2_raw>(x).data)}; } inline __device__ __half2 h2floor(__half2 x) { - return __half2{__ocml_floor_2f16(x)}; + return __half2{__builtin_elementwise_floor(static_cast<__half2_raw>(x).data)}; } inline __device__ __half2 h2rint(__half2 x) { - return __half2{__ocml_rint_2f16(x)}; + return __half2{__builtin_elementwise_rint(static_cast<__half2_raw>(x).data)}; } inline __device__ diff --git a/hipamd/include/hip/amd_detail/device_library_decls.h b/hipamd/include/hip/amd_detail/device_library_decls.h index 81f645052c..48abee3653 100644 --- a/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/hipamd/include/hip/amd_detail/device_library_decls.h @@ -61,11 +61,6 @@ extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort); extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint); extern "C" __device__ __attribute__((const)) uint64_t __ockl_clz_u64(uint64_t); -extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float); -extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float); -extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float); -extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float); - extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float); extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);