https://github.com/arsenm created
https://github.com/llvm/llvm-project/pull/117824
Co-authored-by: Shilei Tian <shilei.t...@amd.com>
>From e511ff9110c445d9b2971bbe43c946d75afaaeff Mon Sep 17 00:00:00 2001
From: Shilei Tian <shilei.t...@amd.com>
Date: Wed, 5 Jun 2024 14:16:20 -0400
Subject: [PATCH] AMDGPU: Builtin & CodeGen support for v_cvt_sr_{bf16|f16}_f32
instructions
Co-authored-by: Shilei Tian <shilei.t...@amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 3 +
clang/test/CodeGenOpenCL/amdgpu-features.cl | 2 +-
.../builtins-amdgcn-gfx950-err.cl | 4 +-
.../CodeGenOpenCL/builtins-amdgcn-gfx950.cl | 66 ++++++++-
.../builtins-amdgcn-error-gfx950-param.cl | 2 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 12 ++
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 2 +
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 7 +
llvm/lib/TargetParser/TargetParser.cpp | 1 +
.../test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sr.ll | 126 ++++++++++++++++++
10 files changed, 221 insertions(+), 4 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sr.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 7ed488d3a7e515..752a751a25a0c7 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -617,5 +617,8 @@
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32, "V6UiV32fUif", "nc
TARGET_BUILTIN(__builtin_amdgcn_bitop3_b32, "iiiiIUc", "nc", "bitop3-insts")
TARGET_BUILTIN(__builtin_amdgcn_bitop3_b16, "ssssIUc", "nc", "bitop3-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf16_f32, "V2yV2yfUiIb", "nc",
"f32-to-f16bf16-cvt-sr-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_f16_f32, "V2hV2hfUiIb", "nc",
"f32-to-f16bf16-cvt-sr-insts")
+
#undef BUILTIN
#undef TARGET_BUILTIN
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl
b/clang/test/CodeGenOpenCL/amdgpu-features.cl
index f629262ac58879..633f1dec5e3705 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl
@@ -89,7 +89,7 @@
// GFX941:
"target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
// GFX942:
"target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
// GFX9_4_Generic:
"target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
-// GFX950:
"target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
+// GFX950:
"target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
// GFX1010:
"target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
// GFX1011:
"target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
// GFX1012:
"target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
index 81195e721c13f6..521121f5e7e543 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
@@ -22,7 +22,7 @@ typedef unsigned int __attribute__((ext_vector_type(6)))
uint6;
typedef half __attribute__((ext_vector_type(32))) half32;
typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
-void test(global uint* out, global uint2* out_v2u32, uint a, uint b, uint c,
global half2* out_v2f16, global float* out_f32, float scale, global short2*
out_v2i16, float src0, float src1,
+void test(global uint* out, global uint2* out_v2u32, uint a, uint b, uint c,
global half2* out_v2f16, global float* out_f32, float scale, unsigned seed,
global short2* out_v2i16, float src0, float src1,
float2 src0_v2f32, global float2* out_v2f32, half2 src0_v2f16,
bfloat2 src0_v2bf16, global bfloat2* out_v2bf16, global float32* out_v32f32,
uint6 src_v6i32,
global half32 *out_v32f16, global bfloat32 *out_v32bf16) {
*out = __builtin_amdgcn_prng_b32(a); //
expected-error{{'__builtin_amdgcn_prng_b32' needs target feature prng-inst}}
@@ -61,4 +61,6 @@ void test(global uint* out, global uint2* out_v2u32, uint a,
uint b, uint c, glo
*out = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(*out, src0_v2f32, 0,
scale, 0); // expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32'
needs target feature fp4-cvt-scale-insts}}
*out = __builtin_amdgcn_bitop3_b32(a, b, c, 1); // expected-error
{{'__builtin_amdgcn_bitop3_b32' needs target feature bitop3-insts}}
*out = __builtin_amdgcn_bitop3_b16((ushort)a, (ushort)b, (ushort)c, 1); //
expected-error {{'__builtin_amdgcn_bitop3_b16' needs target feature
bitop3-insts}}
+ *out_v2bf16 = __builtin_amdgcn_cvt_sr_bf16_f32(*out_v2bf16, src0, seed, 0);
// expected-error{{'__builtin_amdgcn_cvt_sr_bf16_f32' needs target feature
f32-to-f16bf16-cvt-sr-insts}}
+ *out_v2f16 = __builtin_amdgcn_cvt_sr_f16_f32(*out_v2f16, src0, seed, 0); //
expected-error{{'__builtin_amdgcn_cvt_sr_f16_f32' needs target feature
f32-to-f16bf16-cvt-sr-insts}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
index 2e1025e895c636..1eebece1ec22b9 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
@@ -1678,7 +1678,8 @@ void test_cvt_scalef32_sr_fp8_f32(global unsigned *out,
float src, uint seed, fl
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
// CHECK-NEXT: ret void
//
-void test_bitop3_b32(global uint* out, uint a, uint b, uint c) {
+void test_bitop3_b32(global uint* out, uint a, uint b, uint c)
+{
*out = __builtin_amdgcn_bitop3_b32(a, b, c, 1);
}
@@ -1700,6 +1701,67 @@ void test_bitop3_b32(global uint* out, uint a, uint b,
uint c) {
// CHECK-NEXT: store i16 [[TMP3]], ptr addrspace(1) [[TMP4]], align 2
// CHECK-NEXT: ret void
//
-void test_bitop3_b16(global ushort* out, ushort a, ushort b, ushort c) {
+void test_bitop3_b16(global ushort* out, ushort a, ushort b, ushort c)
+{
*out = __builtin_amdgcn_bitop3_b16(a, b, c, 1);
}
+
+// CHECK-LABEL: @test_cvt_sr_bf16_f32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x bfloat>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x bfloat>
@llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat> [[TMP1]], float [[TMP2]], i32
[[TMP3]], i1 false)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x bfloat> [[TMP4]], ptr addrspace(1) [[TMP5]],
align 4
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP7:%.*]] = load <2 x bfloat>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x bfloat>
@llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat> [[TMP7]], float [[TMP8]], i32
[[TMP9]], i1 true)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x bfloat> [[TMP10]], ptr addrspace(1) [[TMP11]],
align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_sr_bf16_f32(global bfloat2 *out, float src, uint seed)
+{
+ *out = __builtin_amdgcn_cvt_sr_bf16_f32(*out, src, seed, 0);
+ *out = __builtin_amdgcn_cvt_sr_bf16_f32(*out, src, seed, 1);
+}
+
+// CHECK-LABEL: @test_cvt_sr_f16_f32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x half>
@llvm.amdgcn.cvt.sr.f16.f32(<2 x half> [[TMP1]], float [[TMP2]], i32 [[TMP3]],
i1 false)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x half> [[TMP4]], ptr addrspace(1) [[TMP5]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP7:%.*]] = load <2 x half>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x half>
@llvm.amdgcn.cvt.sr.f16.f32(<2 x half> [[TMP7]], float [[TMP8]], i32 [[TMP9]],
i1 true)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x half> [[TMP10]], ptr addrspace(1) [[TMP11]],
align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_sr_f16_f32(global half2 *out, float src, uint seed)
+{
+ *out = __builtin_amdgcn_cvt_sr_f16_f32(*out, src, seed, 0);
+ *out = __builtin_amdgcn_cvt_sr_f16_f32(*out, src, seed, 1);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
index c0ea2f961337e1..ce4f6ba8f407fd 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
@@ -198,6 +198,8 @@ void test_cvt_scalef32(global half2* out_v2f16, global
float* out_f32, uint src,
*out = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_fp8_bf16' must be a constant integer}}
*out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_fp8_f16' must be a constant integer}}
*out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_fp8_f32' must be a constant integer}}
+ *out_v2bf16 = __builtin_amdgcn_cvt_sr_bf16_f32(*out_v2bf16, src0, src, X);
// expected-error{{argument to '__builtin_amdgcn_cvt_sr_bf16_f32' must be a
constant integer}}
+ *out_v2f16 = __builtin_amdgcn_cvt_sr_f16_f32(*out_v2f16, src0, src, X); //
expected-error{{argument to '__builtin_amdgcn_cvt_sr_f16_f32' must be a
constant integer}}
}
void test_bitop3_args(global uint* out, uint a, uint b, uint c) {
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index eb3167506faec9..112c26d20db14e 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -694,6 +694,18 @@ class
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<LLVMType Src0Ty, st
[IntrNoMem, IntrWillReturn, ImmArg<ArgIndex<4>>]
>, ClangBuiltin<"__builtin_amdgcn_"#name>;
+class AMDGPUCvtScaleSRF32ToBF16F16TiedInputIntrinsic<LLVMType DstTy, string
name> : DefaultAttrsIntrinsic<
+ [DstTy],
+ [DstTy, // old_vdst
+ llvm_float_ty, // src0
+ llvm_i32_ty, // seed
+ llvm_i1_ty], // dst_lo_hi_sel[true false]
+ [IntrNoMem, IntrWillReturn, ImmArg<ArgIndex<3>>]
+>, ClangBuiltin<"__builtin_amdgcn_"#name>;
+
+def int_amdgcn_cvt_sr_bf16_f32:
AMDGPUCvtScaleSRF32ToBF16F16TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_sr_bf16_f32">;
+def int_amdgcn_cvt_sr_f16_f32 :
AMDGPUCvtScaleSRF32ToBF16F16TiedInputIntrinsic<llvm_v2f16_ty, "cvt_sr_f16_f32">;
+
// llvm.amdgcn.cvt.scalef32.fp16.fp8 v2f16 old_vdst, int src, float scale, int
src_sel_index [0..3], bool dst_lo_hi_sel
def int_amdgcn_cvt_scalef32_f16_fp8 :
AMDGPUCvtScaleFP8BF8ToF16TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_f16_fp8">;
def int_amdgcn_cvt_scalef32_f16_bf8 :
AMDGPUCvtScaleFP8BF8ToF16TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_f16_bf8">;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index ff994aef3951de..e5baffc0f064b2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4543,6 +4543,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_pk_bf8_f32:
case Intrinsic::amdgcn_cvt_sr_fp8_f32:
case Intrinsic::amdgcn_cvt_sr_bf8_f32:
+ case Intrinsic::amdgcn_cvt_sr_bf16_f32:
+ case Intrinsic::amdgcn_cvt_sr_f16_f32:
case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_f16:
case Intrinsic::amdgcn_cvt_scalef32_pk32_bf6_f16:
case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_bf16:
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index 5a78632a3de055..34850e42a3d605 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -1343,11 +1343,18 @@ let SubtargetPredicate = HasBF16ConversionInsts in {
(V_CVT_PK_BF16_F32_e64 0, (f32 (V_CVT_F32_F64_e64
$src0_modifiers, $src0)), 0, (f32 (IMPLICIT_DEF)))>;
}
+class Cvt_Scale_Sr_F32ToBF16F16_Pat<SDPatternOperator node, VOP3_Pseudo inst,
ValueType DstTy> : GCNPat<
+ (DstTy (node DstTy:$vdst_in, f32:$src0, i32:$src1, timm:$word_sel)),
+ (inst (DstSelToOpSelXForm $word_sel), $src0, 0, $src1, VGPR_32:$vdst_in)
+>;
+
let SubtargetPredicate = HasF32ToF16BF16ConversionSRInsts in {
let Constraints = "$vdst = $vdst_in", DisableEncoding = "$vdst_in" in {
defm V_CVT_SR_F16_F32 : VOP3Inst<"v_cvt_sr_f16_f32",
VOP3_CVT_SR_FP16_TiedInput_Profile<VOP_F16_F32_I32>>;
defm V_CVT_SR_BF16_F32 : VOP3Inst<"v_cvt_sr_bf16_f32",
VOP3_CVT_SR_FP16_TiedInput_Profile<VOP_BF16_F32_I32>>;
}
+ def : Cvt_Scale_Sr_F32ToBF16F16_Pat<int_amdgcn_cvt_sr_bf16_f32,
V_CVT_SR_BF16_F32_e64, v2bf16>;
+ def : Cvt_Scale_Sr_F32ToBF16F16_Pat<int_amdgcn_cvt_sr_f16_f32,
V_CVT_SR_F16_F32_e64, v2f16>;
}
let SubtargetPredicate = isGFX12Plus, ReadsModeReg = 0 in {
diff --git a/llvm/lib/TargetParser/TargetParser.cpp
b/llvm/lib/TargetParser/TargetParser.cpp
index ece2df0b13dd33..02295fdb0ecd0d 100644
--- a/llvm/lib/TargetParser/TargetParser.cpp
+++ b/llvm/lib/TargetParser/TargetParser.cpp
@@ -479,6 +479,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const
Triple &T,
Features["bf8-cvt-scale-insts"] = true;
Features["fp8-cvt-scale-insts"] = true;
Features["f16bf16-to-fp6bf6-cvt-scale-insts"] = true;
+ Features["f32-to-f16bf16-cvt-sr-insts"] = true;
Features["prng-inst"] = true;
Features["permlane16-swap"] = true;
Features["permlane32-swap"] = true;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sr.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sr.ll
new file mode 100644
index 00000000000000..4a58d6346fc573
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sr.ll
@@ -0,0 +1,126 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
UTC_ARGS: --version 4
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -o - %s | FileCheck
-check-prefix=GFX950 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx950 -o - %s | FileCheck
-check-prefix=GFX950 %s
+
+declare <2 x half> @llvm.amdgcn.cvt.sr.f16.f32(<2 x half>, float, i32, i1)
+declare <2 x bfloat> @llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat>, float, i32, i1)
+
+define amdgpu_ps void @test_cvt_sr_bf16_f32_word_sel_0(ptr addrspace(1) %out,
float %src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_bf16_f32_word_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_bf16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x bfloat>, ptr addrspace(1) %out, align 4
+ %cvt = tail call <2 x bfloat> @llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat>
%old, float %src0, i32 %src1, i1 false)
+ store <2 x bfloat> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_bf16_f32_word_sel_1(ptr addrspace(1) %out,
float %src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_bf16_f32_word_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_bf16_f32 v4, v2, v3 op_sel:[0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x bfloat>, ptr addrspace(1) %out, align 4
+ %cvt = tail call <2 x bfloat> @llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat>
%old, float %src0, i32 %src1, i1 true)
+ store <2 x bfloat> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_bf16_f32_fabs(ptr addrspace(1) %out, float
%src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_bf16_f32_fabs:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: v_and_b32_e32 v2, 0x7fffffff, v2
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_bf16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x bfloat>, ptr addrspace(1) %out, align 4
+ %src0.fabs = call float @llvm.fabs.f32(float %src0)
+ %cvt = tail call <2 x bfloat> @llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat>
%old, float %src0.fabs, i32 %src1, i1 false)
+ store <2 x bfloat> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_bf16_f32_fneg(ptr addrspace(1) %out, float
%src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_bf16_f32_fneg:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: v_xor_b32_e32 v2, 0x80000000, v2
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_bf16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x bfloat>, ptr addrspace(1) %out, align 4
+ %src0.fneg = fneg float %src0
+ %cvt = tail call <2 x bfloat> @llvm.amdgcn.cvt.sr.bf16.f32(<2 x bfloat>
%old, float %src0.fneg, i32 %src1, i1 false)
+ store <2 x bfloat> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_f16_f32_word_sel_0(ptr addrspace(1) %out,
float %src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_f16_f32_word_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_f16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x half>, ptr addrspace(1) %out, align 4
+ %cvt = tail call <2 x half> @llvm.amdgcn.cvt.sr.f16.f32(<2 x half> %old,
float %src0, i32 %src1, i1 false)
+ store <2 x half> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_f16_f32_word_sel_1(ptr addrspace(1) %out,
float %src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_f16_f32_word_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_f16_f32 v4, v2, v3 op_sel:[0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x half>, ptr addrspace(1) %out, align 4
+ %cvt = tail call <2 x half> @llvm.amdgcn.cvt.sr.f16.f32(<2 x half> %old,
float %src0, i32 %src1, i1 true)
+ store <2 x half> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_f16_f32_fabs(ptr addrspace(1) %out, float
%src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_f16_f32_fabs:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: v_and_b32_e32 v2, 0x7fffffff, v2
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_f16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x half>, ptr addrspace(1) %out, align 4
+ %src0.fabs = call float @llvm.fabs.f32(float %src0)
+ %cvt = tail call <2 x half> @llvm.amdgcn.cvt.sr.f16.f32(<2 x half> %old,
float %src0.fabs, i32 %src1, i1 false)
+ store <2 x half> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_sr_f16_f32_fneg(ptr addrspace(1) %out, float
%src0, i32 %src1) {
+; GFX950-LABEL: test_cvt_sr_f16_f32_fneg:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v4, v[0:1], off
+; GFX950-NEXT: v_xor_b32_e32 v2, 0x80000000, v2
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_sr_f16_f32 v4, v2, v3
+; GFX950-NEXT: global_store_dword v[0:1], v4, off
+; GFX950-NEXT: s_endpgm
+ %old = load <2 x half>, ptr addrspace(1) %out, align 4
+ %src0.fneg = fneg float %src0
+ %cvt = tail call <2 x half> @llvm.amdgcn.cvt.sr.f16.f32(<2 x half> %old,
float %src0.fneg, i32 %src1, i1 false)
+ store <2 x half> %cvt, ptr addrspace(1) %out, align 8
+ ret void
+}
_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
