Skip to content

Commit

Permalink
adjust to adapt ROCm5.4
Browse files Browse the repository at this point in the history
  • Loading branch information
@liligwu
liligwu committed Oct 26, 2022
1 parent 9d181e5 commit d4a12ff
Showing 1 changed file with 2 additions and 2 deletions.
4 changes: 2 additions & 2 deletions fbgemm_gpu/hip_kernel/split_tbe_fwd_hip.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -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,
Expand Down

1 comment on commit d4a12ff

@whchung
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@liligwu From the look of it, it could be a bug in the compiler.

Per discussion with Xiao Hai, a ticket to the compiler team with steps to reproduce could be helpful.

Please sign in to comment.