[llvm-branch-commits] [clang] [llvm] [AMDGPU] Add intrinsics and builtins for v_cvt_scale_pk32_* instructions (PR #222617)
Mariusz Sikora via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 10 04:34:23 PDT 2026
https://github.com/mariusz-sikora-at-amd created https://github.com/llvm/llvm-project/pull/222617
None
>From 462e81737d02725f5c9fbefdcb449c7c589d7a47 Mon Sep 17 00:00:00 2001
From: Mariusz Sikora <mariusz.sikora at amd.com>
Date: Mon, 7 Sep 2026 03:28:53 -0400
Subject: [PATCH] [AMDGPU] Add intrinsics and builtins for v_cvt_scale_pk32_*
instructions
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 24 +
.../include/clang/Basic/BuiltinsAMDGPUDocs.td | 60 +
clang/lib/Sema/SemaAMDGPU.cpp | 6 +
.../builtins-amdgcn-gfx12-err.cl | 12 +-
.../builtins-amdgcn-gfx13-err.cl | 26 +
.../builtins-amdgcn-gfx13-w32-err.cl | 20 +
.../CodeGenOpenCL/builtins-amdgcn-gfx13.cl | 58 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 8 +-
llvm/lib/Target/AMDGPU/AMDGPU.td | 3 +-
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 12 +
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 6 +
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 12 +-
.../llvm.amdgcn.cvt.scale.pk32.gfx13.ll | 1308 +++++++++++++++++
13 files changed, 1546 insertions(+), 9 deletions(-)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-err.cl
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w32-err.cl
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scale.pk32.gfx13.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 9eed0f60a6c50..9b72393e12577 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1080,6 +1080,30 @@ def __builtin_amdgcn_cvt_scale_pk16_f16_bf6 : AMDGPUBuiltin<"_ExtVector<16, _Flo
def __builtin_amdgcn_cvt_scale_pk16_bf16_bf6 : AMDGPUBuiltin<"_ExtVector<16, __bf16>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_scale_pk16_f32_fp6 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_scale_pk16_f32_bf6 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk32_bf16_bf6 : AMDGPUBuiltin<"_ExtVector<32, __bf16>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32BF16BF6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
+def __builtin_amdgcn_cvt_scale_pk32_bf16_fp6 : AMDGPUBuiltin<"_ExtVector<32, __bf16>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32BF16FP6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
+def __builtin_amdgcn_cvt_scale_pk32_f16_bf6 : AMDGPUBuiltin<"_ExtVector<32, _Float16>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32F16BF6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
+def __builtin_amdgcn_cvt_scale_pk32_f16_fp6 : AMDGPUBuiltin<"_ExtVector<32, _Float16>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32F16FP6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
+def __builtin_amdgcn_cvt_scale_pk32_f32_bf6 : AMDGPUBuiltin<"_ExtVector<32, float>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32F32BF6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
+def __builtin_amdgcn_cvt_scale_pk32_f32_fp6 : AMDGPUBuiltin<"_ExtVector<32, float>(_ExtVector<6, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32"> {
+ let Documentation = [DocCvtScalePk32F32FP6];
+ let ArgNames = ["src", "scale", "scale_sel"];
+}
def __builtin_amdgcn_cvt_scalef32_pk8_fp8_bf16 : AMDGPUBuiltin<"_ExtVector<2, unsigned int>(_ExtVector<8, __bf16>, float)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_scalef32_pk8_bf8_bf16 : AMDGPUBuiltin<"_ExtVector<2, unsigned int>(_ExtVector<8, __bf16>, float)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_scalef32_pk8_fp8_f16 : AMDGPUBuiltin<"_ExtVector<2, unsigned int>(_ExtVector<8, _Float16>, float)", [Const], "gfx1250-insts">;
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
index bfccd7de6590a..0126aac588188 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -877,6 +877,66 @@ integers.
}];
}
+def DocCvtScalePk32BF16BF6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component BF6 input ``src`` to packed 32-component BF16
+values, then scales the results using the scale factor ``scale``. ``scale_sel``
+must be a compile-time constant in the range [0, 15]. Only available in
+wavefront size 32.
+}];
+}
+
+def DocCvtScalePk32BF16FP6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component FP6 input ``src`` to packed 32-component BF16
+values, then scales the results using the scale factor ``scale``. ``scale_sel``
+must be a compile-time constant in the range [0, 15]. Only available in
+wavefront size 32.
+}];
+}
+
+def DocCvtScalePk32F16BF6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component BF6 input ``src`` to packed 32-component
+half-precision floating-point values, then scales the results using the scale
+factor ``scale``. ``scale_sel`` must be a compile-time constant in the range
+[0, 15]. Only available in wavefront size 32.
+}];
+}
+
+def DocCvtScalePk32F16FP6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component FP6 input ``src`` to packed 32-component
+half-precision floating-point values, then scales the results using the scale
+factor ``scale``. ``scale_sel`` must be a compile-time constant in the range
+[0, 15]. Only available in wavefront size 32.
+}];
+}
+
+def DocCvtScalePk32F32BF6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component BF6 input ``src`` to packed 32-component
+single-precision floating-point values, then scales the results using the scale
+factor ``scale``. ``scale_sel`` must be a compile-time constant in the range
+[0, 15]. Only available in wavefront size 32.
+}];
+}
+
+def DocCvtScalePk32F32FP6 : Documentation {
+ let Category = DocCatAMDGPUConversion;
+ let Content = [{
+Converts a packed 32-component FP6 input ``src`` to packed 32-component
+single-precision floating-point values, then scales the results using the scale
+factor ``scale``. ``scale_sel`` must be a compile-time constant in the range
+[0, 15]. Only available in wavefront size 32.
+}];
+}
+
//===----------------------------------------------------------------------===//
// Prefetch Builtins
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index ee6d989d2b107..b1f6fa33649bd 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -213,6 +213,12 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI,
case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_bf16_bf6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_bf16_fp6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f16_bf6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f16_fp6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f32_bf6:
+ case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f32_fp6:
return SemaRef.BuiltinConstantArgRange(TheCall, 2, 0, 15);
case AMDGPU::BI__builtin_amdgcn_av_load_b128:
return checkAVLoadStore(TheCall, /*IsStore=*/false);
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
index d8fb243865f76..a7a4705dea514 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
@@ -24,6 +24,8 @@ typedef short v8s __attribute__((ext_vector_type(8)));
typedef short v16s __attribute__((ext_vector_type(16)));
typedef short v32s __attribute__((ext_vector_type(32)));
typedef double v4d __attribute__((ext_vector_type(4)));
+typedef uint v6u __attribute__((ext_vector_type(6)));
+typedef __bf16 v32bf __attribute__((ext_vector_type(32)));
void builtin_test_unsupported(double a_double, float a_float,
int a_int, long a_long,
@@ -31,7 +33,8 @@ void builtin_test_unsupported(double a_double, float a_float,
v2s a_v2s, v4s a_v4s, v8s a_v8s,
v2i a_v2i, v4i a_v4i, v16i a_v16i, v32i a_v32i,
v2f a_v2f, v4f a_v4f, v16f a_v16f, v32f a_v32f,
- v4h a_v4h, v8h a_v8h,
+ v4h a_v4h, v8h a_v8h, v32h a_v32h,
+ v6u a_v6u, v32bf a_v32bf,
uint a, uint b) {
@@ -94,4 +97,11 @@ void builtin_test_unsupported(double a_double, float a_float,
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8' needs target feature fp8-insts}}
+
+ a_v32bf = __builtin_amdgcn_cvt_scale_pk32_bf16_bf6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ a_v32bf = __builtin_amdgcn_cvt_scale_pk32_bf16_fp6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ a_v32h = __builtin_amdgcn_cvt_scale_pk32_f16_bf6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ a_v32h = __builtin_amdgcn_cvt_scale_pk32_f16_fp6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ a_v32f = __builtin_amdgcn_cvt_scale_pk32_f32_bf6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ a_v32f = __builtin_amdgcn_cvt_scale_pk32_f32_fp6(a_v6u, a, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-err.cl
new file mode 100644
index 0000000000000..27c9409be9f86
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-err.cl
@@ -0,0 +1,26 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu13.10-unknown-unknown -verify -emit-llvm -o - %s
+
+typedef unsigned int uint;
+typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
+typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
+typedef half __attribute__((ext_vector_type(32))) half32;
+typedef float __attribute__((ext_vector_type(32))) float32;
+
+void test_cvt_scale_pk32(global bfloat32 *outbf32, global half32 *outhalf32, global float32 *outf32, uint6 src, uint scale, uint scale_sel)
+{
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_bf6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_bf6' must be a constant integer}}
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_fp6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_fp6' must be a constant integer}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_bf6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_bf6' must be a constant integer}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_fp6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_fp6' must be a constant integer}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_bf6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_bf6' must be a constant integer}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_fp6(src, scale, scale_sel); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_fp6' must be a constant integer}}
+
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_bf6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_fp6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_bf6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_fp6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_bf6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_fp6(src, scale, 16); // expected-error {{argument value 16 is outside the valid range [0, 15]}}
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w32-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w32-err.cl
new file mode 100644
index 0000000000000..352c9535ef8d7
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w32-err.cl
@@ -0,0 +1,20 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu13.10-unknown-unknown \
+// RUN: -target-feature +wavefrontsize64 -verify -S -o - %s
+
+typedef unsigned int uint;
+typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
+typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
+typedef half __attribute__((ext_vector_type(32))) half32;
+typedef float __attribute__((ext_vector_type(32))) float32;
+
+void test_cvt_scale_pk32_w64(global bfloat32 *outbf32, global half32 *outhalf32, global float32 *outf32, uint6 src, uint scale)
+{
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_bf6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_fp6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_bf16_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_bf6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_fp6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f16_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_bf6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_bf6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_fp6(src, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk32_f32_fp6' needs target feature fp6bf6-to-f16bf16f32-cvt-scale-insts,wavefrontsize32}}
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index 00e65563a1868..d6247dd88ce7f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -6,7 +6,9 @@
typedef unsigned int __attribute__((ext_vector_type(4))) uint4;
typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
typedef float __attribute__((ext_vector_type(32))) float32;
+typedef half __attribute__((ext_vector_type(32))) half32;
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
+typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
typedef unsigned int uint;
// CHECK-LABEL: @test_cvt_scalef32_pk32_bf6_f32(
@@ -211,3 +213,59 @@ void test_s_prefetch_data(global float *gp, unsigned int len)
__builtin_amdgcn_s_prefetch_data(gp, len);
}
+
+// CHECK-LABEL: @test_cvt_scale_pk32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUTBF32_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[OUTHALF32_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[OUTF32_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 i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SCALE_SEL_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUTBF32:%.*]], ptr addrspace(5) [[OUTBF32_ADDR]], align 8
+// CHECK-NEXT: store ptr addrspace(1) [[OUTHALF32:%.*]], ptr addrspace(5) [[OUTHALF32_ADDR]], align 8
+// CHECK-NEXT: store ptr addrspace(1) [[OUTF32:%.*]], ptr addrspace(5) [[OUTF32_ADDR]], align 8
+// CHECK-NEXT: store <6 x i32> [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: store i32 [[SCALE:%.*]], ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SCALE_SEL:%.*]], ptr addrspace(5) [[SCALE_SEL_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <6 x i32>, ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> [[TMP0]], i32 [[TMP1]], i32 5)
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTBF32_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 i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> [[TMP4]], i32 [[TMP5]], i32 4)
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTBF32_ADDR]], align 8
+// CHECK-NEXT: store <32 x bfloat> [[TMP6]], ptr addrspace(1) [[TMP7]], align 64
+// CHECK-NEXT: [[TMP8:%.*]] = load <6 x i32>, ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> [[TMP8]], i32 [[TMP9]], i32 7)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTHALF32_ADDR]], align 8
+// CHECK-NEXT: store <32 x half> [[TMP10]], ptr addrspace(1) [[TMP11]], align 64
+// CHECK-NEXT: [[TMP12:%.*]] = load <6 x i32>, ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP14:%.*]] = call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> [[TMP12]], i32 [[TMP13]], i32 7)
+// CHECK-NEXT: [[TMP15:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTHALF32_ADDR]], align 8
+// CHECK-NEXT: store <32 x half> [[TMP14]], ptr addrspace(1) [[TMP15]], align 64
+// CHECK-NEXT: [[TMP16:%.*]] = load <6 x i32>, ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP18:%.*]] = call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> [[TMP16]], i32 [[TMP17]], i32 6)
+// CHECK-NEXT: [[TMP19:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTF32_ADDR]], align 8
+// CHECK-NEXT: store <32 x float> [[TMP18]], ptr addrspace(1) [[TMP19]], align 128
+// CHECK-NEXT: [[TMP20:%.*]] = load <6 x i32>, ptr addrspace(5) [[SRC_ADDR]], align 32
+// CHECK-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[SCALE_ADDR]], align 4
+// CHECK-NEXT: [[TMP22:%.*]] = call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> [[TMP20]], i32 [[TMP21]], i32 6)
+// CHECK-NEXT: [[TMP23:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUTF32_ADDR]], align 8
+// CHECK-NEXT: store <32 x float> [[TMP22]], ptr addrspace(1) [[TMP23]], align 128
+// CHECK-NEXT: ret void
+//
+void test_cvt_scale_pk32(global bfloat32 *outbf32, global half32 *outhalf32, global float32 *outf32, uint6 src, uint scale, uint scale_sel)
+{
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_bf6(src, scale, 5);
+ *outbf32 = __builtin_amdgcn_cvt_scale_pk32_bf16_fp6(src, scale, 4);
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_bf6(src, scale, 7);
+ *outhalf32 = __builtin_amdgcn_cvt_scale_pk32_f16_fp6(src, scale, 7);
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_bf6(src, scale, 6);
+ *outf32 = __builtin_amdgcn_cvt_scale_pk32_f32_fp6(src, scale, 6);
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 6b29b8612828b..d1bd24df96630 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -673,7 +673,7 @@ def int_amdgcn_cvt_sr_bf8_f16 : PureIntrinsic<
[ImmArg<ArgIndex<3>>, Range<ArgIndex<3>, 0, 4>]
>, ClangBuiltin<"__builtin_amdgcn_cvt_sr_bf8_f16">;
-// Note: these gfx1250 intrinsics are convergent because they read scales from other lanes.
+// Note: these intrinsics are convergent because they read scales from other lanes.
// llvm.amdgcn.cvt.scale.pk32.f16.bf6 v32f16 vdst, v6i32 src0, i32 scale_sel [0..15]
class AMDGPUCvtScaleIntrinsic<LLVMType DstTy, LLVMType Src0Ty, string name> : DefaultAttrsIntrinsic<
[DstTy], [Src0Ty, llvm_i32_ty, llvm_i32_ty],
@@ -704,6 +704,12 @@ def int_amdgcn_cvt_scale_pk16_f16_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v16f16_ty,
def int_amdgcn_cvt_scale_pk16_bf16_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v16bf16_ty, llvm_v3i32_ty, "cvt_scale_pk16_bf16_fp6">;
def int_amdgcn_cvt_scale_pk16_f32_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v16f32_ty, llvm_v3i32_ty, "cvt_scale_pk16_f32_fp6">;
def int_amdgcn_cvt_scale_pk16_f32_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v16f32_ty, llvm_v3i32_ty, "cvt_scale_pk16_f32_bf6">;
+def int_amdgcn_cvt_scale_pk32_bf16_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v32bf16_ty, llvm_v6i32_ty, "cvt_scale_pk32_bf16_bf6">;
+def int_amdgcn_cvt_scale_pk32_bf16_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v32bf16_ty, llvm_v6i32_ty, "cvt_scale_pk32_bf16_fp6">;
+def int_amdgcn_cvt_scale_pk32_f16_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v32f16_ty, llvm_v6i32_ty, "cvt_scale_pk32_f16_bf6">;
+def int_amdgcn_cvt_scale_pk32_f16_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v32f16_ty, llvm_v6i32_ty, "cvt_scale_pk32_f16_fp6">;
+def int_amdgcn_cvt_scale_pk32_f32_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v32f32_ty, llvm_v6i32_ty, "cvt_scale_pk32_f32_bf6">;
+def int_amdgcn_cvt_scale_pk32_f32_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v32f32_ty, llvm_v6i32_ty, "cvt_scale_pk32_f32_fp6">;
class AMDGPUCvtScaleF32ToFP6BF6Intrinsic<LLVMType DstTy, LLVMType Src0Ty, LLVMType Src1Ty, string name> : PureIntrinsic<
[DstTy], [Src0Ty, Src1Ty, llvm_float_ty]
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 70e29b4513f4a..6f34492d69224 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -3229,7 +3229,8 @@ def AMDGPUFrontendVisibleFeatures {
FeatureDot6Insts, FeatureDot7Insts, FeatureDot8Insts,
FeatureDot9Insts, FeatureExtendedImageInsts, FeatureF16BF16ToFP6BF6ConversionScaleInsts,
FeatureF32ToF16BF16ConversionSRInsts, FeatureF32ToFP6BF6ConversionScaleInsts, FeatureFP4ConversionScaleInsts,
- FeatureFP6BF6ConversionScaleInsts, FeatureFP8ConversionInsts, FeatureFP8ConversionScaleInsts,
+ FeatureFP6BF6ConversionScaleInsts, FeatureFP6BF6ToF16BF16F32ConversionScaleInsts,
+ FeatureFP8ConversionInsts, FeatureFP8ConversionScaleInsts,
FeatureFP8E5M3Insts, FeatureFP8Insts,
FeatureFlatBufferGlobalAtomicFaddF64Inst, FeatureFlatGlobalInsts,
FeatureGFX10Insts, FeatureGFX10_3Insts, FeatureGFX11Insts,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 8aafbbc4dfdf8..8237f2e9ee827 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -2248,6 +2248,18 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
.Any({{DivV16S32}, {{VgprV16S32}, {IntrId, VgprV3S32, Vgpr32}}})
.Any({{UniV16S32}, {{UniInVgprV16S32}, {IntrId, VgprV3S32, Vgpr32}}});
+ addRulesForIOpcs(
+ {amdgcn_cvt_scale_pk32_f16_bf6, amdgcn_cvt_scale_pk32_f16_fp6,
+ amdgcn_cvt_scale_pk32_bf16_bf6, amdgcn_cvt_scale_pk32_bf16_fp6},
+ Standard)
+ .Any({{DivV32S16}, {{VgprV32S16}, {IntrId, VgprV6S32, Vgpr32}}})
+ .Any({{UniV32S16}, {{UniInVgprV32S16}, {IntrId, VgprV6S32, Vgpr32}}});
+
+ addRulesForIOpcs(
+ {amdgcn_cvt_scale_pk32_f32_bf6, amdgcn_cvt_scale_pk32_f32_fp6}, Standard)
+ .Any({{DivV32S32}, {{VgprV32S32}, {IntrId, VgprV6S32, Vgpr32}}})
+ .Any({{UniV32S32}, {{UniInVgprV32S32}, {IntrId, VgprV6S32, Vgpr32}}});
+
addRulesForIOpcs({amdgcn_cvt_scale_pk8_f16_bf8, amdgcn_cvt_scale_pk8_f16_fp8,
amdgcn_cvt_scale_pk8_bf16_bf8,
amdgcn_cvt_scale_pk8_bf16_fp8},
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index a152ce6c20ce5..65032b4cbe9fa 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4668,6 +4668,12 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_scale_pk16_bf16_bf6:
case Intrinsic::amdgcn_cvt_scale_pk16_f32_fp6:
case Intrinsic::amdgcn_cvt_scale_pk16_f32_bf6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_f16_fp6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_bf16_fp6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_f16_bf6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_bf16_bf6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_f32_fp6:
+ case Intrinsic::amdgcn_cvt_scale_pk32_f32_bf6:
case Intrinsic::amdgcn_cvt_scalef32_pk8_fp8_bf16:
case Intrinsic::amdgcn_cvt_scalef32_pk8_bf8_bf16:
case Intrinsic::amdgcn_cvt_scalef32_pk8_fp8_f16:
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index c50ea2067c010..065d893eb9935 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -1935,12 +1935,12 @@ let SubtargetPredicate = isGFX1250Plus in {
let SubtargetPredicate = HasFP6BF6ToF16BF16F32ConversionScaleInsts,
ReadsModeReg = 0, Constraints = "@earlyclobber $vdst" in {
let WaveSizePredicate = isWave32 in {
- defm V_CVT_SCALE_PK32_BF16_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_bf16_bf6", VOP_V32BF16_V6I32_I32, null_frag>;
- defm V_CVT_SCALE_PK32_BF16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_bf16_fp6", VOP_V32BF16_V6I32_I32, null_frag>;
- defm V_CVT_SCALE_PK32_F16_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f16_bf6", VOP_V32F16_V6I32_I32, null_frag>;
- defm V_CVT_SCALE_PK32_F16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f16_fp6", VOP_V32F16_V6I32_I32, null_frag>;
- defm V_CVT_SCALE_PK32_F32_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f32_bf6", VOP_V32F32_V6I32_I32, null_frag>;
- defm V_CVT_SCALE_PK32_F32_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f32_fp6", VOP_V32F32_V6I32_I32, null_frag>;
+ defm V_CVT_SCALE_PK32_BF16_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_bf16_bf6", VOP_V32BF16_V6I32_I32, int_amdgcn_cvt_scale_pk32_bf16_bf6>;
+ defm V_CVT_SCALE_PK32_BF16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_bf16_fp6", VOP_V32BF16_V6I32_I32, int_amdgcn_cvt_scale_pk32_bf16_fp6>;
+ defm V_CVT_SCALE_PK32_F16_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f16_bf6", VOP_V32F16_V6I32_I32, int_amdgcn_cvt_scale_pk32_f16_bf6>;
+ defm V_CVT_SCALE_PK32_F16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f16_fp6", VOP_V32F16_V6I32_I32, int_amdgcn_cvt_scale_pk32_f16_fp6>;
+ defm V_CVT_SCALE_PK32_F32_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f32_bf6", VOP_V32F32_V6I32_I32, int_amdgcn_cvt_scale_pk32_f32_bf6>;
+ defm V_CVT_SCALE_PK32_F32_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk32_f32_fp6", VOP_V32F32_V6I32_I32, int_amdgcn_cvt_scale_pk32_f32_fp6>;
} // End WaveSizePredicate = isWave32
} // End SubtargetPredicate = HasFP6BF6ToF16BF16F32ConversionScaleInsts
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scale.pk32.gfx13.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scale.pk32.gfx13.ll
new file mode 100644
index 0000000000000..2fc72ee91d350
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scale.pk32.gfx13.ll
@@ -0,0 +1,1308 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -global-isel=0 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefix=GFX13-SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefix=GFX13-GISEL %s
+
+declare <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+declare <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+declare <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+declare <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+declare <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+declare <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 %scale, i32 %scale_sel)
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s0 :: v_dual_mov_b32 v20, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s2 :: v_dual_mov_b32 v22, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s4 :: v_dual_mov_b32 v24, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s5 :: v_dual_mov_b32 v23, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s3 :: v_dual_mov_b32 v21, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s1 :: v_dual_mov_b32 v19, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v18, s0 :: v_dual_mov_b32 v19, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v20, s2 :: v_dual_mov_b32 v21, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v22, s4 :: v_dual_mov_b32 v23, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v23, s5 :: v_dual_mov_b32 v22, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v21, s3 :: v_dual_mov_b32 v20, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v19, s1 :: v_dual_mov_b32 v18, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v13
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v9, s7 :: v_dual_mov_b32 v8, s6
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v7, s5 :: v_dual_mov_b32 v6, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v13, s11 :: v_dual_mov_b32 v12, s10
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v11, s9 :: v_dual_mov_b32 v10, s8
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v17, s15 :: v_dual_mov_b32 v16, s14
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v15, s13 :: v_dual_mov_b32 v14, s12
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_bf6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_bf6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_bf6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_bf6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.bf6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s0 :: v_dual_mov_b32 v20, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s2 :: v_dual_mov_b32 v22, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s4 :: v_dual_mov_b32 v24, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s5 :: v_dual_mov_b32 v23, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s3 :: v_dual_mov_b32 v21, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s1 :: v_dual_mov_b32 v19, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v18, s0 :: v_dual_mov_b32 v19, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v20, s2 :: v_dual_mov_b32 v21, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v22, s4 :: v_dual_mov_b32 v23, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v23, s5 :: v_dual_mov_b32 v22, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v21, s3 :: v_dual_mov_b32 v20, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v19, s1 :: v_dual_mov_b32 v18, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v13
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v9, s7 :: v_dual_mov_b32 v8, s6
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v7, s5 :: v_dual_mov_b32 v6, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v13, s11 :: v_dual_mov_b32 v12, s10
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v11, s9 :: v_dual_mov_b32 v10, s8
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v17, s15 :: v_dual_mov_b32 v16, s14
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v15, s13 :: v_dual_mov_b32 v14, s12
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_bf16_fp6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_bf16_fp6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_bf16_fp6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_bf16_fp6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x bfloat> @llvm.amdgcn.cvt.scale.pk32.bf16.fp6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x bfloat> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s0 :: v_dual_mov_b32 v20, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s2 :: v_dual_mov_b32 v22, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s4 :: v_dual_mov_b32 v24, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s5 :: v_dual_mov_b32 v23, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s3 :: v_dual_mov_b32 v21, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s1 :: v_dual_mov_b32 v19, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v18, s0 :: v_dual_mov_b32 v19, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v20, s2 :: v_dual_mov_b32 v21, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v22, s4 :: v_dual_mov_b32 v23, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v23, s5 :: v_dual_mov_b32 v22, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v21, s3 :: v_dual_mov_b32 v20, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v19, s1 :: v_dual_mov_b32 v18, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v13
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v9, s7 :: v_dual_mov_b32 v8, s6
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v7, s5 :: v_dual_mov_b32 v6, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v13, s11 :: v_dual_mov_b32 v12, s10
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v11, s9 :: v_dual_mov_b32 v10, s8
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v17, s15 :: v_dual_mov_b32 v16, s14
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v15, s13 :: v_dual_mov_b32 v14, s12
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_bf6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_bf6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_bf6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_bf6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.bf6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[9:24], v[2:7], v8 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s0 :: v_dual_mov_b32 v20, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s2 :: v_dual_mov_b32 v22, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s4 :: v_dual_mov_b32 v24, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s5 :: v_dual_mov_b32 v23, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s3 :: v_dual_mov_b32 v21, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s1 :: v_dual_mov_b32 v19, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[3:18], v[19:24], v2 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 %scale, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v18, s0 :: v_dual_mov_b32 v19, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v20, s2 :: v_dual_mov_b32 v21, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v22, s4 :: v_dual_mov_b32 v23, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v23, s5 :: v_dual_mov_b32 v22, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v21, s3 :: v_dual_mov_b32 v20, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v19, s1 :: v_dual_mov_b32 v18, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[2:17], v[18:23], 0x64 scale_sel:1
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v13
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v9, s7 :: v_dual_mov_b32 v8, s6
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v7, s5 :: v_dual_mov_b32 v6, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v13, s11 :: v_dual_mov_b32 v12, s10
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v11, s9 :: v_dual_mov_b32 v10, s8
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v17, s15 :: v_dual_mov_b32 v16, s14
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v15, s13 :: v_dual_mov_b32 v14, s12
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 100, i32 1)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f16_fp6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f16_fp6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x3
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f16_fp6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f16_fp6 v[8:23], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x3
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x half> @llvm.amdgcn.cvt.scale.pk32.f16.fp6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x half> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[9:40], v[2:7], v8 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[37:40], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[33:36], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[29:32], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[25:28], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[9:40], v[2:7], v8 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[25:28], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[29:32], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[33:36], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[37:40], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 %scale, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v35, s0 :: v_dual_mov_b32 v36, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v37, s2 :: v_dual_mov_b32 v38, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v39, s4 :: v_dual_mov_b32 v40, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[3:34], v[35:40], v2 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[31:34], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[27:30], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[23:26], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[19:22], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v40, s5 :: v_dual_mov_b32 v39, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v38, s3 :: v_dual_mov_b32 v37, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v36, s1 :: v_dual_mov_b32 v35, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[3:34], v[35:40], v2 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[19:22], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[23:26], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[27:30], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[31:34], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 %scale, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], 0x64 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], 0x64 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 100, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v34, s0 :: v_dual_mov_b32 v35, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v36, s2 :: v_dual_mov_b32 v37, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v38, s4 :: v_dual_mov_b32 v39, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[2:33], v[34:39], 0x64 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[30:33], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[26:29], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[22:25], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[18:21], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v39, s5 :: v_dual_mov_b32 v38, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v37, s3 :: v_dual_mov_b32 v36, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v35, s1 :: v_dual_mov_b32 v34, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[2:33], v[34:39], 0x64 scale_sel:2
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v13
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s8 :: v_dual_mov_b32 v7, s9
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s10 :: v_dual_mov_b32 v9, s11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s16, v18
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s17, v19
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s18, v20
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s19, v21
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s20, v22
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s21, v23
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s22, v24
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s23, v25
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s24, v26
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s25, v27
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s26, v28
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s27, v29
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s28, v30
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s29, v31
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s30, v32
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s31, v33
+; GFX13-GISEL-NEXT: s_clause 0x1
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s4 :: v_dual_mov_b32 v3, s5
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s6 :: v_dual_mov_b32 v5, s7
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s12 :: v_dual_mov_b32 v7, s13
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s14 :: v_dual_mov_b32 v9, s15
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s16 :: v_dual_mov_b32 v11, s17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s18 :: v_dual_mov_b32 v13, s19
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s20 :: v_dual_mov_b32 v15, s21
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s22 :: v_dual_mov_b32 v17, s23
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s24 :: v_dual_mov_b32 v19, s25
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s26 :: v_dual_mov_b32 v21, s27
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s28 :: v_dual_mov_b32 v23, s29
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s30 :: v_dual_mov_b32 v25, s31
+; GFX13-GISEL-NEXT: s_clause 0x5
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[18:21], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[22:25], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 100, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_bf6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_bf6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_bf6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_bf6 v[8:39], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.bf6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_vv(ptr addrspace(1) %out, <6 x i32> %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_vv:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[9:40], v[2:7], v8 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[37:40], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[33:36], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[29:32], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[25:28], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_vv:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[9:40], v[2:7], v8 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[9:12], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[13:16], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[17:20], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[21:24], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[25:28], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[29:32], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[33:36], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[37:40], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 %scale, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_vv_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src, i32 %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_vv_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v35, s0 :: v_dual_mov_b32 v36, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v37, s2 :: v_dual_mov_b32 v38, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v39, s4 :: v_dual_mov_b32 v40, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[3:34], v[35:40], v2 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[31:34], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[27:30], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[23:26], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[19:22], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_vv_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v40, s5 :: v_dual_mov_b32 v39, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v38, s3 :: v_dual_mov_b32 v37, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v36, s1 :: v_dual_mov_b32 v35, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[3:34], v[35:40], v2 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[3:6], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[7:10], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[11:14], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[15:18], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[19:22], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[23:26], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[27:30], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[31:34], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 %scale, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_sl(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_sl:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], 0x64 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_sl:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], 0x64 scale_sel:2
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 100, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_sl_inreg_src(ptr addrspace(1) %out, <6 x i32> inreg %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_sl_inreg_src:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v34, s0 :: v_dual_mov_b32 v35, s1
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v36, s2 :: v_dual_mov_b32 v37, s3
+; GFX13-SDAG-NEXT: v_dual_mov_b32 v38, s4 :: v_dual_mov_b32 v39, s5
+; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[2:33], v[34:39], 0x64 scale_sel:2
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[30:33], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[26:29], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[22:25], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[18:21], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[14:17], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[10:13], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_sl_inreg_src:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v39, s5 :: v_dual_mov_b32 v38, s4
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v37, s3 :: v_dual_mov_b32 v36, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v35, s1 :: v_dual_mov_b32 v34, s0
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[2:33], v[34:39], 0x64 scale_sel:2
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v2
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v3
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v4
+; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4)
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v5
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s8, v6
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s9, v7
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s10, v8
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s11, v9
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s4, v10
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s5, v11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v12
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v13
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v5, s3 :: v_dual_mov_b32 v4, s2
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s12, v14
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s13, v15
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s14, v16
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s15, v17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s8 :: v_dual_mov_b32 v7, s9
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s10 :: v_dual_mov_b32 v9, s11
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s16, v18
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s17, v19
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s18, v20
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s19, v21
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s20, v22
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s21, v23
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s22, v24
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s23, v25
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s24, v26
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s25, v27
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s26, v28
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s27, v29
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s28, v30
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s29, v31
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s30, v32
+; GFX13-GISEL-NEXT: v_readfirstlane_b32 s31, v33
+; GFX13-GISEL-NEXT: s_clause 0x1
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:16
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s4 :: v_dual_mov_b32 v3, s5
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s6 :: v_dual_mov_b32 v5, s7
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s12 :: v_dual_mov_b32 v7, s13
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s14 :: v_dual_mov_b32 v9, s15
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s16 :: v_dual_mov_b32 v11, s17
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s18 :: v_dual_mov_b32 v13, s19
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s20 :: v_dual_mov_b32 v15, s21
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s22 :: v_dual_mov_b32 v17, s23
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s24 :: v_dual_mov_b32 v19, s25
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s26 :: v_dual_mov_b32 v21, s27
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s28 :: v_dual_mov_b32 v23, s29
+; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s30 :: v_dual_mov_b32 v25, s31
+; GFX13-GISEL-NEXT: s_clause 0x5
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[2:5], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[6:9], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[10:13], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[14:17], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[18:21], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[22:25], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 100, i32 2)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_vs(ptr addrspace(1) %out, <6 x i32> %src, i32 inreg %scale) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_vs:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], s0
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_vs:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], s0
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 %scale, i32 0)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_ps void @test_cvt_scale_pk32_f32_fp6_vi(ptr addrspace(1) %out, <6 x i32> %src) {
+; GFX13-SDAG-LABEL: test_cvt_scale_pk32_f32_fp6_vi:
+; GFX13-SDAG: ; %bb.0:
+; GFX13-SDAG-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], 1 scale_sel:15
+; GFX13-SDAG-NEXT: s_clause 0x7
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-SDAG-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-SDAG-NEXT: s_endpgm
+;
+; GFX13-GISEL-LABEL: test_cvt_scale_pk32_f32_fp6_vi:
+; GFX13-GISEL: ; %bb.0:
+; GFX13-GISEL-NEXT: v_cvt_scale_pk32_f32_fp6 v[8:39], v[2:7], 1 scale_sel:15
+; GFX13-GISEL-NEXT: s_clause 0x7
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[8:11], off
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[12:15], off offset:16
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[16:19], off offset:32
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[20:23], off offset:48
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[24:27], off offset:64
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[28:31], off offset:80
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[32:35], off offset:96
+; GFX13-GISEL-NEXT: global_store_b128 v[0:1], v[36:39], off offset:112
+; GFX13-GISEL-NEXT: s_endpgm
+ %cvt = tail call <32 x float> @llvm.amdgcn.cvt.scale.pk32.f32.fp6(<6 x i32> %src, i32 1, i32 15)
+ store <32 x float> %cvt, ptr addrspace(1) %out
+ ret void
+}
More information about the llvm-branch-commits
mailing list