diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h index 735f915bd2..5acf114518 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -30,10 +30,41 @@ THE SOFTWARE. #ifdef HIP_HALF_HW_SUPPORT typedef __fp16 __half; -extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b); +extern "C" __half __hip_hadd_gfx803(__half a, __half b); +extern "C" __half __hip_hfma_gfx803(__half a, __half b); +extern "C" __half __hip_hmul_gfx803(__half a, __half b); +extern "C" __half __hip_hsub_gfx803(__half a, __half b); -__device__ inline __half __hadd(__half a, __half b){ - return __hip_hadd_clang40_gfx803(a, b); +__device__ inline __half __hadd(__half a, __half b) { + return __hip_hadd_gfx803(a, b); +} + +__device__ inline __half __hadd_sat(__half a, __half b) { + return __hip_add_gfx803(a, b); +} + +__device__ inline __half __hfma(__half a, __half b) { + return __hip_hfma_gfx803(a, b); +} + +__device__ inline __half __hfma_sat(__half a, __half b) { + return __hip_hfma_gfx803(a, b); +} + +__device__ inline __half __hmul(__half a, __half b) { + return __hip_hmul_gfx803(a, b); +} + +__device__ inline __half __hmul_sat(__half a, __half b) { + return __hip_hmul_gfx803(a, b); +} + +__device__ inline __half __hsub(__half a, __half b) { + return __hip_hsub_gfx803(a, b); +} + +__device__ inline __half __hsub_sat(__half a, __half b) { + return __hip_hsub_gfx803(a, b); } #else diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll index 078dc3eed5..623be19084 100644 --- a/projects/clr/hipamd/src/hip_ir.ll +++ b/projects/clr/hipamd/src/hip_ir.ll @@ -34,9 +34,24 @@ define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) { ret i32 %val } -define linkonce_odr spir_func half @__hip_hadd_clang40_gfx803(half %a, half %b) { +define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 { %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b) ret half %val } +define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 { + %val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c) + ret half %val +} + +define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 { + %val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b) + ret half %val +} + +define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 { + %val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b) + ret half %val +} + attributes #1 = { alwaysinline nounwind }