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