https://github.com/arsenm updated
https://github.com/llvm/llvm-project/pull/117742
>From eb0e34ab5c455678a7e501d3e3863744a00723d8 Mon Sep 17 00:00:00 2001
From: Pravin Jagtap <pravin.jag...@amd.com>
Date: Fri, 19 Apr 2024 03:09:43 -0400
Subject: [PATCH] Builtins & Codegen support for
v_cvt_scalef32_pk_{fp|bf}8_{f|bf}16 for gfx950
OPSEL[3] determines low/high 16 bits of word to write.
Co-authored-by: Pravin Jagtap <pravin.jag...@amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 +
.../builtins-amdgcn-gfx950-err.cl | 7 +-
.../CodeGenOpenCL/builtins-amdgcn-gfx950.cl | 120 ++++++++++++++++++
.../builtins-amdgcn-error-gfx950-param.cl | 8 +-
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 15 +++
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 4 +
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 1 -
llvm/lib/Target/AMDGPU/SIInstrInfo.td | 2 +
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 22 +++-
.../llvm.amdgcn.cvt.scalef32.pk.gfx950.ll | 84 ++++++++++++
10 files changed, 257 insertions(+), 10 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index f92e5c17188ec9..0b1219d865e0d8 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -578,6 +578,10 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_f32,
"V2sV2sfffIb", "nc", "f
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_f32, "V2sV2sfffIb", "nc",
"bf8-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_fp8, "V2fUifIb", "nc",
"fp8-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_bf8, "V2fUifIb", "nc",
"bf8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_f16, "V2sV2sV2hfIb", "nc",
"fp8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_bf16, "V2sV2sV2yfIb",
"nc", "fp8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_f16, "V2sV2sV2hfIb", "nc",
"bf8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16, "V2sV2sV2yfIb",
"nc", "bf8-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 54b8573a1b4de3..9bcd6a3ea262f6 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl
@@ -15,9 +15,10 @@ typedef unsigned int uint2
__attribute__((ext_vector_type(2)));
typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2;
typedef float __attribute__((ext_vector_type(2))) float2;
+typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
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) {
+ global float2* out_v2f32, half2 src0_v2f16, bfloat2 src0_v2bf16) {
*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}}
@@ -29,4 +30,8 @@ void test(global uint* out, global uint2* out_v2u32, uint a,
uint b, global half
*out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(*out_v2i16, src0,
src1, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_bf8_f32' needs target feature
bf8-cvt-scale-insts}}
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(a, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_f32_fp8' needs target feature
fp8-cvt-scale-insts}}
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(a, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_f32_bf8' needs target feature
bf8-cvt-scale-insts}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(*out_v2i16,
src0_v2f16, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_fp8_f16' needs target feature
fp8-cvt-scale-insts}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(*out_v2i16,
src0_v2bf16, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_fp8_bf16' needs target
feature fp8-cvt-scale-insts}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(*out_v2i16,
src0_v2f16, scale, true); //
expected-error{{'__builtin_amdgcn_cvt_scalef32_pk_bf8_f16' needs target feature
bf8-cvt-scale-insts}}
+ *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}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
index 1313c5ec8d5443..2ea52ad44eaa6d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
@@ -619,3 +619,123 @@ void test_cvt_scalef32_pk_f32_bf8(global float2* out,
unsigned int src, float sc
*out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(src, scale, true);
*out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(src, scale, false);
}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk_fp8_f16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <2 x half>, align 4, 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 <2 x half> [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x i16>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load <2 x half>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.fp8.f16(<2 x i16> [[TMP1]], <2 x half> [[TMP2]],
float [[TMP3]], i1 true)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[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 i16>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load <2 x half>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.fp8.f16(<2 x i16> [[TMP7]], <2 x half> [[TMP8]],
float [[TMP9]], i1 false)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[TMP10]], ptr addrspace(1) [[TMP11]], align
4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk_fp8_f16(global short2* out, half2 src, float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(*out, src, scale, true);
+ *out = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(*out, src, scale, false);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk_fp8_bf16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <2 x bfloat>, align 4, 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 <2 x bfloat> [[SRC:%.*]], ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x i16>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load <2 x bfloat>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.fp8.bf16(<2 x i16> [[TMP1]], <2 x bfloat>
[[TMP2]], float [[TMP3]], i1 true)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[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 i16>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load <2 x bfloat>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.fp8.bf16(<2 x i16> [[TMP7]], <2 x bfloat>
[[TMP8]], float [[TMP9]], i1 false)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[TMP10]], ptr addrspace(1) [[TMP11]], align
4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk_fp8_bf16(global short2* out, bfloat2 src, float
scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(*out, src, scale, true);
+ *out = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(*out, src, scale, false);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk_bf8_f16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <2 x half>, align 4, 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 <2 x half> [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x i16>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load <2 x half>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16> [[TMP1]], <2 x half> [[TMP2]],
float [[TMP3]], i1 true)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[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 i16>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load <2 x half>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16> [[TMP7]], <2 x half> [[TMP8]],
float [[TMP9]], i1 false)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[TMP10]], ptr addrspace(1) [[TMP11]], align
4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk_bf8_f16(global short2* out, half2 src, float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(*out, src, scale, true);
+ *out = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(*out, src, scale, false);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_pk_bf8_bf16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca <2 x bfloat>, align 4, 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 <2 x bfloat> [[SRC:%.*]], ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: store float [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x i16>, ptr addrspace(1) [[TMP0]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load <2 x bfloat>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.bf8.bf16(<2 x i16> [[TMP1]], <2 x bfloat>
[[TMP2]], float [[TMP3]], i1 true)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[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 i16>, ptr addrspace(1) [[TMP6]],
align 4
+// CHECK-NEXT: [[TMP8:%.*]] = load <2 x bfloat>, ptr addrspace(5)
[[SRC_ADDR]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <2 x i16>
@llvm.amdgcn.cvt.scalef32.pk.bf8.bf16(<2 x i16> [[TMP7]], <2 x bfloat>
[[TMP8]], float [[TMP9]], i1 false)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store <2 x i16> [[TMP10]], ptr addrspace(1) [[TMP11]], align
4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_pk_bf8_bf16(global short2* out, bfloat2 src, float
scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(*out, src, scale, true);
+ *out = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(*out, src, scale, false);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
index 7aed9975b5cc22..5ff145bad1af4d 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
@@ -14,6 +14,7 @@ typedef unsigned int uint;
typedef half half2 __attribute__((ext_vector_type(2)));
typedef short short2 __attribute__((ext_vector_type(2)));
typedef float float2 __attribute__((ext_vector_type(2)));
+typedef __bf16 bfloat2 __attribute__((ext_vector_type(2)));
void test_mfma_f32_16x16x32_f16(__global float4* out, half8 a, half8 b, float4
c, int X) {
@@ -163,7 +164,8 @@ void test_permlane32_swap(__global int* out, int old, int
src, bool X) {
}
void test_cvt_scalef32(global half2* out_v2f16, global float* out_f32, uint
src, float scale, int index, bool X,
- global short2* out_v2i16, float src0, float src1,
global float2* out_v2f32) {
+ global short2* out_v2i16, float src0, float src1,
global float2* out_v2f32,
+ half2 src0_v2f16, bfloat2 src0_v2bf16) {
*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}}
*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}}
*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}}
@@ -172,4 +174,8 @@ void test_cvt_scalef32(global half2* out_v2f16, global
float* out_f32, uint src,
*out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(*out_v2i16, src0,
src1, scale, X); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_pk_bf8_f32' must be a constant integer}}
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(src, scale, X); //
expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_f32_fp8' must be
a constant integer}}
*out_v2f32 = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(src, scale, X); //
expected-error{{argument to '__builtin_amdgcn_cvt_scalef32_pk_f32_bf8' must be
a constant integer}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(*out_v2i16,
src0_v2f16, scale, X); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_pk_fp8_f16' must be a constant integer}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(*out_v2i16,
src0_v2bf16, scale, X); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_pk_fp8_bf16' must be a constant integer}}
+ *out_v2i16 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(*out_v2i16,
src0_v2f16, scale, X); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_pk_bf8_f16' must be a constant integer}}
+ *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}}
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 4e02f1d33a1fcc..5ec06ddc6e927e 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -625,6 +625,15 @@ class AMDGPUCvtScale_pk_FP8BF8ToF32Intrinsic<string name>
: DefaultAttrsIntrinsi
[IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<2>>]
>, ClangBuiltin<"__builtin_amdgcn_"#name>;
+class AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<LLVMType SrcTy, string
name> : DefaultAttrsIntrinsic<
+ [llvm_v2i16_ty],
+ [llvm_v2i16_ty, // old_vdst
+ SrcTy, // src
+ llvm_float_ty, // scale
+ llvm_i1_ty], // dst_lo_hi_sel[true false]
+ [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<3>>]
+>, ClangBuiltin<"__builtin_amdgcn_"#name>;
+
class AMDGPUCvtScaleF32ToFP8BF8TiedInputIntrinsic<string name> :
DefaultAttrsIntrinsic<
[llvm_v2i16_ty],
[llvm_v2i16_ty, // old_vdst
@@ -661,6 +670,12 @@ def int_amdgcn_cvt_scalef32_pk_bf8_f32 :
AMDGPUCvtScaleF32ToFP8BF8TiedInputIntri
def int_amdgcn_cvt_scalef32_pk_f32_fp8 :
AMDGPUCvtScale_pk_FP8BF8ToF32Intrinsic<"cvt_scalef32_pk_f32_fp8">;
def int_amdgcn_cvt_scalef32_pk_f32_bf8 :
AMDGPUCvtScale_pk_FP8BF8ToF32Intrinsic<"cvt_scalef32_pk_f32_bf8">;
+// llvm.amdgcn.cvt.scalef32.fp8.fp16 v2i16 old_vdst, v2f16 src, float scale,
bool dst_lo_hi_sel
+def int_amdgcn_cvt_scalef32_pk_fp8_f16 :
AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_pk_fp8_f16">;
+def int_amdgcn_cvt_scalef32_pk_fp8_bf16:
AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_scalef32_pk_fp8_bf16">;
+def int_amdgcn_cvt_scalef32_pk_bf8_f16 :
AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_pk_bf8_f16">;
+def int_amdgcn_cvt_scalef32_pk_bf8_bf16:
AMDGPUCvtScaleF16BF16ToFP8BF8TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_scalef32_pk_bf8_bf16">;
+
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 4c93b6ce4eef1f..277879381b5928 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4555,6 +4555,10 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_scalef32_pk_bf8_f32:
case Intrinsic::amdgcn_cvt_scalef32_pk_f32_fp8:
case Intrinsic::amdgcn_cvt_scalef32_pk_f32_bf8:
+ case Intrinsic::amdgcn_cvt_scalef32_pk_fp8_f16:
+ case Intrinsic::amdgcn_cvt_scalef32_pk_fp8_bf16:
+ case Intrinsic::amdgcn_cvt_scalef32_pk_bf8_f16:
+ case Intrinsic::amdgcn_cvt_scalef32_pk_bf8_bf16:
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/AsmParser/AMDGPUAsmParser.cpp
b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index afd35842ba87f4..f63bc0ea4e7b0c 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -8845,7 +8845,6 @@ void AMDGPUAsmParser::cvtVOP3P(MCInst &Inst, const
OperandVector &Operands,
Opc == AMDGPU::V_CVT_SR_FP8_F32_gfx12_e64_dpp8_gfx12 ||
Opc == AMDGPU::V_CVT_SR_BF8_F32_gfx12_e64_dpp_gfx12 ||
Opc == AMDGPU::V_CVT_SR_BF8_F32_gfx12_e64_dpp8_gfx12)) {
- assert(!IsPacked);
Inst.addOperand(Inst.getOperand(0));
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index a77e96d3986bac..0a91e9676b7249 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2843,6 +2843,8 @@ def VOP_V6I32_V32BF16_F32 : VOPProfile<[v6i32, v32bf16,
f32, untyped]>;
def VOP_V6I32_V16F32_V16F32_F32 : VOPProfile<[v6i32, v16f32, v16f32, f32]>;
def VOP_V2F16_I32_F32 : VOPProfile<[v2f16, i32, f32, untyped]>;
def VOP_V2I16_F32_F32_F32 : VOPProfile<[v2i16, f32, f32, f32]>;
+def VOP_V2I16_V2F16_F32 : VOPProfile<[v2i16, v2f16, f32, untyped]>;
+def VOP_V2I16_V2BF16_F32 : VOPProfile<[v2i16, v2bf16, f32, untyped]>;
def VOP_I64_I64_I32 : VOPProfile <[i64, i64, i32, untyped]>;
def VOP_I64_I32_I64 : VOPProfile <[i64, i32, i64, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index 5bf45b949a79c4..ed33cba5d3cbf3 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -997,11 +997,10 @@ class
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<ValueType DstTy> : VOP3_Pro
let HasOMod = 0;
}
-def VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_Profile : VOP3_Profile<VOPProfile<[i32,
v2f16, f32, untyped]>,
- VOP3_OPSEL> {
+class VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_TiedInput_Profile<VOPProfile P> :
VOP3_Profile<P,VOP3_OPSEL> {
let InsVOP3OpSel = (ins FP32InputMods:$src0_modifiers, Src0RC64:$src0,
FP32InputMods:$src1_modifiers, Src1RC64:$src1,
- op_sel0:$op_sel);
+ VGPR_32:$vdst_in, op_sel0:$op_sel);
let HasClamp = 0;
let HasSrc2 = 0;
let HasSrc2Mods = 0;
@@ -1027,11 +1026,11 @@ let SubtargetPredicate = HasFP8ConversionScaleInsts,
mayRaiseFPException = 0 in
let Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in {
defm V_CVT_SCALEF32_F16_FP8 : VOP3Inst<"v_cvt_scalef32_f16_fp8",
VOP3_CVT_SCALE_F1632_FP8BF8_TiedInput_Profile<VOP_V2F16_I32_F32>>;
defm V_CVT_SCALEF32_PK_FP8_F32 : VOP3Inst<"v_cvt_scalef32_pk_fp8_f32",
VOP3_CVT_SCALE_FP4FP8BF8_F32_TiedInput_Profile<VOP_V2I16_F32_F32_F32>>;
+ defm V_CVT_SCALEF32_PK_FP8_F16 : VOP3Inst<"v_cvt_scalef32_pk_fp8_f16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_TiedInput_Profile<VOP_V2I16_V2F16_F32>>;
+ defm V_CVT_SCALEF32_PK_FP8_BF16 : VOP3Inst<"v_cvt_scalef32_pk_fp8_bf16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_TiedInput_Profile<VOP_V2I16_V2BF16_F32>>;
}
defm V_CVT_SCALEF32_F32_FP8 : VOP3Inst<"v_cvt_scalef32_f32_fp8",
VOP3_CVT_SCALE_F1632_FP8BF8_Profile<f32>>;
defm V_CVT_SCALEF32_PK_F32_FP8 : VOP3Inst<"v_cvt_scalef32_pk_f32_fp8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2f32>>;
- defm V_CVT_SCALEF32_PK_FP8_F16 : VOP3Inst<"v_cvt_scalef32_pk_fp8_f16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_Profile>;
- defm V_CVT_SCALEF32_PK_FP8_BF16 : VOP3Inst<"v_cvt_scalef32_pk_fp8_bf16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_Profile>;
defm V_CVT_SCALEF32_PK_F16_FP8 : VOP3Inst<"v_cvt_scalef32_pk_f16_fp8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2f16>>;
defm V_CVT_SCALEF32_PK_BF16_FP8 : VOP3Inst<"v_cvt_scalef32_pk_bf16_fp8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2f16>>;
}
@@ -1040,11 +1039,11 @@ let SubtargetPredicate = HasBF8ConversionScaleInsts,
mayRaiseFPException = 0 in
let Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in {
defm V_CVT_SCALEF32_F16_BF8 : VOP3Inst<"v_cvt_scalef32_f16_bf8",
VOP3_CVT_SCALE_F1632_FP8BF8_TiedInput_Profile<VOP_V2F16_I32_F32>>;
defm V_CVT_SCALEF32_PK_BF8_F32 : VOP3Inst<"v_cvt_scalef32_pk_bf8_f32",
VOP3_CVT_SCALE_FP4FP8BF8_F32_TiedInput_Profile<VOP_V2I16_F32_F32_F32>>;
+ defm V_CVT_SCALEF32_PK_BF8_F16 : VOP3Inst<"v_cvt_scalef32_pk_bf8_f16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_TiedInput_Profile<VOP_V2I16_V2F16_F32>>;
+ defm V_CVT_SCALEF32_PK_BF8_BF16 : VOP3Inst<"v_cvt_scalef32_pk_bf8_bf16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_TiedInput_Profile<VOP_V2I16_V2BF16_F32>>;
}
defm V_CVT_SCALEF32_F32_BF8 : VOP3Inst<"v_cvt_scalef32_f32_bf8",
VOP3_CVT_SCALE_F1632_FP8BF8_Profile<f32>>;
defm V_CVT_SCALEF32_PK_F32_BF8 : VOP3Inst<"v_cvt_scalef32_pk_f32_bf8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2f32>>;
- defm V_CVT_SCALEF32_PK_BF8_F16 : VOP3Inst<"v_cvt_scalef32_pk_bf8_f16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_Profile>;
- defm V_CVT_SCALEF32_PK_BF8_BF16 : VOP3Inst<"v_cvt_scalef32_pk_bf8_bf16",
VOP3_CVT_SCALE_PK_FP8BF8_F16BF16_Profile>;
defm V_CVT_SCALEF32_PK_F16_BF8 : VOP3Inst<"v_cvt_scalef32_pk_f16_bf8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2bf16>>;
defm V_CVT_SCALEF32_PK_BF16_BF8 : VOP3Inst<"v_cvt_scalef32_pk_bf16_bf8",
VOP3_CVT_SCALE_PK_F16BF16F32_FP4FP8BF8_Profile<v2bf16>>;
}
@@ -1115,6 +1114,15 @@ class Cvt_Scale_PK_FP8BF8ToF32_Pat<SDPatternOperator
node, VOP3_Pseudo inst> : G
def : Cvt_Scale_PK_FP8BF8ToF32_Pat<int_amdgcn_cvt_scalef32_pk_f32_fp8,
V_CVT_SCALEF32_PK_F32_FP8_e64>;
def : Cvt_Scale_PK_FP8BF8ToF32_Pat<int_amdgcn_cvt_scalef32_pk_f32_bf8,
V_CVT_SCALEF32_PK_F32_BF8_e64>;
+class Cvt_Scale_PK_F16BF16ToFP8BF8_Pat<SDPatternOperator node, VOP3_Pseudo
inst, ValueType SrcTy> : GCNPat<
+ (v2i16 (node v2i16:$vdst_in, SrcTy:$src0, f32:$src1, timm:$word_sel)),
+ (inst (DstSelToOpSelXForm $word_sel), $src0, 0, $src1, VGPR_32:$vdst_in)
+>;
+def : Cvt_Scale_PK_F16BF16ToFP8BF8_Pat<int_amdgcn_cvt_scalef32_pk_fp8_f16,
V_CVT_SCALEF32_PK_FP8_F16_e64, v2f16>;
+def : Cvt_Scale_PK_F16BF16ToFP8BF8_Pat<int_amdgcn_cvt_scalef32_pk_fp8_bf16,
V_CVT_SCALEF32_PK_FP8_BF16_e64, v2bf16>;
+def : Cvt_Scale_PK_F16BF16ToFP8BF8_Pat<int_amdgcn_cvt_scalef32_pk_bf8_f16,
V_CVT_SCALEF32_PK_BF8_F16_e64, v2f16>;
+def : Cvt_Scale_PK_F16BF16ToFP8BF8_Pat<int_amdgcn_cvt_scalef32_pk_bf8_bf16,
V_CVT_SCALEF32_PK_BF8_BF16_e64, v2bf16>;
+
let SubtargetPredicate = isGFX10Plus in {
let isCommutable = 1, isReMaterializable = 1 in {
defm V_XOR3_B32 : VOP3Inst <"v_xor3_b32",
VOP3_Profile<VOP_I32_I32_I32_I32>>;
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 f3ad5191491b54..ce4e3206f82ce1 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
@@ -12,6 +12,10 @@ declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.f32(<2 x
i16>, float, float,
declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.f32(<2 x i16>, float,
float, float, i1)
declare <2 x float> @llvm.amdgcn.cvt.scalef32.pk.f32.fp8(i32, float, i1)
declare <2 x float> @llvm.amdgcn.cvt.scalef32.pk.f32.bf8(i32, float, i1)
+declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.f16(<2 x i16>, <2 x half>,
float, i1)
+declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.fb16(<2 x i16>, <2 x
bfloat>, float, i1)
+declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16>, <2 x half>,
float, i1)
+declare <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.fb16(<2 x i16>, <2 x
bfloat>, float, i1)
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:
@@ -560,3 +564,83 @@ define <2 x float> @test_cvt_scalef32_pk_f32_bf8_word1(i32
%src, float %scale) {
%ret = tail call <2 x float> @llvm.amdgcn.cvt.scalef32.pk.f32.bf8(i32 %src,
float %scale, i1 true)
ret <2 x float> %ret
}
+
+define <2 x i16> @test_cvt_scalef32_pk_fp8_f16_word0(<2 x i16> %old, <2 x
half> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_fp8_f16_word0:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_fp8_f16 v0, v1, v2
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.f16(<2 x i16>
%old, <2 x half> %src, float %scale, i1 false)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_fp8_f16_word1(<2 x i16> %old, <2 x
half> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_fp8_f16_word1:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_fp8_f16 v0, v1, v2 op_sel:[0,0,1]
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.f16(<2 x i16>
%old, <2 x half> %src, float %scale, i1 true)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_fp8_bf16_word0(<2 x i16> %old, <2 x
bfloat> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_fp8_bf16_word0:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_fp8_bf16 v0, v1, v2
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.bf16(<2 x i16>
%old, <2 x bfloat> %src, float %scale, i1 false)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_fp8_bf16_word1(<2 x i16> %old, <2 x
bfloat> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_fp8_bf16_word1:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_fp8_bf16 v0, v1, v2 op_sel:[0,0,1]
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.fp8.bf16(<2 x i16>
%old, <2 x bfloat> %src, float %scale, i1 true)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_bf8_f16_word0(<2 x i16> %old, <2 x
half> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_bf8_f16_word0:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_bf8_f16 v0, v1, v2
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16>
%old, <2 x half> %src, float %scale, i1 false)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_bf8_f16_word1(<2 x i16> %old, <2 x
half> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_bf8_f16_word1:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_bf8_f16 v0, v1, v2 op_sel:[0,0,1]
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.f16(<2 x i16>
%old, <2 x half> %src, float %scale, i1 true)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_bf8_bf16_word0(<2 x i16> %old, <2 x
bfloat> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_bf8_bf16_word0:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_bf8_bf16 v0, v1, v2
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.bf16(<2 x i16>
%old, <2 x bfloat> %src, float %scale, i1 false)
+ ret <2 x i16> %ret
+}
+
+define <2 x i16> @test_cvt_scalef32_pk_bf8_bf16_word1(<2 x i16> %old, <2 x
bfloat> %src, float %scale) {
+; GCN-LABEL: test_cvt_scalef32_pk_bf8_bf16_word1:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT: v_cvt_scalef32_pk_bf8_bf16 v0, v1, v2 op_sel:[0,0,1]
+; GCN-NEXT: s_setpc_b64 s[30:31]
+ %ret = tail call <2 x i16> @llvm.amdgcn.cvt.scalef32.pk.bf8.bf16(<2 x i16>
%old, <2 x bfloat> %src, float %scale, i1 true)
+ ret <2 x i16> %ret
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
