https://github.com/arsenm updated
https://github.com/llvm/llvm-project/pull/117821
>From b26cbd2043132438fea529b8ee3eb91ec8f9c064 Mon Sep 17 00:00:00 2001
From: Shilei Tian <shilei.t...@amd.com>
Date: Tue, 4 Jun 2024 16:31:19 -0400
Subject: [PATCH] AMDGPU: Builtin & CodeGen support for
v_cvt_scalef32_sr_{bf8|fp8}_{f16|bf16|f32}
Co-authored-by: Shilei Tian <shilei.t...@amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 6 +
.../CodeGenOpenCL/builtins-amdgcn-gfx950.cl | 258 +++++++++++++
.../builtins-amdgcn-error-gfx950-param.cl | 6 +
.../builtins-amdgcn-error-gfx950.cl | 9 +-
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 14 +-
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 6 +
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 6 +
.../AMDGPU/llvm.amdgcn.cvt.scalef32.sr.ll | 346 ++++++++++++++++++
8 files changed, 646 insertions(+), 5 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.sr.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 61039938267feb..54bbec97b17702 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -601,6 +601,12 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp4_bf16,
"UiUiV2yfIi", "nc", "f
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16, "UiUiV2hUifIi",
"nc", "fp4-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16, "UiUiV2yUifIi",
"nc", "fp4-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32, "UiUiV2fUifIi",
"nc", "fp4-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16, "iiyUifIi", "nc",
"bf8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_f16, "iihUifIi", "nc",
"bf8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_f32, "iifUifIi", "nc",
"bf8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_bf16, "iiyUifIi", "nc",
"fp8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_f16, "iihUifIi", "nc",
"fp8-cvt-scale-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_f32, "iifUifIi", "nc",
"fp8-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_bf16, "V6UiV32yUif",
"nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f16, "V6UiV32hUif",
"nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
index 64403f0bf94ebd..bbfb16e05a53e0 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
@@ -1400,3 +1400,261 @@ void test_cvt_scalef32_sr_pk32(global uint6 *out6,
bfloat32 srcbf32, half32 srch
*out6 = __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f16(srch32, src1, src2);
*out6 = __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32(srcf32, src1, src2);
}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_bf8_bf16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca bfloat, align 2, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, 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 bfloat [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, ptr addrspace(1) [[TMP0]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 [[TMP1]], bfloat [[TMP2]], i32
[[TMP3]], float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 [[TMP8]], bfloat [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 [[TMP15]], bfloat [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_bf8_bf16(global unsigned *out, __bf16 src, uint
seed, float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out, src, seed, scale, 2);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_bf8_f16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, 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 half [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align
2
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, ptr addrspace(1) [[TMP0]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 [[TMP1]], half [[TMP2]], i32 [[TMP3]],
float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 [[TMP8]], half [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 [[TMP15]], half [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_bf8_f16(global unsigned *out, half src, uint seed,
float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(*out, src, seed, scale, 2);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_bf8_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: [[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 float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, 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:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 [[TMP1]], float [[TMP2]], i32
[[TMP3]], float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 [[TMP8]], float [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 [[TMP15]], float [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_bf8_f32(global unsigned *out, float src, uint seed,
float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(*out, src, seed, scale, 2);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_fp8_bf16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca bfloat, align 2, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, 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 bfloat [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, ptr addrspace(1) [[TMP0]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 [[TMP1]], bfloat [[TMP2]], i32
[[TMP3]], float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 [[TMP8]], bfloat [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load bfloat, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 [[TMP15]], bfloat [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_fp8_bf16(global unsigned *out, __bf16 src, uint
seed, float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(*out, src, seed, scale, 2);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_fp8_f16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK-NEXT: [[SEED_ADDR:%.*]] = alloca i32, 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 half [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align
2
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, ptr addrspace(1) [[TMP0]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 [[TMP1]], half [[TMP2]], i32 [[TMP3]],
float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 [[TMP8]], half [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load half, ptr addrspace(5) [[SRC_ADDR]],
align 2
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 [[TMP15]], half [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_fp8_f16(global unsigned *out, half src, uint seed,
float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(*out, src, seed, scale, 2);
+}
+
+// CHECK-LABEL: @test_cvt_scalef32_sr_fp8_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: [[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 float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: store i32 [[SEED:%.*]], ptr addrspace(5) [[SEED_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 i32, 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:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP5:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 [[TMP1]], float [[TMP2]], i32
[[TMP3]], float [[TMP4]], i32 0)
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP5]], ptr addrspace(1) [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP12:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 [[TMP8]], float [[TMP9]], i32
[[TMP10]], float [[TMP11]], i32 1)
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP12]], ptr addrspace(1) [[TMP13]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(1) [[TMP14]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]],
align 4
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SEED_ADDR]],
align 4
+// CHECK-NEXT: [[TMP18:%.*]] = load float, ptr addrspace(5) [[SCALE_ADDR]],
align 4
+// CHECK-NEXT: [[TMP19:%.*]] = call i32
@llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 [[TMP15]], float [[TMP16]], i32
[[TMP17]], float [[TMP18]], i32 2)
+// CHECK-NEXT: [[TMP20:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[TMP20]], align 4
+// CHECK-NEXT: ret void
+//
+void test_cvt_scalef32_sr_fp8_f32(global unsigned *out, float src, uint seed,
float scale)
+{
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(*out, src, seed, scale, 0);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(*out, src, seed, scale, 1);
+ *out = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(*out, src, seed, scale, 2);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
index 7138776e966cfe..65078312a6e8a8 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl
@@ -191,4 +191,10 @@ void test_cvt_scalef32(global half2* out_v2f16, global
float* out_f32, uint src,
*out = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(*out, src0_v2f16, 0,
scale, index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16' must be a constant integer}}
*out = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(*out, src0_v2bf16, 0,
scale, index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16' must be a constant integer}}
*out = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(*out, src0_v2f32, 0,
scale, index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32' must be a constant integer}}
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16' must be a constant integer}}
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_bf8_f16' must be a constant integer}}
+ *out = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(*out, src0, 0, scale,
index); // expected-error{{argument to
'__builtin_amdgcn_cvt_scalef32_sr_bf8_f32' must be a constant integer}}
+ *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}}
}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl
index e0cde1d3ad87bb..d1c134c604dfcd 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl
@@ -28,7 +28,8 @@ void test(__global float4* out0, half8 a0, half8 b0, float4
c0,
__global float16* out13, int4 a13, int8 b13, float16 c13,
__global float4* out14, int8 a14, int8 b14, float4 c14, int d14, int
e14,
__global float16* out15, int8 a15, int8 b15, float16 c15, int d15,
int e15,
- __global uint2* out16, int a16, int b16) {
+ __global uint2* out16, int a16, int b16,
+ __global int *out17, float a17, int b17, float c17) {
*out0 = __builtin_amdgcn_mfma_f32_16x16x32_f16(a0, b0, c0, 0, 0, 0); //
expected-error{{'__builtin_amdgcn_mfma_f32_16x16x32_f16' needs target feature
gfx950-insts}}
*out1 = __builtin_amdgcn_mfma_f32_32x32x16_f16(a1, b1, c1, 0, 0, 0); //
expected-error{{'__builtin_amdgcn_mfma_f32_32x32x16_f16' needs target feature
gfx950-insts}}
*out2 = __builtin_amdgcn_mfma_f32_32x32x16_bf16(a2, b2, c2, 0, 0, 0); //
expected-error{{'__builtin_amdgcn_mfma_f32_32x32x16_bf16' needs target feature
gfx950-insts}}
@@ -53,4 +54,10 @@ void test(__global float4* out0, half8 a0, half8 b0, float4
c0,
*out15 = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a15, b15, c15, 0,
0, 0, d15, 0, e15); //
expected-error{{'__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' needs target
feature gfx950-insts}}
*out16 = __builtin_amdgcn_permlane16_swap(a16, b16, false, false); //
expected-error{{'__builtin_amdgcn_permlane16_swap' needs target feature
permlane16-swap}}
*out16 = __builtin_amdgcn_permlane32_swap(a16, b16, false, false); //
expected-error{{'__builtin_amdgcn_permlane32_swap' needs target feature
permlane32-swap}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out17, a17, b17, c17,
0); // expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16' needs target
feature bf8-cvt-scale-insts}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(*out17, a17, b17, c17, 0);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_bf8_f16' needs target
feature bf8-cvt-scale-insts}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(*out17, a17, b17, c17, 0);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_bf8_f32' needs target
feature bf8-cvt-scale-insts}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(*out17, a17, b17, c17,
0); // expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_fp8_bf16' needs target
feature fp8-cvt-scale-insts}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(*out17, a17, b17, c17, 0);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_fp8_f16' needs target
feature fp8-cvt-scale-insts}}
+ *out17 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(*out17, a17, b17, c17, 0);
// expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_fp8_f32' needs target
feature fp8-cvt-scale-insts}}
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 73f3559ab05a48..1561942c81c67d 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -684,7 +684,7 @@ class AMDGPUCvtScaleF16ToFP4TiedInputIntrinsic<LLVMType
SrcTy, string name> : De
[IntrNoMem, IntrWillReturn, ImmArg<ArgIndex<3>>]
>, ClangBuiltin<"__builtin_amdgcn_"#name>;
-class AMDGPUCvtScaleF32SRToFP4TiedInputIntrinsic<LLVMType Src0Ty, string name>
: DefaultAttrsIntrinsic<
+class AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<LLVMType Src0Ty,
string name> : DefaultAttrsIntrinsic<
[llvm_i32_ty],
[llvm_i32_ty, // old_vdst
Src0Ty, // src0
@@ -746,9 +746,15 @@ def int_amdgcn_cvt_scalef32_pk_bf16_fp8 :
AMDGPUCvtScale_pk_FromFP8BF8Intrinsi
def int_amdgcn_cvt_scalef32_pk_fp4_f16 :
AMDGPUCvtScaleF16ToFP4TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_pk_fp4_f16">;
def int_amdgcn_cvt_scalef32_pk_fp4_bf16:
AMDGPUCvtScaleF16ToFP4TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_scalef32_pk_fp4_bf16">;
-def int_amdgcn_cvt_scalef32_sr_pk_fp4_f16:
AMDGPUCvtScaleF32SRToFP4TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_sr_pk_fp4_f16">;
-def int_amdgcn_cvt_scalef32_sr_pk_fp4_bf16:
AMDGPUCvtScaleF32SRToFP4TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_scalef32_sr_pk_fp4_bf16">;
-def int_amdgcn_cvt_scalef32_sr_pk_fp4_f32:
AMDGPUCvtScaleF32SRToFP4TiedInputIntrinsic<llvm_v2f32_ty,
"cvt_scalef32_sr_pk_fp4_f32">;
+def int_amdgcn_cvt_scalef32_sr_pk_fp4_f16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_v2f16_ty,
"cvt_scalef32_sr_pk_fp4_f16">;
+def int_amdgcn_cvt_scalef32_sr_pk_fp4_bf16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_v2bf16_ty,
"cvt_scalef32_sr_pk_fp4_bf16">;
+def int_amdgcn_cvt_scalef32_sr_pk_fp4_f32:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_v2f32_ty,
"cvt_scalef32_sr_pk_fp4_f32">;
+def int_amdgcn_cvt_scalef32_sr_bf8_bf16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_bfloat_ty,
"cvt_scalef32_sr_bf8_bf16">;
+def int_amdgcn_cvt_scalef32_sr_bf8_f16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_half_ty,
"cvt_scalef32_sr_bf8_f16">;
+def int_amdgcn_cvt_scalef32_sr_bf8_f32:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_float_ty,
"cvt_scalef32_sr_bf8_f32">;
+def int_amdgcn_cvt_scalef32_sr_fp8_bf16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_bfloat_ty,
"cvt_scalef32_sr_fp8_bf16">;
+def int_amdgcn_cvt_scalef32_sr_fp8_f16:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_half_ty,
"cvt_scalef32_sr_fp8_f16">;
+def int_amdgcn_cvt_scalef32_sr_fp8_f32:
AMDGPUCvtScaleBF16F16F32SRToFP4BF8F8TiedInputIntrinsic<llvm_float_ty,
"cvt_scalef32_sr_fp8_f32">;
def int_amdgcn_prng_b32 : DefaultAttrsIntrinsic<
[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 158603a7aff879..8d389d2fb9ea9c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4584,6 +4584,12 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_scalef32_sr_pk32_fp6_bf16:
case Intrinsic::amdgcn_cvt_scalef32_sr_pk32_fp6_f16:
case Intrinsic::amdgcn_cvt_scalef32_sr_pk32_fp6_f32:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_bf8_bf16:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_bf8_f16:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_bf8_f32:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_fp8_bf16:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_fp8_f16:
+ case Intrinsic::amdgcn_cvt_scalef32_sr_fp8_f32:
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 3a79532cecb917..00caea1f923391 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -1197,6 +1197,12 @@ class
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<SDPatternOperator node, VOP3_Pse
def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_pk_fp4_f16,
V_CVT_SCALEF32_SR_PK_FP4_F16_e64, v2f16>;
def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_pk_fp4_bf16,
V_CVT_SCALEF32_SR_PK_FP4_BF16_e64, v2bf16>;
def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_pk_fp4_f32,
V_CVT_SCALEF32_SR_PK_FP4_F32_e64, v2f32>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_bf8_bf16,
V_CVT_SCALEF32_SR_BF8_BF16_e64, bf16>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_bf8_f16,
V_CVT_SCALEF32_SR_BF8_F16_e64, f16>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_bf8_f32,
V_CVT_SCALEF32_SR_BF8_F32_e64, f32>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_fp8_bf16,
V_CVT_SCALEF32_SR_FP8_BF16_e64, bf16>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_fp8_f16,
V_CVT_SCALEF32_SR_FP8_F16_e64, f16>;
+def :
Cvt_Scale_SR_PK_BF16F16F32ToFP4BF8FP8_Pat<int_amdgcn_cvt_scalef32_sr_fp8_f32,
V_CVT_SCALEF32_SR_FP8_F32_e64, f32>;
let SubtargetPredicate = isGFX10Plus in {
let isCommutable = 1, isReMaterializable = 1 in {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.sr.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.sr.ll
new file mode 100644
index 00000000000000..e1bf9f0daa1efa
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.sr.ll
@@ -0,0 +1,346 @@
+; 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 -verify-machineinstrs
-o - %s | FileCheck -check-prefix=GFX950 %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs
-o - %s | FileCheck -check-prefix=GFX950 %s
+
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 %old, bfloat %src, i32
%seed, float %scale, i32 %dst_sel)
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 %old, half %src, i32
%seed, float %scale, i32 %dst_sel)
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 %old, float %src, i32
%seed, float %scale, i32 %dst_sel)
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 %old, bfloat %src, i32
%seed, float %scale, i32 %dst_sel)
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 %old, half %src, i32
%seed, float %scale, i32 %dst_sel)
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src, i32
%seed, float %scale, i32 %dst_sel)
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_bf16_dst_sel_0(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_bf16_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_bf16 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_bf16_dst_sel_1(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_bf16_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_bf16 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_bf16_dst_sel_2(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_bf16_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_bf16 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_bf16_dst_sel_3(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_bf16_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_bf16 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f16_dst_sel_0(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f16_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f16 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f16_dst_sel_1(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f16_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f16 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f16_dst_sel_2(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f16_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f16 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f16_dst_sel_3(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f16_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f16 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f32_dst_sel_0(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f32_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f32 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f32_dst_sel_1(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f32_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f32 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f32_dst_sel_2(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f32_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f32 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_bf8_f32_dst_sel_3(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_bf8_f32_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_bf8_f32 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.bf8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_bf16_dst_sel_0(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_bf16_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_bf16 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_bf16_dst_sel_1(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_bf16_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_bf16 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_bf16_dst_sel_2(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_bf16_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_bf16 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_bf16_dst_sel_3(ptr
addrspace(1) %out, bfloat %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_bf16_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_bf16 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.bf16(i32 %old, bfloat %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f16_dst_sel_0(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f16_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f16 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f16_dst_sel_1(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f16_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f16 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f16_dst_sel_2(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f16_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f16 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f16_dst_sel_3(ptr addrspace(1)
%out, half %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f16_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f16 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f16(i32 %old, half %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f32_dst_sel_0(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f32_dst_sel_0:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f32 v5, v2, v3, v4
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 0)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f32_dst_sel_1(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f32_dst_sel_1:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f32 v5, v2, v3, v4 op_sel:[0,0,1,0]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 1)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f32_dst_sel_2(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f32_dst_sel_2:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f32 v5, v2, v3, v4 op_sel:[0,0,0,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 2)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scalef32_sr_fp8_f32_dst_sel_3(ptr addrspace(1)
%out, float %src, i32 %seed, float %scale) {
+; GFX950-LABEL: test_cvt_scalef32_sr_fp8_f32_dst_sel_3:
+; GFX950: ; %bb.0:
+; GFX950-NEXT: global_load_dword v5, v[0:1], off
+; GFX950-NEXT: s_waitcnt vmcnt(0)
+; GFX950-NEXT: v_cvt_scalef32_sr_fp8_f32 v5, v2, v3, v4 op_sel:[0,0,1,1]
+; GFX950-NEXT: global_store_dword v[0:1], v5, off
+; GFX950-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src,
i32 %seed, float %scale, i32 3)
+ store i32 %cvt, ptr addrspace(1) %out, align 4
+ ret void
+}
_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
