Skip to content

Commit eeb7688

Browse files
arsenmpravinjagtap
arsenm
and
pravinjagtap
authored
AMDGPU: Builtins & CodeGen support for v_cvt_scalef32_pk_{f|bf}16_fp4 for gfx950 (#117744)
OPSEL ASM Syntax for v_cvt_scalef32_pk_{f|bf}16_fp4 : opsel:[x,y,z] where, x & y i.e. OPSEL[1 : 0] selects which src_byte to read. Note: Conventional Inst{13} i.e. OPSEL[2] is ignored in asm syntax. Co-authored-by: Pravin Jagtap <[email protected]>
1 parent 2b9e947 commit eeb7688

File tree

8 files changed

+178
-6
lines changed

8 files changed

+178
-6
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -584,6 +584,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_f16, "V2sV2sV2hfIb", "nc", "
584584
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16, "V2sV2sV2yfIb", "nc", "bf8-cvt-scale-insts")
585585
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_fp4, "V2fUifIi", "nc", "fp4-cvt-scale-insts")
586586
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp4_f32, "UiUifffIi", "nc", "fp4-cvt-scale-insts")
587+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f16_fp4, "V2hUifIi", "nc", "fp4-cvt-scale-insts")
588+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4, "V2yUifIi", "nc", "fp4-cvt-scale-insts")
587589

588590
#undef BUILTIN
589591
#undef TARGET_BUILTIN

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@ typedef float __attribute__((ext_vector_type(2))) float2;
1818
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
1919

