https://github.com/arsenm updated
https://github.com/llvm/llvm-project/pull/117747
>From dfbc9e972c079db142962f6c2ba0bcfc3a705e52 Mon Sep 17 00:00:00 2001
From: Pravin Jagtap <pravin.jag...@amd.com>
Date: Mon, 22 Apr 2024 05:29:12 -0400
Subject: [PATCH] AMDGPU: Builtin & codegen support for
v_cvt_scalef32_pk32_{bf|f}16_{bf|fp}6 for gfx950
Co-authored-by: Pravin Jagtap <pravin.jag...@amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 +
.../builtins-amdgcn-gfx950-err.cl | 13 +-
.../CodeGenOpenCL/builtins-amdgcn-gfx950.cl | 52 ++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 +
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 4 +
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 8 +-
.../llvm.amdgcn.cvt.scalef32.pk.gfx950.ll | 236 ++++++++++++++++++
7 files changed, 316 insertions(+), 7 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 073ae9cd1ab89f..943466e37194ae 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -588,6 +588,10 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f16_fp4,
"V2hUifIi", "nc", "fp4-
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4, "V2yUifIi", "nc",
"fp4-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f32_fp6, "V32fV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f32_bf6, "V32fV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f16_fp6, "V32hV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6, "V32yV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f16_bf6, "V32hV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6, "V32yV6Uif", "nc",
"fp6bf6-cvt-scale-insts")
#undef BUILTIN
#undef TARGET_BUILTIN
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
index 678ef633ee6fee..58d7ec7d2fc714 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
@@ -18,9 +18,12 @@ typedef float __attribute__((ext_vector_type(2))) float2;
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
typedef float __attribute__((ext_vector_type(32))) float32;
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, global
half2* out_v2f16, global float* out_f32, float scale, global short2* out_v2i16,
float src0, float src1,
- global float2* out_v2f32, half2 src0_v2f16, bfloat2 src0_v2bf16,
global bfloat2* out_v2bf16, global float32* out_v36f32, uint6 src_v6i32) {
+ 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}}
*out_v2u32 = __builtin_amdgcn_permlane16_swap(a, b, false, false); //
expected-error{{'__builtin_amdgcn_permlane16_swap' needs target feature
permlane16-swap}}
*out_v2u32 = __builtin_amdgcn_permlane32_swap(a, b, false, false); //
expected-error{{'__builtin_amdgcn_permlane32_swap' needs target feature
permlane32-swap}}
@@ -40,6 +43,10 @@ void test(global uint* out, global uint2* out_v2u32, uint a,
uint b, global half
*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}}
*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}}
*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}}
- *out_v36f32 = __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f32_fp6' needs target
feature fp6bf6-cvt-scale-insts}}
- *out_v36f32 = __builtin_amdgcn_cvt_scalef32_pk32_f32_bf6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f32_bf6' needs target
feature fp6bf6-cvt-scale-insts}}
+ *out_v32f32 = __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f32_fp6' needs target
feature fp6bf6-cvt-scale-insts}}
+ *out_v32f32 = __builtin_amdgcn_cvt_scalef32_pk32_f32_bf6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f32_bf6' needs target
feature fp6bf6-cvt-scale-insts}}
+ *out_v32f16 = __builtin_amdgcn_cvt_scalef32_pk32_f16_fp6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f16_fp6' needs target
feature fp6bf6-cvt-scale-insts}}
+ *out_v32f16 = __builtin_amdgcn_cvt_scalef32_pk32_f16_bf6(src_v6i32, scale);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_f16_bf6' needs target
feature fp6bf6-cvt-scale-insts}}
+ *out_v32bf16 = __builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6(src_v6i32,
scale); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6' needs
target feature fp6bf6-cvt-scale-insts}}
+ *out_v32bf16 = __builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6(src_v6i32,
scale); // expected-error{{'__builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6' needs
target feature fp6bf6-cvt-scale-insts}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
index 2bb0f3357e8d21..ab2b810bdb26e5 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
@@ -932,3 +932,55 @@ void test_cvt_scalef32_pk_f32_fp6(global float32* out,
uint6 src, float scale)
*out = __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(src, scale);
*out = __builtin_amdgcn_cvt_scalef32_pk32_f32_bf6(src, scale);
}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk32_f16_fpbf6(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <6 x i32>, align 32, addrspace(5)
+// CHECK-NEXT: [[SCALE_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <6 x i32> [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 32
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <6 x i32>, ptr addrspace(5)
[[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call <32 x half>
@llvm.amdgcn.cvt.scalef32.pk32.f16.fp6(<6 x i32> [[TMP0]], float [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <32 x half> [[TMP2]], ptr addrspace(1) [[TMP3]], align
64
+// CHECK-NEXT: [[TMP4:%.*]] = load <6 x i32>, ptr addrspace(5)
[[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP6:%.*]] = call <32 x half>
@llvm.amdgcn.cvt.scalef32.pk32.f16.bf6(<6 x i32> [[TMP4]], float [[TMP5]])
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <32 x half> [[TMP6]], ptr addrspace(1) [[TMP7]], align
64
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk32_f16_fpbf6(global half32 *out, uint6 src, float
scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk32_f16_fp6(src, scale);
+ *out = __builtin_amdgcn_cvt_scalef32_pk32_f16_bf6(src, scale);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk32_bf16_fpbf6(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <6 x i32>, align 32, addrspace(5)
+// CHECK-NEXT: [[SCALE_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <6 x i32> [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 32
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <6 x i32>, ptr addrspace(5)
[[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call <32 x bfloat>
@llvm.amdgcn.cvt.scalef32.pk32.bf16.fp6(<6 x i32> [[TMP0]], float [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <32 x bfloat> [[TMP2]], ptr addrspace(1) [[TMP3]],
align 64
+// CHECK-NEXT: [[TMP4:%.*]] = load <6 x i32>, ptr addrspace(5)
[[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP6:%.*]] = call <32 x bfloat>
@llvm.amdgcn.cvt.scalef32.pk32.bf16.bf6(<6 x i32> [[TMP4]], float [[TMP5]])
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <32 x bfloat> [[TMP6]], ptr addrspace(1) [[TMP7]],
align 64
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk32_bf16_fpbf6(global bfloat32 *out, uint6 src, float
scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6(src, scale);
+ *out = __builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6(src, scale);
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 47a55abf318e54..7db1b2a4b620de 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -700,6 +700,12 @@ def int_amdgcn_cvt_scalef32_pk_bf16_fp4:
AMDGPUCvtScaleFP4FP8BF8ToF1632Intrinsic
def int_amdgcn_cvt_scalef32_pk32_f32_fp6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32f32_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_f32_fp6">;
def int_amdgcn_cvt_scalef32_pk32_f32_bf6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32f32_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_f32_bf6">;
+// llvm.amdgcn.cvt.scalef32.pk32.f16.fp6 v6i32 src, float scale
+def int_amdgcn_cvt_scalef32_pk32_f16_bf6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32f16_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_f16_bf6">;
+def int_amdgcn_cvt_scalef32_pk32_bf16_bf6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32bf16_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_bf16_bf6">;
+def int_amdgcn_cvt_scalef32_pk32_f16_fp6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32f16_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_f16_fp6">;
+def int_amdgcn_cvt_scalef32_pk32_bf16_fp6 :
AMDGPUCvtScaleF32Intrinsic<llvm_v32bf16_ty, llvm_v6i32_ty,
"cvt_scalef32_pk32_bf16_fp6">;
+
def int_amdgcn_prng_b32 : DefaultAttrsIntrinsic<
[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]
>, ClangBuiltin<"__builtin_amdgcn_prng_b32">;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index dc434c390c2227..ff3bc605dab2cf 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4565,6 +4565,10 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_scalef32_pk_bf16_fp4:
case Intrinsic::amdgcn_cvt_scalef32_pk32_f32_fp6:
case Intrinsic::amdgcn_cvt_scalef32_pk32_f32_bf6:
+ case Intrinsic::amdgcn_cvt_scalef32_pk32_f16_bf6:
+ case Intrinsic::amdgcn_cvt_scalef32_pk32_bf16_bf6:
+ case Intrinsic::amdgcn_cvt_scalef32_pk32_f16_fp6:
+ case Intrinsic::amdgcn_cvt_scalef32_pk32_bf16_fp6:
case Intrinsic::amdgcn_ashr_pk_i8_i32:
case Intrinsic::amdgcn_ashr_pk_u8_i32:
case Intrinsic::amdgcn_cvt_scalef32_2xpk16_fp6_f32:
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index f52baa9a596e9b..61578af892d4b1 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -1066,10 +1066,10 @@ let SubtargetPredicate = HasFP4ConversionScaleInsts,
mayRaiseFPException = 0 in
let SubtargetPredicate = HasFP6BF6ConversionScaleInsts, mayRaiseFPException =
0 in {
defm V_CVT_SCALEF32_PK32_F32_FP6 : VOP3Inst<"v_cvt_scalef32_pk32_f32_fp6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F32_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_f32_fp6>;
defm V_CVT_SCALEF32_PK32_F32_BF6 : VOP3Inst<"v_cvt_scalef32_pk32_f32_bf6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F32_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_f32_bf6>;
- defm V_CVT_SCALEF32_PK32_F16_FP6 : VOP3Inst<"v_cvt_scalef32_pk32_f16_fp6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F16_V6I32_F32>>;
- defm V_CVT_SCALEF32_PK32_BF16_FP6 : VOP3Inst<"v_cvt_scalef32_pk32_bf16_fp6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32BF16_V6I32_F32>>;
- defm V_CVT_SCALEF32_PK32_F16_BF6 : VOP3Inst<"v_cvt_scalef32_pk32_f16_bf6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F16_V6I32_F32>>;
- defm V_CVT_SCALEF32_PK32_BF16_BF6 : VOP3Inst<"v_cvt_scalef32_pk32_bf16_bf6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32BF16_V6I32_F32>>;
+ defm V_CVT_SCALEF32_PK32_F16_FP6 : VOP3Inst<"v_cvt_scalef32_pk32_f16_fp6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F16_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_f16_fp6>;
+ defm V_CVT_SCALEF32_PK32_BF16_FP6 : VOP3Inst<"v_cvt_scalef32_pk32_bf16_fp6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32BF16_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_bf16_fp6>;
+ defm V_CVT_SCALEF32_PK32_F16_BF6 : VOP3Inst<"v_cvt_scalef32_pk32_f16_bf6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32F16_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_f16_bf6>;
+ defm V_CVT_SCALEF32_PK32_BF16_BF6 : VOP3Inst<"v_cvt_scalef32_pk32_bf16_bf6",
VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V32BF16_V6I32_F32>,
int_amdgcn_cvt_scalef32_pk32_bf16_bf6>;
}
let SubtargetPredicate = HasF16BF16ToFP6BF6ConversionScaleInsts,
mayRaiseFPException = 0 in {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx950.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx950.ll
index 7d78d0b739b9f9..25f6c454640bab 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx950.ll
@@ -22,6 +22,10 @@ declare <2 x half> @llvm.amdgcn.cvt.scalef32.pk.f16.fp4(i32,
float, i32)
declare <2 x bfloat> @llvm.amdgcn.cvt.scalef32.pk.bf16.fp4(i32, float, i32)
declare <32 x float> @llvm.amdgcn.cvt.scalef32.pk32.f32.fp6(<6 x i32>, float)
declare <32 x float> @llvm.amdgcn.cvt.scalef32.pk32.f32.bf6(<6 x i32>, float)
+declare <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.fp6(<6 x i32>, float)
+declare <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.fp6(<6 x i32>, float)
+declare <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.bf6(<6 x i32>, float)
+declare <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.bf6(<6 x i32>, float)
define amdgpu_ps void @test_scalef32_pk32_fp6_f32_vv(<16 x float> %src, float
%scale, ptr addrspace(1) %out) {
; GFX950-SDAG-LABEL: test_scalef32_pk32_fp6_f32_vv:
@@ -872,3 +876,235 @@ define <32 x float> @test_cvt_scale_pk32_f32_bf6(<6 x
i32> %src, float %scale) {
%ret = tail call <32 x float> @llvm.amdgcn.cvt.scalef32.pk32.f32.bf6(<6 x
i32> %src, float %scale)
ret <32 x float> %ret
}
+
+define <32 x half> @test_cvt_scalef32_pk32_f16_fp6_vv(<6 x i32> %src, float
%scale) {
+; GCN-LABEL: test_cvt_scalef32_pk32_f16_fp6_vv:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk32_f16_fp6 v[0:15], v[0:5], v6
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.fp6(<6 x
i32> %src, float %scale)
+ ret <32 x half> %ret
+}
+
+define <32 x half> @test_cvt_scalef32_pk32_f16_fp6_sl(<6 x i32> inreg %src) {
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_f16_fp6_sl:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v2, s2
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v3, s3
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v4, s16
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v5, s17
+; GFX950-SDAG-NEXT: s_mov_b32 s0, 0x42c80000
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_f16_fp6 v[0:15], v[0:5], s0
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_f16_fp6_sl:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: s_mov_b32 s4, s16
+; GFX950-GISEL-NEXT: s_mov_b32 s5, s17
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, 0x42c80000
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_f16_fp6 v[0:15], v[0:5], v6
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.fp6(<6 x
i32> %src, float 100.0)
+ ret <32 x half> %ret
+}
+
+define <32 x bfloat> @test_cvt_scalef32_pk32_bf16_fp6_vv(<6 x i32> %src, float
%scale) {
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_bf16_fp6_vv:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_bf16_fp6 v[0:15], v[0:5], v6
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_bf16_fp6_vv:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_bf16_fp6 v[16:31], v[0:5], v6
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v1, 16, v16
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v3, 16, v17
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v5, 16, v18
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v7, 16, v19
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v9, 16, v20
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v11, 16, v21
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v13, 16, v22
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v15, 16, v23
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v0, v16
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v2, v17
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v4, v18
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, v19
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v8, v20
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v10, v21
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v12, v22
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v14, v23
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.fp6(<6 x
i32> %src, float %scale)
+ ret <32 x bfloat> %ret
+}
+
+define <32 x bfloat> @test_cvt_scalef32_pk32_bf16_fp6_sl(<6 x i32> inreg %src)
{
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_bf16_fp6_sl:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v2, s2
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v3, s3
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v4, s16
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v5, s17
+; GFX950-SDAG-NEXT: s_mov_b32 s0, 0x42c80000
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_bf16_fp6 v[0:15], v[0:5], s0
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_bf16_fp6_sl:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: s_mov_b32 s4, s16
+; GFX950-GISEL-NEXT: s_mov_b32 s5, s17
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, 0x42c80000
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_bf16_fp6 v[16:31], v[0:5], v6
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v1, 16, v16
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v3, 16, v17
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v5, 16, v18
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v7, 16, v19
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v9, 16, v20
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v11, 16, v21
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v13, 16, v22
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v15, 16, v23
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v0, v16
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v2, v17
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v4, v18
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, v19
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v8, v20
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v10, v21
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v12, v22
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v14, v23
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.fp6(<6 x
i32> %src, float 100.0)
+ ret <32 x bfloat> %ret
+}
+
+define <32 x half> @test_cvt_scalef32_pk32_f16_bf6_vv(<6 x i32> %src, float
%scale) {
+; GCN-LABEL: test_cvt_scalef32_pk32_f16_bf6_vv:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk32_f16_bf6 v[0:15], v[0:5], v6
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.bf6(<6 x
i32> %src, float %scale)
+ ret <32 x half> %ret
+}
+
+define <32 x half> @test_cvt_scalef32_pk32_f16_bf6_sl(<6 x i32> inreg %src) {
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_f16_bf6_sl:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v2, s2
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v3, s3
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v4, s16
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v5, s17
+; GFX950-SDAG-NEXT: s_mov_b32 s0, 0x42c80000
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_f16_bf6 v[0:15], v[0:5], s0
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_f16_bf6_sl:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: s_mov_b32 s4, s16
+; GFX950-GISEL-NEXT: s_mov_b32 s5, s17
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, 0x42c80000
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_f16_bf6 v[0:15], v[0:5], v6
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x half> @llvm.amdgcn.cvt.scalef32.pk32.f16.bf6(<6 x
i32> %src, float 100.0)
+ ret <32 x half> %ret
+}
+
+define <32 x bfloat> @test_cvt_scalef32_pk32_bf16_bf6_vv(<6 x i32> %src, float
%scale) {
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_bf16_bf6_vv:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_bf16_bf6 v[0:15], v[0:5], v6
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_bf16_bf6_vv:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_bf16_bf6 v[16:31], v[0:5], v6
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v1, 16, v16
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v3, 16, v17
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v5, 16, v18
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v7, 16, v19
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v9, 16, v20
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v11, 16, v21
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v13, 16, v22
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v15, 16, v23
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v0, v16
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v2, v17
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v4, v18
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, v19
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v8, v20
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v10, v21
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v12, v22
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v14, v23
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.bf6(<6 x
i32> %src, float %scale)
+ ret <32 x bfloat> %ret
+}
+
+define <32 x bfloat> @test_cvt_scalef32_pk32_bf16_bf6_sl(<6 x i32> inreg %src)
{
+; GFX950-SDAG-LABEL: test_cvt_scalef32_pk32_bf16_bf6_sl:
+; GFX950-SDAG: ; %bb.0:
+; GFX950-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v2, s2
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v3, s3
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v4, s16
+; GFX950-SDAG-NEXT: v_mov_b32_e32 v5, s17
+; GFX950-SDAG-NEXT: s_mov_b32 s0, 0x42c80000
+; GFX950-SDAG-NEXT: v_cvt_scalef32_pk32_bf16_bf6 v[0:15], v[0:5], s0
+; GFX950-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX950-GISEL-LABEL: test_cvt_scalef32_pk32_bf16_bf6_sl:
+; GFX950-GISEL: ; %bb.0:
+; GFX950-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX950-GISEL-NEXT: s_mov_b32 s4, s16
+; GFX950-GISEL-NEXT: s_mov_b32 s5, s17
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX950-GISEL-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, 0x42c80000
+; GFX950-GISEL-NEXT: v_cvt_scalef32_pk32_bf16_bf6 v[16:31], v[0:5], v6
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v1, 16, v16
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v3, 16, v17
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v5, 16, v18
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v7, 16, v19
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v9, 16, v20
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v11, 16, v21
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v13, 16, v22
+; GFX950-GISEL-NEXT: v_lshrrev_b32_e32 v15, 16, v23
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v0, v16
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v2, v17
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v4, v18
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v6, v19
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v8, v20
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v10, v21
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v12, v22
+; GFX950-GISEL-NEXT: v_mov_b32_e32 v14, v23
+; GFX950-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <32 x bfloat> @llvm.amdgcn.cvt.scalef32.pk32.bf16.bf6(<6 x
i32> %src, float 100.0)
+ ret <32 x bfloat> %ret
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
