[llvm-branch-commits] [clang] [llvm] [AMDGPU] Add intrinsics and builtins for v_cvt_scale_pk32_* instructions (PR #222617)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 10 04:34:57 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Mariusz Sikora (mariusz-sikora-at-amd)
<details>
<summary>Changes</summary>
---
Patch is 100.86 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/222617.diff
13 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+24)
- (modified) clang/include/clang/Basic/BuiltinsAMDGPUDocs.td (+60)
- (modified) clang/lib/Sema/SemaAMDGPU.cpp (+6)
- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl (+11-1)
- (added) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-err.cl (+26)
- (added) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w32-err.cl (+20)
- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl (+58)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+7-1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+2-1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp (+12)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+6)
- (modified) llvm/lib/Target/AMDGPU/VOP3Instructions.td (+6-6)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scale.pk32.gfx13.ll (+1308)
``````````diff
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_cv...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/222617
More information about the llvm-branch-commits
mailing list