llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Zach Goldthorpe (zGoldthorpe) <details> <summary>Changes</summary> Fixes LCOMPILER-3009. --- Full diff: https://github.com/llvm/llvm-project/pull/233883.diff 6 Files Affected: - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx1250-err.cl (+32) - (modified) llvm/lib/Target/AMDGPU/GCNFeatures.td (-3) - (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll (+1-1) - (modified) llvm/test/MC/AMDGPU/gfx1250-strict_err.s (+44) - (modified) llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s (+35) - (modified) llvm/test/MC/AMDGPU/gfx1251_err.s (+55) ``````````diff diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx1250-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx1250-err.cl new file mode 100644 index 00000000000000..5b62f7b5a66265 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx1250-err.cl @@ -0,0 +1,32 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.50-unknown-unknown -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu12.50s-unknown-unknown -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu12.51-unknown-unknown -verify -S -o - %s + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef int v4i __attribute__((ext_vector_type(4))); +typedef float v8f __attribute__((ext_vector_type(8))); +typedef half v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); +typedef int v8i __attribute__((ext_vector_type(8))); +typedef half v16h __attribute__((ext_vector_type(16))); +typedef short v16s __attribute__((ext_vector_type(16))); + +void test_amdgcn_swmmac_w32(global v8f* out8f, global v8h* out8h, global v8s* out8s, global v8i* out8i, + v8h a8h, v8s a8s, v2i a2i, int ai, + v16h b16h, v16s b16s, v4i b4i, v2i b2i, + v8f c8f, v8h c8h, v8s c8s, v8i c8i, + int index) +{ + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32(a8h, b16h, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32(a8s, b16s, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8h = __builtin_amdgcn_swmmac_f16_16x16x32_f16_w32(a8h, b16h, c8h, index); // expected-error{{'__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8s = __builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32(a8s, b16s, c8s, index); // expected-error{{'__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8i = __builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32(true, a2i, true, b4i, c8i, index, true); // expected-error{{'__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8i = __builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32(true, ai, true, b2i, c8i, index, true); // expected-error{{'__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8i = __builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32(true, a2i, true, b4i, c8i, index, true); // expected-error{{'__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32(a2i, b4i, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32(a2i, b4i, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32(a2i, b4i, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} + *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32(a2i, b4i, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} +} diff --git a/llvm/lib/Target/AMDGPU/GCNFeatures.td b/llvm/lib/Target/AMDGPU/GCNFeatures.td index 23c79f5949122b..67a21f0357a72b 100644 --- a/llvm/lib/Target/AMDGPU/GCNFeatures.td +++ b/llvm/lib/Target/AMDGPU/GCNFeatures.td @@ -2587,7 +2587,6 @@ def FeatureISAVersion12_50_STRICT : FeatureSet< FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, - FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureTransCoexecutionHazard, FeatureGFX1250VALUBlockingCycles, @@ -2626,7 +2625,6 @@ def FeatureISAVersion12_50 : FeatureSet< FeatureCvtPkNormVOP3Insts, FeatureBlock16ConversionScaleInsts, FeatureFP8E5M3Insts, - FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureTransCoexecutionHazard, FeatureGFX1250VALUBlockingCycles, @@ -2669,7 +2667,6 @@ def FeatureISAVersion12_51 : FeatureSet< FeatureCvtPkNormVOP3Insts, FeatureBlock16ConversionScaleInsts, FeatureFP8E5M3Insts, - FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureUseAddPC64Inst, FeatureTransCoexecutionHazard, diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll index e7ddfc73a8d60f..5bac28fa5f116e 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll @@ -1,7 +1,7 @@ ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py ; RUN: llc -mtriple=amdgpu12.00 -misched-cluster=0 < %s | FileCheck -check-prefix=GCN %s ; RUN: llc -mtriple=amdgpu12.00 -misched-cluster=0 -amdgpu-igrouplp-exact-solver-max-branches=250000 < %s | FileCheck -check-prefix=EXACTCUTOFF %s -; RUN: llc -mtriple=amdgpu12.50 -amdgpu-sched-strategy=coexec -enable-post-misched=0 < %s | FileCheck -check-prefix=COEXEC %s +; RUN: llc -mtriple=amdgpu12.50 -mattr=+swmmac-gfx1200-insts -amdgpu-sched-strategy=coexec -enable-post-misched=0 < %s | FileCheck -check-prefix=COEXEC %s declare <8 x half> @llvm.amdgcn.swmmac.f16.16x16x32.f16.v8f16.v8f16.v16f16..i16(<8 x half>, <16 x half>, <8 x half>, i16) diff --git a/llvm/test/MC/AMDGPU/gfx1250-strict_err.s b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s index 17c7337c8cb25f..bb74fc0766894e 100644 --- a/llvm/test/MC/AMDGPU/gfx1250-strict_err.s +++ b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s @@ -35,6 +35,50 @@ v_wmma_f32_16x16x128_bf8_fp8 v[16:23], v[0:15], v[8:23], v[16:23] v_wmma_f32_16x16x128_bf8_bf8 v[16:23], v[0:15], v[8:23], v[16:23] // GFX1250S-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_wmma_f32_16x16x128_bf8_bf8 +v_swmmac_f32_16x16x32_f16 v[12:19], v[0:3], v[4:11], v20 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_f16 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_f16 + +v_swmmac_f32_16x16x32_bf16 v[12:19], v[0:3], v[4:11], v20 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_bf16 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_bf16 + +v_swmmac_f16_16x16x32_f16 v[12:15], v[0:3], v[4:11], v16 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f16_16x16x32_f16 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f16_16x16x32_f16 + +v_swmmac_bf16_16x16x32_bf16 v[12:15], v[0:3], v[4:11], v16 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_bf16_16x16x32_bf16 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_bf16_16x16x32_bf16 + +v_swmmac_i32_16x16x32_iu8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_i32_16x16x32_iu8 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_i32_16x16x32_iu8 + +v_swmmac_i32_16x16x32_iu4 v[3:10], v0, v[1:2], v11 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_i32_16x16x32_iu4 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_i32_16x16x32_iu4 + +v_swmmac_i32_16x16x64_iu4 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_i32_16x16x64_iu4 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_i32_16x16x64_iu4 + +v_swmmac_f32_16x16x32_fp8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_fp8_fp8 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_fp8_fp8 + +v_swmmac_f32_16x16x32_fp8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_fp8_bf8 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_fp8_bf8 + +v_swmmac_f32_16x16x32_bf8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_bf8_fp8 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_bf8_fp8 + +v_swmmac_f32_16x16x32_bf8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250): v_swmmac_f32_16x16x32_bf8_bf8 +// GFX1250S-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_swmmac_f32_16x16x32_bf8_bf8 + v_cvt_scale_pk16_f16_fp6 v[10:17], v[20:22], v8 scale_sel:8 // GFX1250-ERR: :[[@LINE-1]]:49: error: scale_sel maximum supported value is 7 // GFX1250S-ERR: :[[@LINE-2]]:49: error: scale_sel maximum supported value is 3 diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s b/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s index 4d91c7ddb0b662..36e211065f4d61 100644 --- a/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s +++ b/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s @@ -17,6 +17,41 @@ v_dual_dot2acc_f32_bf16 v0, v1, v2 :: v_dual_add_f32 v3, v4, v5 v_dual_add_f32 v0, v1, v2 :: v_dual_dot2acc_f32_bf16 v3, v4, v5 // GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode +;; GFX12 SWMMAC instructions + +v_swmmac_f32_16x16x32_f16 v[12:19], v[0:3], v[4:11], v20 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f32_16x16x32_bf16 v[12:19], v[0:3], v[4:11], v20 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f16_16x16x32_f16 v[12:15], v[0:3], v[4:11], v16 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_bf16_16x16x32_bf16 v[12:15], v[0:3], v[4:11], v16 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_i32_16x16x32_iu8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_i32_16x16x32_iu4 v[3:10], v0, v[1:2], v11 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_i32_16x16x64_iu4 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f32_16x16x32_fp8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f32_16x16x32_fp8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f32_16x16x32_bf8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + +v_swmmac_f32_16x16x32_bf8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU + ;; DOT4_F32_*, DOT2_F32_*, DOT2_F16 and DOT2_BF16 v_dot4_f32_fp8_fp8 v0, v1, v2, v3 diff --git a/llvm/test/MC/AMDGPU/gfx1251_err.s b/llvm/test/MC/AMDGPU/gfx1251_err.s index a44dc08ac6e76b..474954a294d587 100644 --- a/llvm/test/MC/AMDGPU/gfx1251_err.s +++ b/llvm/test/MC/AMDGPU/gfx1251_err.s @@ -286,3 +286,58 @@ v_pk_lshl_add_u64 v[4:7], v[8:11], v[12:13], v[16:19] neg_lo:[1,0,0] v_pk_lshl_add_u64 v[4:7], v[17:20], v[6:7], null // GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX1251-ERR: v_pk_lshl_add_u64 v[4:7], v[17:20], v[6:7], null + +v_swmmac_f32_16x16x32_f16 v[12:19], v[0:3], v[4:11], v20 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_f16 +// GFX1251-ERR: v_swmmac_f32_16x16x32_f16 v[12:19], v[0:3], v[4:11], v20 +// GFX1251-ERR: ^ + +v_swmmac_f32_16x16x32_bf16 v[12:19], v[0:3], v[4:11], v20 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_bf16 +// GFX1251-ERR: v_swmmac_f32_16x16x32_bf16 v[12:19], v[0:3], v[4:11], v20 +// GFX1251-ERR: ^ + +v_swmmac_f16_16x16x32_f16 v[12:15], v[0:3], v[4:11], v16 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f16_16x16x32_f16 +// GFX1251-ERR: v_swmmac_f16_16x16x32_f16 v[12:15], v[0:3], v[4:11], v16 +// GFX1251-ERR: ^ + +v_swmmac_bf16_16x16x32_bf16 v[12:15], v[0:3], v[4:11], v16 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_bf16_16x16x32_bf16 +// GFX1251-ERR: v_swmmac_bf16_16x16x32_bf16 v[12:15], v[0:3], v[4:11], v16 +// GFX1251-ERR: ^ + +v_swmmac_i32_16x16x32_iu8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_i32_16x16x32_iu8 +// GFX1251-ERR: v_swmmac_i32_16x16x32_iu8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ + +v_swmmac_i32_16x16x32_iu4 v[3:10], v0, v[1:2], v11 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_i32_16x16x32_iu4 +// GFX1251-ERR: v_swmmac_i32_16x16x32_iu4 v[3:10], v0, v[1:2], v11 +// GFX1251-ERR: ^ + +v_swmmac_i32_16x16x64_iu4 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_i32_16x16x64_iu4 +// GFX1251-ERR: v_swmmac_i32_16x16x64_iu4 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ + +v_swmmac_f32_16x16x32_fp8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_fp8_fp8 +// GFX1251-ERR: v_swmmac_f32_16x16x32_fp8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ + +v_swmmac_f32_16x16x32_fp8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_fp8_bf8 +// GFX1251-ERR: v_swmmac_f32_16x16x32_fp8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ + +v_swmmac_f32_16x16x32_bf8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_bf8_fp8 +// GFX1251-ERR: v_swmmac_f32_16x16x32_bf8_fp8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ + +v_swmmac_f32_16x16x32_bf8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU (gfx1251): v_swmmac_f32_16x16x32_bf8_bf8 +// GFX1251-ERR: v_swmmac_f32_16x16x32_bf8_bf8 v[6:13], v[0:1], v[2:5], v14 +// GFX1251-ERR: ^ `````````` </details> https://github.com/llvm/llvm-project/pull/233883 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
