From d4a12ff5e59e843027e0ead64dd49aebe02afdb9 Mon Sep 17 00:00:00 2001 From: Li Li Date: Wed, 26 Oct 2022 16:16:16 +0000 Subject: [PATCH] adjust to adapt ROCm5.4 --- fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp b/fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp index 03d0471e19..66755c0ce7 100644 --- a/fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp +++ b/fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp @@ -60,7 +60,7 @@ __device__ half llvm_amdgcn_raw_buffer_load_fp16(int32x4_t srsrc, int32_t voffset, int32_t soffset, - int32_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16"); + int32_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16"); __device__ float llvm_amdgcn_raw_buffer_load_fp32(int32x4_t srsrc, @@ -72,7 +72,7 @@ __device__ half2 llvm_amdgcn_raw_buffer_load_fp16x2(int32x4_t srsrc, int32_t voffset, int32_t soffset, - int32_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16"); + int32_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32"); __device__ floatx2_t llvm_amdgcn_raw_buffer_load_fp32x2(int32x4_t srsrc,