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 <[email protected]> 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 +} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