2020
void test(global uint* out, global uint2* out_v2u32, uint a, uint b, global half2* out_v2f16, global float* out_f32, float scale, global short2* out_v2i16, float src0, float src1,
21-
global float2* out_v2f32, half2 src0_v2f16, bfloat2 src0_v2bf16) {
21+
global float2* out_v2f32, half2 src0_v2f16, bfloat2 src0_v2bf16, global bfloat2* out_v2bf16) {
2222
*out = __builtin_amdgcn_prng_b32(a); // expected-error{{'__builtin_amdgcn_prng_b32' needs target feature prng-inst}}
2323
*out_v2u32 = __builtin_amdgcn_permlane16_swap(a, b, false, false); // expected-error{{'__builtin_amdgcn_permlane16_swap' needs target feature permlane16-swap}}
2424
*out_v2u32 = __builtin_amdgcn_permlane32_swap(a, b, false, false); // expected-error{{'__builtin_amdgcn_permlane32_swap' needs target feature permlane32-swap}}
@@ -36,4 +36,6 @@ void test(global uint* out, global uint2* out_v2u32, uint a, uint b, global half
3636
*out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(*out_v2i16, src0_v2bf16, scale, true); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16' needs target feature bf8-cvt-scale-insts}}
3737
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(a, scale, 3); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_f32_fp4' needs target feature fp4-cvt-scale-insts}}
3838
*out = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(*out, src0, src1, scale, 3); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_fp4_f32' needs target feature fp4-cvt-scale-insts}}
39+
*out_v2f16 = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(a, scale, 3); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_f16_fp4' needs target feature fp4-cvt-scale-insts}}
40+
*out_v2bf16 = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(a, scale, 3); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4' needs target feature fp4-cvt-scale-insts}}
3941
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -829,3 +829,79 @@ void test_cvt_scalef32_pk_fp4_f32(global unsigned int* out, float src0, float sr
829829
*out = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(*out, src0, src1, scale, 2);
830830
*out = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(*out, src0, src1, scale, 3);
831831
}
832+
833+
// CHECK-LABEL: @test_cvt_scalef32_pk_f16_fp4(
834+
// CHECK-NEXT: entry:
835+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
836+
// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
837+
// CHECK-NEXT: [[SCALE_ADDR:%.*]] = alloca float, align 4, addrspace(5)
838+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
839+
// CHECK-NEXT: store i32 [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align 4
840+
// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]], align 4
841+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
842+
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
843+
// CHECK-NEXT: [[TMP2:%.*]] = call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 [[TMP0]], float [[TMP1]], i32 0)
844+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
845+
// CHECK-NEXT: store <2 x half> [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
846+
// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
847+
// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
848+
// CHECK-NEXT: [[TMP6:%.*]] = call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 [[TMP4]], float [[TMP5]], i32 1)
849+
// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
850+
// CHECK-NEXT: store <2 x half> [[TMP6]], ptr addrspace(1) [[TMP7]], align 4
851+
// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
852+
// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
853+
// CHECK-NEXT: [[TMP10:%.*]] = call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 [[TMP8]], float [[TMP9]], i32 2)
854+
// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
855+
// CHECK-NEXT: store <2 x half> [[TMP10]], ptr addrspace(1) [[TMP11]], align 4
856+
// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
857+
// CHECK-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
858+
// CHECK-NEXT: [[TMP14:%.*]] = call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 [[TMP12]], float [[TMP13]], i32 3)
859+
// CHECK-NEXT: [[TMP15:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
860+
// CHECK-NEXT: store <2 x half> [[TMP14]], ptr addrspace(1) [[TMP15]], align 4
861+
// CHECK-NEXT: ret void
862+
//
863+
void test_cvt_scalef32_pk_f16_fp4(global half2* out, uint src, float scale)
864+
{
865+
*out = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 0);
866+
*out = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 1);
867+
*out = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 2);
868+
*out = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 3);
869+
}
870+
871+
// CHECK-LABEL: @test_cvt_scalef32_pk_bf16_fp4(
872+
// CHECK-NEXT: entry:
873+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
874+
// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
875+
// CHECK-NEXT: [[SCALE_ADDR:%.*]] = alloca float, align 4, addrspace(5)
876+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
877+
// CHECK-NEXT: store i32 [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align 4
878+
// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]], align 4
879+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
880+
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
881+
// CHECK-NEXT: [[TMP2:%.*]] = call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 [[TMP0]], float [[TMP1]], i32 0)
882+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
883+
// CHECK-NEXT: store <2 x bfloat> [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
884+
// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
885+
// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
886+
// CHECK-NEXT: [[TMP6:%.*]] = call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 [[TMP4]], float [[TMP5]], i32 1)
887+
// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
888+
// CHECK-NEXT: store <2 x bfloat> [[TMP6]], ptr addrspace(1) [[TMP7]], align 4
889+
// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
890+
// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
891+
// CHECK-NEXT: [[TMP10:%.*]] = call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 [[TMP8]], float [[TMP9]], i32 2)
892+
// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
893+
// CHECK-NEXT: store <2 x bfloat> [[TMP10]], ptr addrspace(1) [[TMP11]], align 4
894+
// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr addrspace(5) [[SRC_ADDR]], align 4
895+
// CHECK-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]], align 4
896+
// CHECK-NEXT: [[TMP14:%.*]] = call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 [[TMP12]], float [[TMP13]], i32 3)
897+
// CHECK-NEXT: [[TMP15:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
898+
// CHECK-NEXT: store <2 x bfloat> [[TMP14]], ptr addrspace(1) [[TMP15]], align 4
899+
// CHECK-NEXT: ret void
900+
//
901+
void test_cvt_scalef32_pk_bf16_fp4(global bfloat2* out, uint src, float scale)
902+
{
903+
*out = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 0);
904+
*out = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 1);
905+
*out = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 2);
906+
*out = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 3);
907+
}

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -165,7 +165,7 @@ void test_permlane32_swap(__global int* out, int old, int src, bool X) {
165165

166166
void test_cvt_scalef32(global half2* out_v2f16, global float* out_f32, uint src, float scale, int index, bool X,
167167
global short2* out_v2i16, float src0, float src1, global float2* out_v2f32,
168-
half2 src0_v2f16, bfloat2 src0_v2bf16, global uint* out) {
168+
half2 src0_v2f16, bfloat2 src0_v2bf16, global uint* out, global bfloat2* out_v2bf16) {
169169
*out_v2f16 = __builtin_amdgcn_cvt_scalef32_f16_fp8(*out_v2f16, src, scale, index, X); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_f16_fp8' must be a constant integer}}
170170
*out_f32 = __builtin_amdgcn_cvt_scalef32_f32_fp8(src, scale, index); // // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_f32_fp8' must be a constant integer}}
171171
*out_v2f16 = __builtin_amdgcn_cvt_scalef32_f16_bf8(*out_v2f16, src, scale, index, X); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_f16_bf8' must be a constant integer}}
@@ -180,4 +180,6 @@ void test_cvt_scalef32(global half2* out_v2f16, global float* out_f32, uint src,
180180
*out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(*out_v2i16, src0_v2bf16, scale, X); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16' must be a constant integer}}
181181
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(src, scale, index); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_f32_fp4' must be a constant integer}}
182182
*out = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(*out, src0, src1, scale, index); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_fp4_f32' must be a constant integer}}
183+
*out_v2f16 = __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, index); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_f16_fp4' must be a constant integer}}
184+
*out_v2bf16 = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, index); // expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4' must be a constant integer}}
183185
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -609,7 +609,7 @@ def int_amdgcn_cvt_scalef32_pk32_bf6_bf16 : AMDGPUCvtScaleF32Intrinsic<llvm_v6i3
609609
def int_amdgcn_cvt_scalef32_2xpk16_fp6_f32 : AMDGPUCvtScaleF32ToFP6BF6Intrinsic<llvm_v6i32_ty, llvm_v16f32_ty, llvm_v16f32_ty, "cvt_scalef32_2xpk16_fp6_f32">;
610610
def int_amdgcn_cvt_scalef32_2xpk16_bf6_f32 : AMDGPUCvtScaleF32ToFP6BF6Intrinsic<llvm_v6i32_ty, llvm_v16f32_ty, llvm_v16f32_ty, "cvt_scalef32_2xpk16_bf6_f32">;
611611

612-
class AMDGPUCvtScaleFP8BF8ToF32Intrinsic<LLVMType DstTy, string name> : DefaultAttrsIntrinsic<
612+
class AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<LLVMType DstTy, string name> : DefaultAttrsIntrinsic<
613613
[DstTy],
614614
[llvm_i32_ty, // src
615615
llvm_float_ty, // scale
@@ -669,8 +669,8 @@ def int_amdgcn_cvt_scalef32_f16_fp8 : AMDGPUCvtScaleFP8BF8ToF16TiedInputIntrins
669669
def int_amdgcn_cvt_scalef32_f16_bf8 : AMDGPUCvtScaleFP8BF8ToF16TiedInputIntrinsic<llvm_v2f16_ty, "cvt_scalef32_f16_bf8">;
670670

671671
// llvm.amdgcn.cvt.scalef32.f32.fp8 int src, float scale, int src_sel_index [0..3]
672-
def int_amdgcn_cvt_scalef32_f32_fp8 : AMDGPUCvtScaleFP8BF8ToF32Intrinsic<llvm_float_ty, "cvt_scalef32_f32_fp8">;
673-
def int_amdgcn_cvt_scalef32_f32_bf8 : AMDGPUCvtScaleFP8BF8ToF32Intrinsic<llvm_float_ty, "cvt_scalef32_f32_bf8">;
672+
def int_amdgcn_cvt_scalef32_f32_fp8 : AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<llvm_float_ty, "cvt_scalef32_f32_fp8">;
673+
def int_amdgcn_cvt_scalef32_f32_bf8 : AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<llvm_float_ty, "cvt_scalef32_f32_bf8">;
674674

675675
// llvm.amdgcn.cvt.scalef32.pk.fp8.f32 v2i16 old_vdst, float srcA, float srcB, float scale, bool dst_lo_hi_sel
676676
def int_amdgcn_cvt_scalef32_pk_fp8_f32 : AMDGPUCvtScaleF32ToFP8BF8TiedInputIntrinsic<"cvt_scalef32_pk_fp8_f32">;
@@ -687,11 +687,15 @@ def int_amdgcn_cvt_scalef32_pk_bf8_f16 : AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputI
687687
def int_amdgcn_cvt_scalef32_pk_bf8_bf16: AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<llvm_v2bf16_ty, "cvt_scalef32_pk_bf8_bf16">;
688688

689689
// llvm.amdgcn.cvt.scalef32.pk.f32.fp4 int src, float scale, int src_sel_index [0..3]
690-
def int_amdgcn_cvt_scalef32_pk_f32_fp4 : AMDGPUCvtScaleFP8BF8ToF32Intrinsic<llvm_v2f32_ty, "cvt_scalef32_pk_f32_fp4">;
690+
def int_amdgcn_cvt_scalef32_pk_f32_fp4 : AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<llvm_v2f32_ty, "cvt_scalef32_pk_f32_fp4">;
691691

692692
// llvm.amdgcn.cvt.scalef32.pk.fp4.f32 i32 old_vdst, float srcA, float srcB, float scale, int dst_sel_index[0..3]
693693
def int_amdgcn_cvt_scalef32_pk_fp4_f32 : AMDGPUCvtScaleF32ToFP4Intrinsic<"cvt_scalef32_pk_fp4_f32">;
694694

695+
// llvm.amdgcn.cvt.scalef32.pk.f32.fp4 int src, float scale, int src_sel_index [0..3]
696+
def int_amdgcn_cvt_scalef32_pk_f16_fp4 : AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<llvm_v2f16_ty, "cvt_scalef32_pk_f16_fp4">;
697+
def int_amdgcn_cvt_scalef32_pk_bf16_fp4: AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic<llvm_v2bf16_ty, "cvt_scalef32_pk_bf16_fp4">;
698+
695699
def int_amdgcn_prng_b32 : DefaultAttrsIntrinsic<
696700
[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]
697701
>, ClangBuiltin<"__builtin_amdgcn_prng_b32">;

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4561,6 +4561,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
45614561
case Intrinsic::amdgcn_cvt_scalef32_pk_bf8_bf16:
45624562
case Intrinsic::amdgcn_cvt_scalef32_pk_f32_fp4:
45634563
case Intrinsic::amdgcn_cvt_scalef32_pk_fp4_f32:
4564+
case Intrinsic::amdgcn_cvt_scalef32_pk_f16_fp4:
4565+
case Intrinsic::amdgcn_cvt_scalef32_pk_bf16_fp4:
45644566
case Intrinsic::amdgcn_ashr_pk_i8_i32:
45654567
case Intrinsic::amdgcn_ashr_pk_u8_i32:
45664568
case Intrinsic::amdgcn_cvt_scalef32_2xpk16_fp6_f32:

‎llvm/lib/Target/AMDGPU/VOP3Instructions.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1090,6 +1090,8 @@ class Cvt_Scale_FP4FP8BF8ToF16F32_Pat<SDPatternOperator node, VOP3_Pseudo inst,
10901090
>;
10911091
def : Cvt_Scale_FP4FP8BF8ToF16F32_Pat<int_amdgcn_cvt_scalef32_f32_fp8, V_CVT_SCALEF32_F32_FP8_e64, f32>;
10921092
def : Cvt_Scale_FP4FP8BF8ToF16F32_Pat<int_amdgcn_cvt_scalef32_f32_bf8, V_CVT_SCALEF32_F32_BF8_e64, f32>;
1093+
def : Cvt_Scale_FP4FP8BF8ToF16F32_Pat<int_amdgcn_cvt_scalef32_pk_f16_fp4, V_CVT_SCALEF32_PK_F16_FP4_e64, v2f16>;
1094+
def : Cvt_Scale_FP4FP8BF8ToF16F32_Pat<int_amdgcn_cvt_scalef32_pk_bf16_fp4, V_CVT_SCALEF32_PK_BF16_FP4_e64, v2bf16>;
10931095

10941096
class Cvt_Scale_FP8BF8ToF16_Pat<SDPatternOperator node, VOP3_Pseudo inst, int dst_sel> : GCNPat<
10951097
(v2f16 (node v2f16:$vdst_in, i32:$src0, f32:$src1, timm:$src_sel, dst_sel)),

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx950.ll

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,8 @@ declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16>, <2 x half>, fl
1818
declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.fb16(<2 x i16>, <2 x bfloat>, float, i1)
1919
declare <2 x float> @llvm.amdgcn.cvt.scalef32.pk.f32.fp4(i32, float, i32)
2020
declare i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32, float, float, float, i32)
21+
declare <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32, float, i32)
22+
declare <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32, float, i32)
2123

