665 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mcpu=gfx950 -amdgpu-mfma-vgpr-form < %s | FileCheck %s3 4target triple = "amdgcn-amd-amdhsa"5 6define void @test_rewrite_mfma_f32_16x16x32_f16(<8 x half> %arg0, <8 x half> %arg1, ptr addrspace(1) %ptr) #0 {7; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_f16:8; CHECK: ; %bb.0:9; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)10; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off11; CHECK-NEXT: s_waitcnt vmcnt(0)12; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[4:7], a[0:3]13; CHECK-NEXT: ;;#ASMSTART14; CHECK-NEXT: ; use a[0:3]15; CHECK-NEXT: ;;#ASMEND16; CHECK-NEXT: s_setpc_b64 s[30:31]17 %src2 = load <4 x float>, ptr addrspace(1) %ptr18 %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)19 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)20 ret void21}22 23define void @test_rewrite_mfma_f32_32x32x16_f16(<8 x half> %arg0, <8 x half> %arg1, ptr addrspace(1) %ptr) #0 {24; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_f16:25; CHECK: ; %bb.0:26; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)27; CHECK-NEXT: global_load_dwordx4 a[12:15], v[8:9], off offset:4828; CHECK-NEXT: global_load_dwordx4 a[8:11], v[8:9], off offset:3229; CHECK-NEXT: global_load_dwordx4 a[4:7], v[8:9], off offset:1630; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off31; CHECK-NEXT: s_waitcnt vmcnt(0)32; CHECK-NEXT: v_mfma_f32_32x32x16_f16 a[0:15], v[0:3], v[4:7], a[0:15]33; CHECK-NEXT: ;;#ASMSTART34; CHECK-NEXT: ; use a[0:15]35; CHECK-NEXT: ;;#ASMEND36; CHECK-NEXT: s_setpc_b64 s[30:31]37 %src2 = load <16 x float>, ptr addrspace(1) %ptr38 %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.f16(<8 x half> %arg0, <8 x half> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)39 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)40 ret void41}42 43define void @test_rewrite_mfma_i32_16x16x64_i8(<4 x i32> %arg0, <4 x i32> %arg1, ptr addrspace(1) %ptr) #0 {44; CHECK-LABEL: test_rewrite_mfma_i32_16x16x64_i8:45; CHECK: ; %bb.0:46; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)47; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off48; CHECK-NEXT: s_waitcnt vmcnt(0)49; CHECK-NEXT: v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], v[4:7], a[0:3]50; CHECK-NEXT: ;;#ASMSTART51; CHECK-NEXT: ; use a[0:3]52; CHECK-NEXT: ;;#ASMEND53; CHECK-NEXT: s_setpc_b64 s[30:31]54 %src2 = load <4 x i32>, ptr addrspace(1) %ptr55 %mai = call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x64.i8(<4 x i32> %arg0, <4 x i32> %arg1, <4 x i32> %src2, i32 0, i32 0, i32 0)56 call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)57 ret void58}59 60define void @test_rewrite_mfma_i32_32x32x32_i8(<4 x i32> %arg0, <4 x i32> %arg1, ptr addrspace(1) %ptr) #0 {61; CHECK-LABEL: test_rewrite_mfma_i32_32x32x32_i8:62; CHECK: ; %bb.0:63; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)64; CHECK-NEXT: global_load_dwordx4 a[12:15], v[8:9], off offset:4865; CHECK-NEXT: global_load_dwordx4 a[8:11], v[8:9], off offset:3266; CHECK-NEXT: global_load_dwordx4 a[4:7], v[8:9], off offset:1667; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off68; CHECK-NEXT: s_waitcnt vmcnt(0)69; CHECK-NEXT: v_mfma_i32_32x32x32_i8 a[0:15], v[0:3], v[4:7], a[0:15]70; CHECK-NEXT: ;;#ASMSTART71; CHECK-NEXT: ; use a[0:15]72; CHECK-NEXT: ;;#ASMEND73; CHECK-NEXT: s_setpc_b64 s[30:31]74 %src2 = load <16 x i32>, ptr addrspace(1) %ptr75 %mai = call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x32.i8(<4 x i32> %arg0, <4 x i32> %arg1, <16 x i32> %src2, i32 0, i32 0, i32 0)76 call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)77 ret void78}79 80define void @test_rewrite_mfma_f32_16x16x32_bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, ptr addrspace(1) %ptr) #0 {81; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_bf16:82; CHECK: ; %bb.0:83; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)84; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off85; CHECK-NEXT: s_waitcnt vmcnt(0)86; CHECK-NEXT: v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], v[4:7], a[0:3]87; CHECK-NEXT: ;;#ASMSTART88; CHECK-NEXT: ; use a[0:3]89; CHECK-NEXT: ;;#ASMEND90; CHECK-NEXT: s_setpc_b64 s[30:31]91 %src2 = load <4 x float>, ptr addrspace(1) %ptr92 %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)93 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)94 ret void95}96 97define void @test_rewrite_mfma_f32_32x32x16_bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, ptr addrspace(1) %ptr) #0 {98; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_bf16:99; CHECK: ; %bb.0:100; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)101; CHECK-NEXT: global_load_dwordx4 a[12:15], v[8:9], off offset:48102; CHECK-NEXT: global_load_dwordx4 a[8:11], v[8:9], off offset:32103; CHECK-NEXT: global_load_dwordx4 a[4:7], v[8:9], off offset:16104; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off105; CHECK-NEXT: s_waitcnt vmcnt(0)106; CHECK-NEXT: v_mfma_f32_32x32x16_bf16 a[0:15], v[0:3], v[4:7], a[0:15]107; CHECK-NEXT: ;;#ASMSTART108; CHECK-NEXT: ; use a[0:15]109; CHECK-NEXT: ;;#ASMEND110; CHECK-NEXT: s_setpc_b64 s[30:31]111 %src2 = load <16 x float>, ptr addrspace(1) %ptr112 %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)113 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)114 ret void115}116 117; TODO: Full cross product of src0/src1 sizes not tested118define void @test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {119; CHECK-LABEL: test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0:120; CHECK: ; %bb.0:121; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)122; CHECK-NEXT: global_load_dwordx4 a[0:3], v[18:19], off123; CHECK-NEXT: s_waitcnt vmcnt(0)124; CHECK-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[0,0,0]125; CHECK-NEXT: ;;#ASMSTART126; CHECK-NEXT: ; use a[0:3]127; CHECK-NEXT: ;;#ASMEND128; CHECK-NEXT: s_setpc_b64 s[30:31]129 %src2 = load <4 x float>, ptr addrspace(1) %ptr130 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2,131 i32 0, ; cbsz132 i32 0, ; blgp133 i32 0, i32 %scale0, i32 0, i32 %scale1)134 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)135 ret void136}137 138define void @test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, ptr addrspace(1) %ptr) #0 {139; CHECK-LABEL: test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0:140; CHECK: ; %bb.0:141; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)142; CHECK-NEXT: global_load_dwordx4 a[0:3], v[16:17], off143; CHECK-NEXT: s_waitcnt vmcnt(0)144; CHECK-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3]145; CHECK-NEXT: ;;#ASMSTART146; CHECK-NEXT: ; use a[0:3]147; CHECK-NEXT: ;;#ASMEND148; CHECK-NEXT: s_setpc_b64 s[30:31]149 %src2 = load <4 x float>, ptr addrspace(1) %ptr150 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2,151 i32 0, ; cbsz152 i32 0, ; blgp153 i32 0, i32 0, i32 0, i32 0)154 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)155 ret void156}157 158define void @test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {159; CHECK-LABEL: test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2:160; CHECK: ; %bb.0:161; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)162; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off163; CHECK-NEXT: s_waitcnt vmcnt(0)164; CHECK-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3], v12, v13 op_sel_hi:[0,0,0] cbsz:2 blgp:2165; CHECK-NEXT: ;;#ASMSTART166; CHECK-NEXT: ; use a[0:3]167; CHECK-NEXT: ;;#ASMEND168; CHECK-NEXT: s_setpc_b64 s[30:31]169 %src2 = load <4 x float>, ptr addrspace(1) %ptr170 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %src2,171 i32 2, ; cbsz172 i32 2, ; blgp173 i32 0, i32 %scale0, i32 0, i32 %scale1)174 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)175 ret void176}177 178define void @test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, ptr addrspace(1) %ptr) #0 {179; CHECK-LABEL: test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2:180; CHECK: ; %bb.0:181; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)182; CHECK-NEXT: global_load_dwordx4 a[0:3], v[12:13], off183; CHECK-NEXT: s_waitcnt vmcnt(0)184; CHECK-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3] cbsz:2 blgp:2185; CHECK-NEXT: ;;#ASMSTART186; CHECK-NEXT: ; use a[0:3]187; CHECK-NEXT: ;;#ASMEND188; CHECK-NEXT: s_setpc_b64 s[30:31]189 %src2 = load <4 x float>, ptr addrspace(1) %ptr190 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %src2,191 i32 2, ; cbsz192 i32 2, ; blgp193 i32 0, i32 0, i32 0, i32 0)194 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)195 ret void196}197 198define void @test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4(<4 x i32> %arg0, <4 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {199; CHECK-LABEL: test_rewrite_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4:200; CHECK: ; %bb.0:201; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)202; CHECK-NEXT: global_load_dwordx4 a[0:3], v[10:11], off203; CHECK-NEXT: s_waitcnt vmcnt(0)204; CHECK-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:7], a[0:3], v8, v9 op_sel_hi:[0,0,0] cbsz:4 blgp:4205; CHECK-NEXT: ;;#ASMSTART206; CHECK-NEXT: ; use a[0:3]207; CHECK-NEXT: ;;#ASMEND208; CHECK-NEXT: s_setpc_b64 s[30:31]209 %src2 = load <4 x float>, ptr addrspace(1) %ptr210 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2,211 i32 4, ; cbsz212 i32 4, ; blgp213 i32 0, i32 %scale0, i32 0, i32 %scale1)214 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)215 ret void216}217 218define void @test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4(<4 x i32> %arg0, <4 x i32> %arg1, ptr addrspace(1) %ptr) #0 {219; CHECK-LABEL: test_rewrite_mfma_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4:220; CHECK: ; %bb.0:221; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)222; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off223; CHECK-NEXT: s_waitcnt vmcnt(0)224; CHECK-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:7], a[0:3] cbsz:4 blgp:4225; CHECK-NEXT: ;;#ASMSTART226; CHECK-NEXT: ; use a[0:3]227; CHECK-NEXT: ;;#ASMEND228; CHECK-NEXT: s_setpc_b64 s[30:31]229 %src2 = load <4 x float>, ptr addrspace(1) %ptr230 %mai = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2,231 i32 4, ; cbsz232 i32 4, ; blgp233 i32 0, i32 0, i32 0, i32 0)234 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)235 ret void236}237 238 239define void @test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz0__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {240; CHECK-LABEL: test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz0__blgp0:241; CHECK: ; %bb.0:242; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)243; CHECK-NEXT: global_load_dwordx4 a[12:15], v[18:19], off offset:48244; CHECK-NEXT: global_load_dwordx4 a[8:11], v[18:19], off offset:32245; CHECK-NEXT: global_load_dwordx4 a[4:7], v[18:19], off offset:16246; CHECK-NEXT: global_load_dwordx4 a[0:3], v[18:19], off247; CHECK-NEXT: s_waitcnt vmcnt(0)248; CHECK-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], v16, v17 op_sel_hi:[0,0,0]249; CHECK-NEXT: ;;#ASMSTART250; CHECK-NEXT: ; use a[0:15]251; CHECK-NEXT: ;;#ASMEND252; CHECK-NEXT: s_setpc_b64 s[30:31]253 %src2 = load <16 x float>, ptr addrspace(1) %ptr254 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2,255 i32 0, ; cbsz256 i32 0, ; blgp257 i32 0, i32 %scale0, i32 0, i32 %scale1)258 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)259 ret void260}261 262define void @test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz0__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, ptr addrspace(1) %ptr) #0 {263; CHECK-LABEL: test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz0__blgp0:264; CHECK: ; %bb.0:265; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)266; CHECK-NEXT: global_load_dwordx4 a[12:15], v[16:17], off offset:48267; CHECK-NEXT: global_load_dwordx4 a[8:11], v[16:17], off offset:32268; CHECK-NEXT: global_load_dwordx4 a[4:7], v[16:17], off offset:16269; CHECK-NEXT: global_load_dwordx4 a[0:3], v[16:17], off270; CHECK-NEXT: s_waitcnt vmcnt(0)271; CHECK-NEXT: v_mfma_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15]272; CHECK-NEXT: ;;#ASMSTART273; CHECK-NEXT: ; use a[0:15]274; CHECK-NEXT: ;;#ASMEND275; CHECK-NEXT: s_setpc_b64 s[30:31]276 %src2 = load <16 x float>, ptr addrspace(1) %ptr277 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2,278 i32 0, ; cbsz279 i32 0, ; blgp280 i32 0, i32 0, i32 0, i32 0)281 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)282 ret void283}284 285define void @test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz2__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {286; CHECK-LABEL: test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz2__blgp2:287; CHECK: ; %bb.0:288; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)289; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48290; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32291; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16292; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off293; CHECK-NEXT: s_waitcnt vmcnt(0)294; CHECK-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:5], v[6:11], a[0:15], v12, v13 op_sel_hi:[0,0,0] cbsz:2 blgp:2295; CHECK-NEXT: ;;#ASMSTART296; CHECK-NEXT: ; use a[0:15]297; CHECK-NEXT: ;;#ASMEND298; CHECK-NEXT: s_setpc_b64 s[30:31]299 %src2 = load <16 x float>, ptr addrspace(1) %ptr300 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <16 x float> %src2,301 i32 2, ; cbsz302 i32 2, ; blgp303 i32 0, i32 %scale0, i32 0, i32 %scale1)304 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)305 ret void306}307 308define void @test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz2__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, ptr addrspace(1) %ptr) #0 {309; CHECK-LABEL: test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz2__blgp2:310; CHECK: ; %bb.0:311; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)312; CHECK-NEXT: global_load_dwordx4 a[12:15], v[12:13], off offset:48313; CHECK-NEXT: global_load_dwordx4 a[8:11], v[12:13], off offset:32314; CHECK-NEXT: global_load_dwordx4 a[4:7], v[12:13], off offset:16315; CHECK-NEXT: global_load_dwordx4 a[0:3], v[12:13], off316; CHECK-NEXT: s_waitcnt vmcnt(0)317; CHECK-NEXT: v_mfma_f32_32x32x64_f8f6f4 a[0:15], v[0:5], v[6:11], a[0:15] cbsz:2 blgp:2318; CHECK-NEXT: ;;#ASMSTART319; CHECK-NEXT: ; use a[0:15]320; CHECK-NEXT: ;;#ASMEND321; CHECK-NEXT: s_setpc_b64 s[30:31]322 %src2 = load <16 x float>, ptr addrspace(1) %ptr323 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <16 x float> %src2,324 i32 2, ; cbsz325 i32 2, ; blgp326 i32 0, i32 0, i32 0, i32 0)327 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)328 ret void329}330 331define void @test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz4__blgp4(<4 x i32> %arg0, <4 x i32> %arg1, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {332; CHECK-LABEL: test_rewrite_mfma_scale_f32_32x32x64_f8f6f4_0_0__cbsz4__blgp4:333; CHECK: ; %bb.0:334; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)335; CHECK-NEXT: global_load_dwordx4 a[12:15], v[10:11], off offset:48336; CHECK-NEXT: global_load_dwordx4 a[8:11], v[10:11], off offset:32337; CHECK-NEXT: global_load_dwordx4 a[4:7], v[10:11], off offset:16338; CHECK-NEXT: global_load_dwordx4 a[0:3], v[10:11], off339; CHECK-NEXT: s_waitcnt vmcnt(0)340; CHECK-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:3], v[4:7], a[0:15], v8, v9 op_sel_hi:[0,0,0] cbsz:4 blgp:4341; CHECK-NEXT: ;;#ASMSTART342; CHECK-NEXT: ; use a[0:15]343; CHECK-NEXT: ;;#ASMEND344; CHECK-NEXT: s_setpc_b64 s[30:31]345 %src2 = load <16 x float>, ptr addrspace(1) %ptr346 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2,347 i32 4, ; cbsz348 i32 4, ; blgp349 i32 0, i32 %scale0, i32 0, i32 %scale1)350 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)351 ret void352}353 354define void @test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz4__blgp4(<4 x i32> %arg0, <4 x i32> %arg1, ptr addrspace(1) %ptr) #0 {355; CHECK-LABEL: test_rewrite_mfma_f32_32x32x64_f8f6f4_0_0__cbsz4__blgp4:356; CHECK: ; %bb.0:357; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)358; CHECK-NEXT: global_load_dwordx4 a[12:15], v[8:9], off offset:48359; CHECK-NEXT: global_load_dwordx4 a[8:11], v[8:9], off offset:32360; CHECK-NEXT: global_load_dwordx4 a[4:7], v[8:9], off offset:16361; CHECK-NEXT: global_load_dwordx4 a[0:3], v[8:9], off362; CHECK-NEXT: s_waitcnt vmcnt(0)363; CHECK-NEXT: v_mfma_f32_32x32x64_f8f6f4 a[0:15], v[0:3], v[4:7], a[0:15] cbsz:4 blgp:4364; CHECK-NEXT: ;;#ASMSTART365; CHECK-NEXT: ; use a[0:15]366; CHECK-NEXT: ;;#ASMEND367; CHECK-NEXT: s_setpc_b64 s[30:31]368 %src2 = load <16 x float>, ptr addrspace(1) %ptr369 %mai = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2,370 i32 4, ; cbsz371 i32 4, ; blgp372 i32 0, i32 0, i32 0, i32 0)373 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)374 ret void375}376 377define void @test_rewrite_smfmac_f32_16x16x64_f16(<8 x half> %arg0, <16 x half> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {378; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x64_f16:379; CHECK: ; %bb.0:380; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)381; CHECK-NEXT: v_mov_b32_e32 v15, v14382; CHECK-NEXT: v_mov_b32_e32 v14, v13383; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off384; CHECK-NEXT: s_waitcnt vmcnt(0)385; CHECK-NEXT: v_smfmac_f32_16x16x64_f16 a[0:3], v[0:3], v[4:11], v12386; CHECK-NEXT: ;;#ASMSTART387; CHECK-NEXT: ; use a[0:3]388; CHECK-NEXT: ;;#ASMEND389; CHECK-NEXT: s_setpc_b64 s[30:31]390 %src2 = load <4 x float>, ptr addrspace(1) %ptr391 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.f16(<8 x half> %arg0, <16 x half> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)392 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)393 ret void394}395 396define void @test_rewrite_smfmac_f32_32x32x32_f16(<8 x half> %arg0, <16 x half> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {397; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x32_f16:398; CHECK: ; %bb.0:399; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)400; CHECK-NEXT: v_mov_b32_e32 v15, v14401; CHECK-NEXT: v_mov_b32_e32 v14, v13402; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48403; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32404; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16405; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off406; CHECK-NEXT: s_waitcnt vmcnt(0)407; CHECK-NEXT: v_smfmac_f32_32x32x32_f16 a[0:15], v[0:3], v[4:11], v12408; CHECK-NEXT: ;;#ASMSTART409; CHECK-NEXT: ; use a[0:15]410; CHECK-NEXT: ;;#ASMEND411; CHECK-NEXT: s_setpc_b64 s[30:31]412 %src2 = load <16 x float>, ptr addrspace(1) %ptr413 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.f16(<8 x half> %arg0, <16 x half> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)414 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)415 ret void416}417 418define void @test_rewrite_smfmac_f32_16x16x64_bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {419; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x64_bf16:420; CHECK: ; %bb.0:421; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)422; CHECK-NEXT: v_mov_b32_e32 v15, v14423; CHECK-NEXT: v_mov_b32_e32 v14, v13424; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off425; CHECK-NEXT: s_waitcnt vmcnt(0)426; CHECK-NEXT: v_smfmac_f32_16x16x64_bf16 a[0:3], v[0:3], v[4:11], v12427; CHECK-NEXT: ;;#ASMSTART428; CHECK-NEXT: ; use a[0:3]429; CHECK-NEXT: ;;#ASMEND430; CHECK-NEXT: s_setpc_b64 s[30:31]431 %src2 = load <4 x float>, ptr addrspace(1) %ptr432 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)433 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)434 ret void435}436 437define void @test_rewrite_smfmac_f32_32x32x32_bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {438; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x32_bf16:439; CHECK: ; %bb.0:440; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)441; CHECK-NEXT: v_mov_b32_e32 v15, v14442; CHECK-NEXT: v_mov_b32_e32 v14, v13443; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48444; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32445; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16446; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off447; CHECK-NEXT: s_waitcnt vmcnt(0)448; CHECK-NEXT: v_smfmac_f32_32x32x32_bf16 a[0:15], v[0:3], v[4:11], v12449; CHECK-NEXT: ;;#ASMSTART450; CHECK-NEXT: ; use a[0:15]451; CHECK-NEXT: ;;#ASMEND452; CHECK-NEXT: s_setpc_b64 s[30:31]453 %src2 = load <16 x float>, ptr addrspace(1) %ptr454 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)455 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)456 ret void457}458 459define void @test_rewrite_smfmac_i32_16x16x128_i8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {460; CHECK-LABEL: test_rewrite_smfmac_i32_16x16x128_i8:461; CHECK: ; %bb.0:462; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)463; CHECK-NEXT: v_mov_b32_e32 v15, v14464; CHECK-NEXT: v_mov_b32_e32 v14, v13465; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off466; CHECK-NEXT: s_waitcnt vmcnt(0)467; CHECK-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v12468; CHECK-NEXT: ;;#ASMSTART469; CHECK-NEXT: ; use a[0:3]470; CHECK-NEXT: ;;#ASMEND471; CHECK-NEXT: s_setpc_b64 s[30:31]472 %src2 = load <4 x i32>, ptr addrspace(1) %ptr473 %mai = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %src2, i32 %arg2, i32 0, i32 0)474 call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)475 ret void476}477 478define void @test_rewrite_smfmac_i32_32x32x64_i8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {479; CHECK-LABEL: test_rewrite_smfmac_i32_32x32x64_i8:480; CHECK: ; %bb.0:481; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)482; CHECK-NEXT: v_mov_b32_e32 v15, v14483; CHECK-NEXT: v_mov_b32_e32 v14, v13484; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48485; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32486; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16487; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off488; CHECK-NEXT: s_waitcnt vmcnt(0)489; CHECK-NEXT: v_smfmac_i32_32x32x64_i8 a[0:15], v[0:3], v[4:11], v12490; CHECK-NEXT: ;;#ASMSTART491; CHECK-NEXT: ; use a[0:15]492; CHECK-NEXT: ;;#ASMEND493; CHECK-NEXT: s_setpc_b64 s[30:31]494 %src2 = load <16 x i32>, ptr addrspace(1) %ptr495 %mai = call <16 x i32> @llvm.amdgcn.smfmac.i32.32x32x64.i8(<4 x i32> %arg0, <8 x i32> %arg1, <16 x i32> %src2, i32 %arg2, i32 0, i32 0)496 call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)497 ret void498}499 500define void @test_rewrite_smfmac_f32_16x16x128_bf8_bf8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {501; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x128_bf8_bf8:502; CHECK: ; %bb.0:503; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)504; CHECK-NEXT: v_mov_b32_e32 v15, v14505; CHECK-NEXT: v_mov_b32_e32 v14, v13506; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off507; CHECK-NEXT: s_waitcnt vmcnt(0)508; CHECK-NEXT: v_smfmac_f32_16x16x128_bf8_bf8 a[0:3], v[0:3], v[4:11], v12 cbsz:1 abid:2509; CHECK-NEXT: ;;#ASMSTART510; CHECK-NEXT: ; use a[0:3]511; CHECK-NEXT: ;;#ASMEND512; CHECK-NEXT: s_setpc_b64 s[30:31]513 %src2 = load <4 x float>, ptr addrspace(1) %ptr514 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x128.bf8.bf8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 1, i32 2)515 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)516 ret void517}518 519define void @test_rewrite_smfmac_f32_16x16x128_bf8_fp8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {520; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x128_bf8_fp8:521; CHECK: ; %bb.0:522; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)523; CHECK-NEXT: v_mov_b32_e32 v15, v14524; CHECK-NEXT: v_mov_b32_e32 v14, v13525; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off526; CHECK-NEXT: s_waitcnt vmcnt(0)527; CHECK-NEXT: v_smfmac_f32_16x16x128_bf8_fp8 a[0:3], v[0:3], v[4:11], v12 cbsz:1 abid:2528; CHECK-NEXT: ;;#ASMSTART529; CHECK-NEXT: ; use a[0:3]530; CHECK-NEXT: ;;#ASMEND531; CHECK-NEXT: s_setpc_b64 s[30:31]532 %src2 = load <4 x float>, ptr addrspace(1) %ptr533 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x128.bf8.fp8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 1, i32 2)534 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)535 ret void536}537 538define void @test_rewrite_smfmac_f32_16x16x128_fp8_bf8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {539; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x128_fp8_bf8:540; CHECK: ; %bb.0:541; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)542; CHECK-NEXT: v_mov_b32_e32 v15, v14543; CHECK-NEXT: v_mov_b32_e32 v14, v13544; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off545; CHECK-NEXT: s_waitcnt vmcnt(0)546; CHECK-NEXT: v_smfmac_f32_16x16x128_fp8_bf8 a[0:3], v[0:3], v[4:11], v12 cbsz:1 abid:2547; CHECK-NEXT: ;;#ASMSTART548; CHECK-NEXT: ; use a[0:3]549; CHECK-NEXT: ;;#ASMEND550; CHECK-NEXT: s_setpc_b64 s[30:31]551 %src2 = load <4 x float>, ptr addrspace(1) %ptr552 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x128.fp8.bf8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 1, i32 2)553 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)554 ret void555}556 557define void @test_rewrite_smfmac_f32_16x16x128_fp8_fp8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {558; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x128_fp8_fp8:559; CHECK: ; %bb.0:560; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)561; CHECK-NEXT: v_mov_b32_e32 v15, v14562; CHECK-NEXT: v_mov_b32_e32 v14, v13563; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off564; CHECK-NEXT: s_waitcnt vmcnt(0)565; CHECK-NEXT: v_smfmac_f32_16x16x128_fp8_fp8 a[0:3], v[0:3], v[4:11], v12 cbsz:1 abid:2566; CHECK-NEXT: ;;#ASMSTART567; CHECK-NEXT: ; use a[0:3]568; CHECK-NEXT: ;;#ASMEND569; CHECK-NEXT: s_setpc_b64 s[30:31]570 %src2 = load <4 x float>, ptr addrspace(1) %ptr571 %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x128.fp8.fp8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 1, i32 2)572 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)573 ret void574}575 576define void @test_rewrite_smfmac_f32_32x32x64_bf8_bf8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {577; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x64_bf8_bf8:578; CHECK: ; %bb.0:579; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)580; CHECK-NEXT: v_mov_b32_e32 v15, v14581; CHECK-NEXT: v_mov_b32_e32 v14, v13582; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48583; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32584; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16585; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off586; CHECK-NEXT: s_waitcnt vmcnt(0)587; CHECK-NEXT: v_smfmac_f32_32x32x64_bf8_bf8 a[0:15], v[0:3], v[4:11], v12 cbsz:1 abid:2588; CHECK-NEXT: ;;#ASMSTART589; CHECK-NEXT: ; use a[0:15]590; CHECK-NEXT: ;;#ASMEND591; CHECK-NEXT: s_setpc_b64 s[30:31]592 %src2 = load <16 x float>, ptr addrspace(1) %ptr593 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x64.bf8.bf8(<4 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 1, i32 2)594 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)595 ret void596}597 598define void @test_rewrite_smfmac_f32_32x32x64_bf8_fp8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {599; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x64_bf8_fp8:600; CHECK: ; %bb.0:601; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)602; CHECK-NEXT: v_mov_b32_e32 v15, v14603; CHECK-NEXT: v_mov_b32_e32 v14, v13604; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48605; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32606; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16607; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off608; CHECK-NEXT: s_waitcnt vmcnt(0)609; CHECK-NEXT: v_smfmac_f32_32x32x64_bf8_fp8 a[0:15], v[0:3], v[4:11], v12 cbsz:1 abid:2610; CHECK-NEXT: ;;#ASMSTART611; CHECK-NEXT: ; use a[0:15]612; CHECK-NEXT: ;;#ASMEND613; CHECK-NEXT: s_setpc_b64 s[30:31]614 %src2 = load <16 x float>, ptr addrspace(1) %ptr615 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x64.bf8.fp8(<4 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 1, i32 2)616 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)617 ret void618}619 620define void @test_rewrite_smfmac_f32_32x32x64_fp8_bf8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {621; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x64_fp8_bf8:622; CHECK: ; %bb.0:623; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)624; CHECK-NEXT: v_mov_b32_e32 v15, v14625; CHECK-NEXT: v_mov_b32_e32 v14, v13626; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48627; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32628; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16629; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off630; CHECK-NEXT: s_waitcnt vmcnt(0)631; CHECK-NEXT: v_smfmac_f32_32x32x64_fp8_bf8 a[0:15], v[0:3], v[4:11], v12 cbsz:1 abid:2632; CHECK-NEXT: ;;#ASMSTART633; CHECK-NEXT: ; use a[0:15]634; CHECK-NEXT: ;;#ASMEND635; CHECK-NEXT: s_setpc_b64 s[30:31]636 %src2 = load <16 x float>, ptr addrspace(1) %ptr637 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x64.fp8.bf8(<4 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 1, i32 2)638 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)639 ret void640}641 642define void @test_rewrite_smfmac_f32_32x32x64_fp8_fp8(<4 x i32> %arg0, <8 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {643; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x64_fp8_fp8:644; CHECK: ; %bb.0:645; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)646; CHECK-NEXT: v_mov_b32_e32 v15, v14647; CHECK-NEXT: v_mov_b32_e32 v14, v13648; CHECK-NEXT: global_load_dwordx4 a[12:15], v[14:15], off offset:48649; CHECK-NEXT: global_load_dwordx4 a[8:11], v[14:15], off offset:32650; CHECK-NEXT: global_load_dwordx4 a[4:7], v[14:15], off offset:16651; CHECK-NEXT: global_load_dwordx4 a[0:3], v[14:15], off652; CHECK-NEXT: s_waitcnt vmcnt(0)653; CHECK-NEXT: v_smfmac_f32_32x32x64_fp8_fp8 a[0:15], v[0:3], v[4:11], v12 cbsz:1 abid:2654; CHECK-NEXT: ;;#ASMSTART655; CHECK-NEXT: ; use a[0:15]656; CHECK-NEXT: ;;#ASMEND657; CHECK-NEXT: s_setpc_b64 s[30:31]658 %src2 = load <16 x float>, ptr addrspace(1) %ptr659 %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x64.fp8.fp8(<4 x i32> %arg0, <8 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 1, i32 2)660 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)661 ret void662}663 664attributes #0 = { nounwind }665