From 4830979f0e98b9dac9de06fdaa708a5503fe31b9 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Thu, 13 Nov 2025 16:20:27 -0800 Subject: [PATCH] SWDEV-548892 - Stop using ocml fma wrappers (#1702) Directly use elementwise builtin --- .../clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) 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 8869a44723..84dd14c1a9 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 @@ -821,8 +821,9 @@ inline __HOST_DEVICE__ __half __hadd_sat(__half x, __half y) { return __clamp_01 inline __HOST_DEVICE__ __half __hsub_sat(__half x, __half y) { return __clamp_01(__hsub(x, y)); } inline __HOST_DEVICE__ __half __hmul_sat(__half x, __half y) { return __clamp_01(__hmul(x, y)); } inline __device__ __half __hfma(__half x, __half y, __half z) { - return __half_raw{__ocml_fma_f16(static_cast<__half_raw>(x).data, static_cast<__half_raw>(y).data, - static_cast<__half_raw>(z).data)}; + return __half_raw{__builtin_elementwise_fma(static_cast<__half_raw>(x).data, + static_cast<__half_raw>(y).data, + static_cast<__half_raw>(z).data)}; } inline __device__ __half __hfma_sat(__half x, __half y, __half z) { return __clamp_01(__hfma(x, y, z)); @@ -867,7 +868,9 @@ inline __HOST_DEVICE__ __half2 __hmul2_sat(__half2 x, __half2 y) { return __half2{__clamp_01(__half_raw{r.data.x}), __clamp_01(__half_raw{r.data.y})}; } inline __device__ __half2 __hfma2(__half2 x, __half2 y, __half2 z) { - return __half2{__ocml_fma_2f16(x, y, z)}; + return __half2{__builtin_elementwise_fma(static_cast<__half2_raw>(x).data, + static_cast<__half2_raw>(y).data, + static_cast<__half2_raw>(z).data)}; } inline __device__ __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z) { auto r = static_cast<__half2_raw>(__hfma2(x, y, z));