2224
define amdgpu_ps void @test_scalef32_pk32_fp6_f32_vv(<16 x float> %src, float %scale, ptr addrspace(1) %out) {
2325
; GFX950-SDAG-LABEL: test_scalef32_pk32_fp6_f32_vv:
@@ -768,3 +770,83 @@ define i32 @test_cvt_scale_fp4_f32_byte3_fabs_fneg(i32 %old, float %src0, float
768770
%ret = tail call i32 @llvm.amdgcn.cvt.scalef32.pk.fp4.f32(i32 %old, float %fabs.src0, float %fneg.src1, float %scale, i32 3)
769771
ret i32 %ret
770772
}
773+
774+
define <2 x half> @test_cvt_scale_f16_fp4_byte0(i32 %src, float %scale) {
775+
; GCN-LABEL: test_cvt_scale_f16_fp4_byte0:
776+
; GCN: ; %bb.0:
777+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
778+
; GCN-NEXT: v_cvt_scalef32_pk_f16_fp4 v0, v0, v1
779+
; GCN-NEXT: s_setpc_b64 s[30:31]
780+
%ret = tail call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 %src, float %scale, i32 0)
781+
ret <2 x half> %ret
782+
}
783+
784+
define <2 x half> @test_cvt_scale_f16_fp4_byte1(i32 %src, float %scale) {
785+
; GCN-LABEL: test_cvt_scale_f16_fp4_byte1:
786+
; GCN: ; %bb.0:
787+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
788+
; GCN-NEXT: v_cvt_scalef32_pk_f16_fp4 v0, v0, v1 op_sel:[0,1,0]
789+
; GCN-NEXT: s_setpc_b64 s[30:31]
790+
%ret = tail call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 %src, float %scale, i32 1)
791+
ret <2 x half> %ret
792+
}
793+
794+
define <2 x half> @test_cvt_scale_f16_fp4_byte2(i32 %src, float %scale) {
795+
; GCN-LABEL: test_cvt_scale_f16_fp4_byte2:
796+
; GCN: ; %bb.0:
797+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
798+
; GCN-NEXT: v_cvt_scalef32_pk_f16_fp4 v0, v0, v1 op_sel:[1,0,0]
799+
; GCN-NEXT: s_setpc_b64 s[30:31]
800+
%ret = tail call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 %src, float %scale, i32 2)
801+
ret <2 x half> %ret
802+
}
803+
804+
define <2 x half> @test_cvt_scale_f16_fp4_byte3(i32 %src, float %scale) {
805+
; GCN-LABEL: test_cvt_scale_f16_fp4_byte3:
806+
; GCN: ; %bb.0:
807+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
808+
; GCN-NEXT: v_cvt_scalef32_pk_f16_fp4 v0, v0, v1 op_sel:[1,1,0]
809+
; GCN-NEXT: s_setpc_b64 s[30:31]
810+
%ret = tail call <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32 %src, float %scale, i32 3)
811+
ret <2 x half> %ret
812+
}
813+
814+
define <2 x bfloat> @test_cvt_scale_bf16_fp4_byte0(i32 %src, float %scale) {
815+
; GCN-LABEL: test_cvt_scale_bf16_fp4_byte0:
816+
; GCN: ; %bb.0:
817+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
818+
; GCN-NEXT: v_cvt_scalef32_pk_bf16_fp4 v0, v0, v1
819+
; GCN-NEXT: s_setpc_b64 s[30:31]
820+
%ret = tail call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 %src, float %scale, i32 0)
821+
ret <2 x bfloat> %ret
822+
}
823+
824+
define <2 x bfloat> @test_cvt_scale_bf16_fp4_byte1(i32 %src, float %scale) {
825+
; GCN-LABEL: test_cvt_scale_bf16_fp4_byte1:
826+
; GCN: ; %bb.0:
827+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
828+
; GCN-NEXT: v_cvt_scalef32_pk_bf16_fp4 v0, v0, v1 op_sel:[0,1,0]
829+
; GCN-NEXT: s_setpc_b64 s[30:31]
830+
%ret = tail call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 %src, float %scale, i32 1)
831+
ret <2 x bfloat> %ret
832+
}
833+
834+
define <2 x bfloat> @test_cvt_scale_bf16_fp4_byte2(i32 %src, float %scale) {
835+
; GCN-LABEL: test_cvt_scale_bf16_fp4_byte2:
836+
; GCN: ; %bb.0:
837+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
838+
; GCN-NEXT: v_cvt_scalef32_pk_bf16_fp4 v0, v0, v1 op_sel:[1,0,0]
839+
; GCN-NEXT: s_setpc_b64 s[30:31]
840+
%ret = tail call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 %src, float %scale, i32 2)
841+
ret <2 x bfloat> %ret
842+
}
843+
844+
define <2 x bfloat> @test_cvt_scale_bf16_fp4_byte3(i32 %src, float %scale) {
845+
; GCN-LABEL: test_cvt_scale_bf16_fp4_byte3:
846+
; GCN: ; %bb.0:
847+
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
848+
; GCN-NEXT: v_cvt_scalef32_pk_bf16_fp4 v0, v0, v1 op_sel:[1,1,0]
849+
; GCN-NEXT: s_setpc_b64 s[30:31]
850+
%ret = tail call <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32 %src, float %scale, i32 3)
851+
ret <2 x bfloat> %ret
852+
}

0 commit comments

Comments
 (0)