142 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mcpu=gfx90a -amdgpu-mfma-vgpr-form < %s | FileCheck %s3 4target triple = "amdgcn-amd-amdhsa"5 6define void @test_rewrite_mfma_i32_32x32x8i8(i32 %arg0, i32 %arg1, ptr addrspace(1) %ptr) #0 {7; CHECK-LABEL: test_rewrite_mfma_i32_32x32x8i8:8; CHECK: ; %bb.0:9; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)10; CHECK-NEXT: global_load_dwordx4 a[12:15], v[2:3], off offset:4811; CHECK-NEXT: global_load_dwordx4 a[8:11], v[2:3], off offset:3212; CHECK-NEXT: global_load_dwordx4 a[4:7], v[2:3], off offset:1613; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off14; CHECK-NEXT: s_waitcnt vmcnt(0)15; CHECK-NEXT: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15]16; CHECK-NEXT: ;;#ASMSTART17; CHECK-NEXT: ; use a[0:15]18; CHECK-NEXT: ;;#ASMEND19; CHECK-NEXT: s_setpc_b64 s[30:31]20 %src2 = load <16 x i32>, ptr addrspace(1) %ptr21 %mai = call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x8i8(i32 %arg0, i32 %arg1, <16 x i32> %src2, i32 0, i32 0, i32 0)22 call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)23 ret void24}25 26define void @test_rewrite_mfma_i32_16x16x16i8(i32 %arg0, i32 %arg1, ptr addrspace(1) %ptr) #0 {27; CHECK-LABEL: test_rewrite_mfma_i32_16x16x16i8:28; CHECK: ; %bb.0:29; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)30; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off31; CHECK-NEXT: s_waitcnt vmcnt(0)32; CHECK-NEXT: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[0:3]33; CHECK-NEXT: ;;#ASMSTART34; CHECK-NEXT: ; use a[0:3]35; CHECK-NEXT: ;;#ASMEND36; CHECK-NEXT: s_setpc_b64 s[30:31]37 %src2 = load <4 x i32>, ptr addrspace(1) %ptr38 %mai = call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32 %arg0, i32 %arg1, <4 x i32> %src2, i32 0, i32 0, i32 0)39 call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)40 ret void41}42 43define void @test_rewrite_mfma_f32_32x32x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, ptr addrspace(1) %ptr) #0 {44; CHECK-LABEL: test_rewrite_mfma_f32_32x32x2bf16:45; CHECK: ; %bb.0:46; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)47; CHECK-NEXT: global_load_dwordx4 a[28:31], v[2:3], off offset:11248; CHECK-NEXT: global_load_dwordx4 a[24:27], v[2:3], off offset:9649; CHECK-NEXT: global_load_dwordx4 a[20:23], v[2:3], off offset:8050; CHECK-NEXT: global_load_dwordx4 a[16:19], v[2:3], off offset:6451; CHECK-NEXT: global_load_dwordx4 a[12:15], v[2:3], off offset:4852; CHECK-NEXT: global_load_dwordx4 a[8:11], v[2:3], off offset:3253; CHECK-NEXT: global_load_dwordx4 a[4:7], v[2:3], off offset:1654; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off55; CHECK-NEXT: s_waitcnt vmcnt(0)56; CHECK-NEXT: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[0:31]57; CHECK-NEXT: ;;#ASMSTART58; CHECK-NEXT: ; use a[0:31]59; CHECK-NEXT: ;;#ASMEND60; CHECK-NEXT: s_setpc_b64 s[30:31]61 %src2 = load <32 x float>, ptr addrspace(1) %ptr62 %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)63 call void asm sideeffect "; use $0", "a"(<32 x float> %mai)64 ret void65}66 67define void @test_rewrite_mfma_f32_16x16x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, ptr addrspace(1) %ptr) #0 {68; CHECK-LABEL: test_rewrite_mfma_f32_16x16x2bf16:69; CHECK: ; %bb.0:70; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)71; CHECK-NEXT: global_load_dwordx4 a[12:15], v[2:3], off offset:4872; CHECK-NEXT: global_load_dwordx4 a[8:11], v[2:3], off offset:3273; CHECK-NEXT: global_load_dwordx4 a[4:7], v[2:3], off offset:1674; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off75; CHECK-NEXT: s_waitcnt vmcnt(0)76; CHECK-NEXT: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[0:15]77; CHECK-NEXT: ;;#ASMSTART78; CHECK-NEXT: ; use a[0:15]79; CHECK-NEXT: ;;#ASMEND80; CHECK-NEXT: s_setpc_b64 s[30:31]81 %src2 = load <16 x float>, ptr addrspace(1) %ptr82 %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)83 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)84 ret void85}86 87define void @test_rewrite_mfma_f32_4x4x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, ptr addrspace(1) %ptr) #0 {88; CHECK-LABEL: test_rewrite_mfma_f32_4x4x2bf16:89; CHECK: ; %bb.0:90; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)91; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off92; CHECK-NEXT: s_waitcnt vmcnt(0)93; CHECK-NEXT: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[0:3]94; CHECK-NEXT: ;;#ASMSTART95; CHECK-NEXT: ; use a[0:3]96; CHECK-NEXT: ;;#ASMEND97; CHECK-NEXT: s_setpc_b64 s[30:31]98 %src2 = load <4 x float>, ptr addrspace(1) %ptr99 %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %arg0, <2 x i16> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)100 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)101 ret void102}103 104define void @test_rewrite_mfma_f32_32x32x4bf16(<2 x i16> %arg0, <2 x i16> %arg1, ptr addrspace(1) %ptr) #0 {105; CHECK-LABEL: test_rewrite_mfma_f32_32x32x4bf16:106; CHECK: ; %bb.0:107; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)108; CHECK-NEXT: global_load_dwordx4 a[12:15], v[2:3], off offset:48109; CHECK-NEXT: global_load_dwordx4 a[8:11], v[2:3], off offset:32110; CHECK-NEXT: global_load_dwordx4 a[4:7], v[2:3], off offset:16111; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off112; CHECK-NEXT: s_waitcnt vmcnt(0)113; CHECK-NEXT: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[0:15]114; CHECK-NEXT: ;;#ASMSTART115; CHECK-NEXT: ; use a[0:15]116; CHECK-NEXT: ;;#ASMEND117; CHECK-NEXT: s_setpc_b64 s[30:31]118 %src2 = load <16 x float>, ptr addrspace(1) %ptr119 %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %arg0, <2 x i16> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)120 call void asm sideeffect "; use $0", "a"(<16 x float> %mai)121 ret void122}123 124define void @test_rewrite_mfma_f32_16x16x8bf16(<2 x i16> %arg0, <2 x i16> %arg1, ptr addrspace(1) %ptr) #0 {125; CHECK-LABEL: test_rewrite_mfma_f32_16x16x8bf16:126; CHECK: ; %bb.0:127; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)128; CHECK-NEXT: global_load_dwordx4 a[0:3], v[2:3], off129; CHECK-NEXT: s_waitcnt vmcnt(0)130; CHECK-NEXT: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[0:3]131; CHECK-NEXT: ;;#ASMSTART132; CHECK-NEXT: ; use a[0:3]133; CHECK-NEXT: ;;#ASMEND134; CHECK-NEXT: s_setpc_b64 s[30:31]135 %src2 = load <4 x float>, ptr addrspace(1) %ptr136 %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %arg0, <2 x i16> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)137 call void asm sideeffect "; use $0", "a"(<4 x float> %mai)138 ret void139}140 141attributes #0 = { nounwind "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="4,4" }142