Author: Zach Goldthorpe Date: 2026-10-20T12:55:21-06:00 New Revision: a1c90088f87aad987924bdf4f06b1d13fe60c10c
URL: https://github.com/llvm/llvm-project/commit/a1c90088f87aad987924bdf4f06b1d13fe60c10c DIFF: https://github.com/llvm/llvm-project/commit/a1c90088f87aad987924bdf4f06b1d13fe60c10c.diff LOG: [AMDGPU][MC] Remove SWMMAC opcodes not supported by gfx1250 (#233883) Fixes LCOMPILER-3009. Added: clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx1250-err.cl llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx1250.ll Modified: llvm/lib/Target/AMDGPU/GCNFeatures.td llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll llvm/test/MC/AMDGPU/gfx1250-strict_err.s llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s llvm/test/MC/AMDGPU/gfx1251_err.s Removed: ################################################################################ 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 5937933bb63704..d92cca149fcc76 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..d44ef464049faa 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,6 @@ ; 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 declare <8 x half> @llvm.amdgcn.swmmac.f16.16x16x32.f16.v8f16.v8f16.v16f16..i16(<8 x half>, <16 x half>, <8 x half>, i16) @@ -101,58 +100,6 @@ define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_cluster(ptr ; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(5) SyncID(0) ; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(5) SyncID(0) ; EXACTCUTOFF-NEXT: s_endpgm -; -; COEXEC-LABEL: test_sched_group_barrier_pipeline_SWMMAC_cluster: -; COEXEC: ; %bb.0: ; %entry -; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; COEXEC-NEXT: s_mov_b64 s[64:65], 0 -; COEXEC-NEXT: v_nop -; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv -; COEXEC-NEXT: v_dual_mov_b32 v48, 0 :: v_dual_lshlrev_b32 v0, 4, v0 -; COEXEC-NEXT: s_mov_b32 s2, 0x3ff0 -; COEXEC-NEXT: s_wait_kmcnt 0x0 -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) -; COEXEC-NEXT: v_dual_mov_b32 v49, s1 :: v_dual_bitop2_b32 v0, s2, v0 bitop3:0x40 -; COEXEC-NEXT: v_dual_add_nc_u32 v4, s0, v0 :: v_dual_add_nc_u32 v50, s1, v0 -; COEXEC-NEXT: ds_load_b128 v[8:11], v4 -; COEXEC-NEXT: ds_load_b128 v[12:15], v4 offset:512 -; COEXEC-NEXT: ds_load_b128 v[16:19], v4 offset:5120 -; COEXEC-NEXT: ds_load_b128 v[20:23], v4 offset:3072 -; COEXEC-NEXT: ds_load_b128 v[24:27], v4 offset:1536 -; COEXEC-NEXT: ds_load_b128 v[0:3], v4 offset:11264 -; COEXEC-NEXT: ds_load_b128 v[4:7], v4 offset:11280 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(7) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x6 -; COEXEC-NEXT: v_mov_b64_e32 v[30:31], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[28:29], v[8:9] -; COEXEC-NEXT: s_wait_dscnt 0x5 -; COEXEC-NEXT: v_mov_b64_e32 v[34:35], v[14:15] -; COEXEC-NEXT: v_mov_b64_e32 v[32:33], v[12:13] -; COEXEC-NEXT: s_wait_dscnt 0x4 -; COEXEC-NEXT: v_mov_b64_e32 v[38:39], v[18:19] -; COEXEC-NEXT: v_mov_b64_e32 v[36:37], v[16:17] -; COEXEC-NEXT: s_wait_dscnt 0x3 -; COEXEC-NEXT: v_mov_b64_e32 v[42:43], v[22:23] -; COEXEC-NEXT: v_mov_b64_e32 v[40:41], v[20:21] -; COEXEC-NEXT: s_wait_dscnt 0x2 -; COEXEC-NEXT: v_mov_b64_e32 v[46:47], v[26:27] -; COEXEC-NEXT: v_mov_b64_e32 v[44:45], v[24:25] -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[28:31], v[8:11], v[0:7], v48 -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[32:35], v[12:15], v[0:7], v48 -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_3) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[44:47], v[24:27], v[0:7], v48 -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[40:43], v[20:23], v[0:7], v48 -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[36:39], v[16:19], v[0:7], v48 -; COEXEC-NEXT: ds_store_b128 v50, v[28:31] -; COEXEC-NEXT: ds_store_b128 v49, v[32:35] offset:512 -; COEXEC-NEXT: ds_store_b128 v49, v[44:47] offset:1024 -; COEXEC-NEXT: ds_store_b128 v49, v[40:43] offset:1536 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(5) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v49, v[36:39] offset:2048 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(5) SyncID(0) -; COEXEC-NEXT: s_endpgm entry: %idx = call i32 @llvm.amdgcn.workitem.id.x() %load.0.addr = getelementptr <8 x half>, ptr addrspace(3) %in, i32 %idx @@ -321,75 +268,6 @@ define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_interleaved( ; EXACTCUTOFF-NEXT: ds_store_b128 v16, v[12:15] offset:2048 ; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) ; EXACTCUTOFF-NEXT: s_endpgm -; -; COEXEC-LABEL: test_sched_group_barrier_pipeline_SWMMAC_interleaved: -; COEXEC: ; %bb.0: ; %entry -; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; COEXEC-NEXT: s_mov_b64 s[64:65], 0 -; COEXEC-NEXT: v_nop -; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv -; COEXEC-NEXT: s_mov_b32 s2, 0x3ff -; COEXEC-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2) -; COEXEC-NEXT: v_dual_mov_b32 v16, 0 :: v_dual_bitop2_b32 v0, s2, v0 bitop3:0x40 -; COEXEC-NEXT: s_wait_kmcnt 0x0 -; COEXEC-NEXT: v_mov_b32_e32 v17, s1 -; COEXEC-NEXT: v_lshl_add_u32 v18, v0, 5, s0 -; COEXEC-NEXT: v_lshl_add_u32 v19, v0, 4, s1 -; COEXEC-NEXT: ds_load_b128 v[8:11], v18 offset:1024 -; COEXEC-NEXT: ds_load_b128 v[0:3], v18 -; COEXEC-NEXT: ds_load_b128 v[4:7], v18 offset:16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(3) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x2 -; COEXEC-NEXT: v_mov_b64_e32 v[14:15], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[12:13], v[8:9] -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[12:15], v[8:11], v[0:7], v16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v19, v[12:15] -; COEXEC-NEXT: ds_load_b128 v[8:11], v18 offset:2560 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(1) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: v_mov_b64_e32 v[14:15], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[12:13], v[8:9] -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[12:15], v[8:11], v[0:7], v16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v17, v[12:15] offset:512 -; COEXEC-NEXT: ds_load_b128 v[8:11], v18 offset:4608 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(1) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: v_mov_b64_e32 v[14:15], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[12:13], v[8:9] -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[12:15], v[8:11], v[0:7], v16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v17, v[12:15] offset:1024 -; COEXEC-NEXT: ds_load_b128 v[8:11], v18 offset:7168 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(1) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: v_mov_b64_e32 v[14:15], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[12:13], v[8:9] -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[12:15], v[8:11], v[0:7], v16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v17, v[12:15] offset:1536 -; COEXEC-NEXT: ds_load_b128 v[8:11], v18 offset:10240 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(1) SyncID(0) -; COEXEC-NEXT: s_wait_dscnt 0x0 -; COEXEC-NEXT: v_mov_b64_e32 v[14:15], v[10:11] -; COEXEC-NEXT: v_mov_b64_e32 v[12:13], v[8:9] -; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) -; COEXEC-NEXT: v_swmmac_f16_16x16x32_f16 v[12:15], v[8:11], v[0:7], v16 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) -; COEXEC-NEXT: ds_store_b128 v17, v[12:15] offset:2048 -; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) -; COEXEC-NEXT: s_endpgm entry: %idx = call i32 @llvm.amdgcn.workitem.id.x() %load.b.addr = getelementptr <16 x half>, ptr addrspace(3) %in, i32 %idx diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx1250.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx1250.ll new file mode 100644 index 00000000000000..b5e09c368af4b2 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx1250.ll @@ -0,0 +1,255 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc -mtriple=amdgpu12.50 -amdgpu-sched-strategy=coexec -enable-post-misched=0 < %s | FileCheck -check-prefix=COEXEC %s + +define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_cluster(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 { +; COEXEC-LABEL: test_sched_group_barrier_pipeline_SWMMAC_cluster: +; COEXEC: ; %bb.0: ; %entry +; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; COEXEC-NEXT: s_mov_b64 s[64:65], 0 +; COEXEC-NEXT: v_nop +; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; COEXEC-NEXT: s_mov_b32 s2, 0x3ff +; COEXEC-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2) +; COEXEC-NEXT: v_dual_mov_b32 v76, 0 :: v_dual_bitop2_b32 v0, s2, v0 bitop3:0x40 +; COEXEC-NEXT: s_wait_kmcnt 0x0 +; COEXEC-NEXT: v_mov_b32_e32 v77, s1 +; COEXEC-NEXT: v_lshl_add_u32 v20, v0, 5, s0 +; COEXEC-NEXT: v_lshl_add_u32 v78, v0, 4, s1 +; COEXEC-NEXT: ds_load_b128 v[40:43], v20 +; COEXEC-NEXT: ds_load_b128 v[16:19], v20 offset:10240 +; COEXEC-NEXT: ds_load_b128 v[24:27], v20 offset:6144 +; COEXEC-NEXT: ds_load_b128 v[32:35], v20 offset:3072 +; COEXEC-NEXT: ds_load_b128 v[48:51], v20 offset:1024 +; COEXEC-NEXT: ds_load_b128 v[44:47], v20 offset:16 +; COEXEC-NEXT: ds_load_b128 v[0:3], v20 offset:22528 +; COEXEC-NEXT: ds_load_b128 v[52:55], v20 offset:1040 +; COEXEC-NEXT: ds_load_b128 v[4:7], v20 offset:22544 +; COEXEC-NEXT: ds_load_b128 v[36:39], v20 offset:3088 +; COEXEC-NEXT: ds_load_b128 v[8:11], v20 offset:22560 +; COEXEC-NEXT: ds_load_b128 v[28:31], v20 offset:6160 +; COEXEC-NEXT: ds_load_b128 v[12:15], v20 offset:22576 +; COEXEC-NEXT: ds_load_b128 v[20:23], v20 offset:10256 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(14) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0xd +; COEXEC-NEXT: v_mov_b64_e32 v[58:59], v[42:43] +; COEXEC-NEXT: v_mov_b64_e32 v[56:57], v[40:41] +; COEXEC-NEXT: s_wait_dscnt 0xc +; COEXEC-NEXT: v_mov_b64_e32 v[62:63], v[18:19] +; COEXEC-NEXT: v_mov_b64_e32 v[60:61], v[16:17] +; COEXEC-NEXT: s_wait_dscnt 0xb +; COEXEC-NEXT: v_mov_b64_e32 v[66:67], v[26:27] +; COEXEC-NEXT: v_mov_b64_e32 v[64:65], v[24:25] +; COEXEC-NEXT: s_wait_dscnt 0xa +; COEXEC-NEXT: v_mov_b64_e32 v[70:71], v[34:35] +; COEXEC-NEXT: v_mov_b64_e32 v[68:69], v[32:33] +; COEXEC-NEXT: s_wait_dscnt 0x9 +; COEXEC-NEXT: v_mov_b64_e32 v[74:75], v[50:51] +; COEXEC-NEXT: v_mov_b64_e32 v[72:73], v[48:49] +; COEXEC-NEXT: s_wait_dscnt 0x1 +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[56:59], v[40:47], v[0:15], v76 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[72:75], v[48:55], v[0:15], v76 +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[68:71], v[32:39], v[0:15], v76 +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[64:67], v[24:31], v[0:15], v76 +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[60:63], v[16:23], v[0:15], v76 +; COEXEC-NEXT: ds_store_b128 v78, v[56:59] +; COEXEC-NEXT: ds_store_b128 v77, v[72:75] offset:512 +; COEXEC-NEXT: ds_store_b128 v77, v[68:71] offset:1024 +; COEXEC-NEXT: ds_store_b128 v77, v[64:67] offset:1536 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(5) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v77, v[60:63] offset:2048 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(5) SyncID(0) +; COEXEC-NEXT: s_endpgm +entry: + %idx = call i32 @llvm.amdgcn.workitem.id.x() + %load.0.addr = getelementptr <16 x half>, ptr addrspace(3) %in, i32 %idx + %load.0 = load <16 x half>, ptr addrspace(3) %load.0.addr + %load.1.addr = getelementptr <16 x half>, ptr addrspace(3) %load.0.addr, i32 32 + %load.1 = load <16 x half>, ptr addrspace(3) %load.1.addr + %load.2.addr = getelementptr <16 x half>, ptr addrspace(3) %load.1.addr, i32 64 + %load.2 = load <16 x half>, ptr addrspace(3) %load.2.addr + %load.3.addr = getelementptr <16 x half>, ptr addrspace(3) %load.2.addr, i32 96 + %load.3 = load <16 x half>, ptr addrspace(3) %load.3.addr + %load.4.addr = getelementptr <16 x half>, ptr addrspace(3) %load.3.addr, i32 128 + %load.4 = load <16 x half>, ptr addrspace(3) %load.4.addr + %load.b.addr = getelementptr <32 x half>, ptr addrspace(3) %load.4.addr, i32 192 + %load.b = load <32 x half>, ptr addrspace(3) %load.b.addr + %c.0 = shufflevector <16 x half> %load.0, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.1 = shufflevector <16 x half> %load.1, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.2 = shufflevector <16 x half> %load.2, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.3 = shufflevector <16 x half> %load.3, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.4 = shufflevector <16 x half> %load.4, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %mai.0 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.0, i1 0, <32 x half> %load.b, <8 x half> %c.0, i16 0, i1 false, i1 false) + %mai.1 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.1, i1 0, <32 x half> %load.b, <8 x half> %c.1, i16 0, i1 false, i1 false) + %mai.2 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.2, i1 0, <32 x half> %load.b, <8 x half> %c.2, i16 0, i1 false, i1 false) + %mai.3 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.3, i1 0, <32 x half> %load.b, <8 x half> %c.3, i16 0, i1 false, i1 false) + %mai.4 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.4, i1 0, <32 x half> %load.b, <8 x half> %c.4, i16 0, i1 false, i1 false) + %store.0.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 %idx + store <8 x half> %mai.0, ptr addrspace(3) %store.0.addr + %store.1.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 32 + store <8 x half> %mai.1, ptr addrspace(3) %store.1.addr + %store.2.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 64 + store <8 x half> %mai.2, ptr addrspace(3) %store.2.addr + %store.3.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 96 + store <8 x half> %mai.3, ptr addrspace(3) %store.3.addr + %store.4.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 128 + store <8 x half> %mai.4, ptr addrspace(3) %store.4.addr + ; 14 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 14, i32 0) + ; 5 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 5, i32 0) + ; 5 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 5, i32 0) + ret void +} + +define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_interleaved(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 { +; COEXEC-LABEL: test_sched_group_barrier_pipeline_SWMMAC_interleaved: +; COEXEC: ; %bb.0: ; %entry +; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; COEXEC-NEXT: s_mov_b64 s[64:65], 0 +; COEXEC-NEXT: v_nop +; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; COEXEC-NEXT: s_mov_b32 s2, 0x3ff +; COEXEC-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2) +; COEXEC-NEXT: v_dual_mov_b32 v40, 0 :: v_dual_bitop2_b32 v0, s2, v0 bitop3:0x40 +; COEXEC-NEXT: s_wait_kmcnt 0x0 +; COEXEC-NEXT: v_mov_b32_e32 v41, s1 +; COEXEC-NEXT: v_lshl_add_u32 v42, v0, 6, s0 +; COEXEC-NEXT: v_lshl_add_u32 v28, v0, 4, s1 +; COEXEC-NEXT: ds_load_b128 v[16:19], v42 offset:2048 +; COEXEC-NEXT: ds_load_b128 v[0:3], v42 +; COEXEC-NEXT: ds_load_b128 v[4:7], v42 offset:16 +; COEXEC-NEXT: ds_load_b128 v[8:11], v42 offset:32 +; COEXEC-NEXT: ds_load_b128 v[12:15], v42 offset:48 +; COEXEC-NEXT: ds_load_b128 v[20:23], v42 offset:2064 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(6) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0x5 +; COEXEC-NEXT: v_mov_b64_e32 v[26:27], v[18:19] +; COEXEC-NEXT: v_mov_b64_e32 v[24:25], v[16:17] +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[24:27], v[16:23], v[0:15], v40 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v28, v[24:27] +; COEXEC-NEXT: ds_load_b128 v[24:27], v42 offset:5120 +; COEXEC-NEXT: ds_load_b128 v[28:31], v42 offset:5136 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(2) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0x1 +; COEXEC-NEXT: v_nop +; COEXEC-NEXT: v_nop +; COEXEC-NEXT: v_mov_b64_e32 v[34:35], v[26:27] +; COEXEC-NEXT: v_mov_b64_e32 v[32:33], v[24:25] +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[32:35], v[24:31], v[0:15], v40 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v41, v[32:35] offset:512 +; COEXEC-NEXT: ds_load_b128 v[16:19], v42 offset:9216 +; COEXEC-NEXT: ds_load_b128 v[20:23], v42 offset:9232 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(2) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0x1 +; COEXEC-NEXT: v_mov_b64_e32 v[38:39], v[18:19] +; COEXEC-NEXT: v_mov_b64_e32 v[36:37], v[16:17] +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[36:39], v[16:23], v[0:15], v40 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v41, v[36:39] offset:1024 +; COEXEC-NEXT: ds_load_b128 v[24:27], v42 offset:14336 +; COEXEC-NEXT: ds_load_b128 v[28:31], v42 offset:14352 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(2) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0x1 +; COEXEC-NEXT: v_mov_b64_e32 v[34:35], v[26:27] +; COEXEC-NEXT: v_mov_b64_e32 v[32:33], v[24:25] +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[32:35], v[24:31], v[0:15], v40 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v41, v[32:35] offset:1536 +; COEXEC-NEXT: ds_load_b128 v[16:19], v42 offset:20480 +; COEXEC-NEXT: ds_load_b128 v[20:23], v42 offset:20496 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000100) size(2) SyncID(0) +; COEXEC-NEXT: s_wait_dscnt 0x1 +; COEXEC-NEXT: v_mov_b64_e32 v[38:39], v[18:19] +; COEXEC-NEXT: v_mov_b64_e32 v[36:37], v[16:17] +; COEXEC-NEXT: s_wait_dscnt 0x0 +; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) +; COEXEC-NEXT: v_swmmac_f16_16x16x64_f16 v[36:39], v[16:23], v[0:15], v40 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0) +; COEXEC-NEXT: ds_store_b128 v41, v[36:39] offset:2048 +; COEXEC-NEXT: ; sched_group_barrier mask(0x00000200) size(1) SyncID(0) +; COEXEC-NEXT: s_endpgm +entry: + %idx = call i32 @llvm.amdgcn.workitem.id.x() + %load.b.addr = getelementptr <32 x half>, ptr addrspace(3) %in, i32 %idx + %load.b = load <32 x half>, ptr addrspace(3) %load.b.addr + %load.0.addr = getelementptr <16 x half>, ptr addrspace(3) %load.b.addr, i32 64 + %load.0 = load <16 x half>, ptr addrspace(3) %load.0.addr + %load.1.addr = getelementptr <16 x half>, ptr addrspace(3) %load.0.addr, i32 96 + %load.1 = load <16 x half>, ptr addrspace(3) %load.1.addr + %load.2.addr = getelementptr <16 x half>, ptr addrspace(3) %load.1.addr, i32 128 + %load.2 = load <16 x half>, ptr addrspace(3) %load.2.addr + %load.3.addr = getelementptr <16 x half>, ptr addrspace(3) %load.2.addr, i32 160 + %load.3 = load <16 x half>, ptr addrspace(3) %load.3.addr + %load.4.addr = getelementptr <16 x half>, ptr addrspace(3) %load.3.addr, i32 192 + %load.4 = load <16 x half>, ptr addrspace(3) %load.4.addr + %c.0 = shufflevector <16 x half> %load.0, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.1 = shufflevector <16 x half> %load.1, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.2 = shufflevector <16 x half> %load.2, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.3 = shufflevector <16 x half> %load.3, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %c.4 = shufflevector <16 x half> %load.4, <16 x half> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> + %mai.0 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.0, i1 0, <32 x half> %load.b, <8 x half> %c.0, i16 0, i1 false, i1 false) + %mai.1 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.1, i1 0, <32 x half> %load.b, <8 x half> %c.1, i16 0, i1 false, i1 false) + %mai.2 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.2, i1 0, <32 x half> %load.b, <8 x half> %c.2, i16 0, i1 false, i1 false) + %mai.3 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.3, i1 0, <32 x half> %load.b, <8 x half> %c.3, i16 0, i1 false, i1 false) + %mai.4 = call <8 x half> @llvm.amdgcn.swmmac.f16.16x16x64.f16.v8f16.v16f16.v32f16.i16(i1 0, <16 x half> %load.4, i1 0, <32 x half> %load.b, <8 x half> %c.4, i16 0, i1 false, i1 false) + %store.0.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 %idx + store <8 x half> %mai.0, ptr addrspace(3) %store.0.addr + %store.1.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 32 + store <8 x half> %mai.1, ptr addrspace(3) %store.1.addr + %store.2.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 64 + store <8 x half> %mai.2, ptr addrspace(3) %store.2.addr + %store.3.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 96 + store <8 x half> %mai.3, ptr addrspace(3) %store.3.addr + %store.4.addr = getelementptr <8 x half>, ptr addrspace(3) %out, i32 128 + store <8 x half> %mai.4, ptr addrspace(3) %store.4.addr + ; 6 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 6, i32 0) + ; 1 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0) + ; 1 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 1, i32 0) + ; 2 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 2, i32 0) + ; 1 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0) + ; 1 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 1, i32 0) + ; 2 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 2, i32 0) + ; 1 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0) + ; 1 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 1, i32 0) + ; 2 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 2, i32 0) + ; 1 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0) + ; 1 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 1, i32 0) + ; 2 DS read + call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 2, i32 0) + ; 1 SWMMAC + call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0) + ; 1 DS write + call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 1, i32 0) + ret void +} 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: ^ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
