11; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2- ; RUN: llc -O3 - mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942
3- ; RUN: llc -O3 - mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908
2+ ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942
3+ ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908
44
5- define amdgpu_kernel void @test ( i1 %arg , i1 %arg1 ) {
6- ; GFX942-LABEL: test :
7- ; GFX942: ; %bb.0: ; %bb
8- ; GFX942-NEXT: s_load_dword s0 , s[4:5], 0x0
5+ define amdgpu_kernel void @matmul_kernel ( i32 %a0 , i32 %a1 ) {
6+ ; GFX942-LABEL: matmul_kernel :
7+ ; GFX942: ; %bb.0: ; %entry
8+ ; GFX942-NEXT: s_load_dwordx2 s[0:1] , s[4:5], 0x0
99; GFX942-NEXT: v_mov_b32_e32 v1, 0
1010; GFX942-NEXT: s_mov_b32 s2, 0
1111; GFX942-NEXT: v_accvgpr_write_b32 a0, v1
1212; GFX942-NEXT: s_mov_b32 s3, 0
1313; GFX942-NEXT: s_waitcnt lgkmcnt(0)
14- ; GFX942-NEXT: s_bitcmp1_b32 s0, 0
14+ ; GFX942-NEXT: s_cmp_lg_u32 s0, 0
1515; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0
1616; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
1717; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
1818; GFX942-NEXT: s_branch .LBB0_2
19- ; GFX942-NEXT: .LBB0_1: ; %Flow
20- ; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
21- ; GFX942-NEXT: s_andn2_b64 vcc, exec, s[4:5]
22- ; GFX942-NEXT: s_cbranch_vccz .LBB0_4
23- ; GFX942-NEXT: .LBB0_2: ; %bb2
24- ; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
25- ; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1]
26- ; GFX942-NEXT: s_mov_b64 s[4:5], -1
27- ; GFX942-NEXT: s_cbranch_vccnz .LBB0_1
28- ; GFX942-NEXT: ; %bb.3: ; %bb4
19+ ; GFX942-NEXT: .LBB0_1: ; %bb2
2920; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
3021; GFX942-NEXT: s_or_b32 s4, s3, 1
3122; GFX942-NEXT: s_ashr_i32 s5, s3, 31
@@ -39,39 +30,37 @@ define amdgpu_kernel void @test(i1 %arg, i1 %arg1) {
3930; GFX942-NEXT: v_accvgpr_write_b32 a2, v2
4031; GFX942-NEXT: v_accvgpr_write_b32 a3, v3
4132; GFX942-NEXT: s_and_b32 s3, s5, s4
42- ; GFX942-NEXT: s_mov_b64 s[4:5], 0
33+ ; GFX942-NEXT: s_nop 0
4334; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3]
44- ; GFX942-NEXT: s_branch .LBB0_1
35+ ; GFX942-NEXT: s_cbranch_execz .LBB0_4
36+ ; GFX942-NEXT: .LBB0_2: ; %bb
37+ ; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
38+ ; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1]
39+ ; GFX942-NEXT: s_cbranch_vccz .LBB0_1
40+ ; GFX942-NEXT: ; %bb.3:
41+ ; GFX942-NEXT: ; implicit-def: $sgpr3
4542; GFX942-NEXT: .LBB0_4: ; %common.ret
4643; GFX942-NEXT: s_endpgm
4744;
48- ; GFX908-LABEL: test :
49- ; GFX908: ; %bb.0: ; %bb
50- ; GFX908-NEXT: s_load_dword s0 , s[8:9], 0x0
45+ ; GFX908-LABEL: matmul_kernel :
46+ ; GFX908: ; %bb.0: ; %entry
47+ ; GFX908-NEXT: s_load_dwordx2 s[0:1] , s[8:9], 0x0
5148; GFX908-NEXT: v_mov_b32_e32 v1, 0
5249; GFX908-NEXT: s_mov_b32 s2, 0
5350; GFX908-NEXT: s_mov_b32 s3, 0
5451; GFX908-NEXT: v_accvgpr_write_b32 a0, v1
5552; GFX908-NEXT: s_waitcnt lgkmcnt(0)
56- ; GFX908-NEXT: s_bitcmp1_b32 s0, 0
53+ ; GFX908-NEXT: s_cmp_lg_u32 s0, 0
5754; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0
5855; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
5956; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
6057; GFX908-NEXT: s_branch .LBB0_2
61- ; GFX908-NEXT: .LBB0_1: ; %Flow
62- ; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
63- ; GFX908-NEXT: s_andn2_b64 vcc, exec, s[4:5]
64- ; GFX908-NEXT: s_cbranch_vccz .LBB0_4
65- ; GFX908-NEXT: .LBB0_2: ; %bb2
66- ; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1
67- ; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
68- ; GFX908-NEXT: s_mov_b64 s[4:5], -1
69- ; GFX908-NEXT: s_cbranch_vccnz .LBB0_1
70- ; GFX908-NEXT: ; %bb.3: ; %bb4
58+ ; GFX908-NEXT: .LBB0_1: ; %bb2
7159; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
7260; GFX908-NEXT: s_or_b32 s4, s3, 1
7361; GFX908-NEXT: s_ashr_i32 s5, s3, 31
7462; GFX908-NEXT: s_mov_b32 s3, s2
63+ ; GFX908-NEXT: s_nop 3
7564; GFX908-NEXT: v_accvgpr_read_b32 v0, a0
7665; GFX908-NEXT: v_mov_b32_e32 v5, s3
7766; GFX908-NEXT: v_mov_b32_e32 v4, s2
@@ -83,42 +72,49 @@ define amdgpu_kernel void @test(i1 %arg, i1 %arg1) {
8372; GFX908-NEXT: v_accvgpr_write_b32 a3, v3
8473; GFX908-NEXT: s_and_b32 s3, s5, s4
8574; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3]
86- ; GFX908-NEXT: s_mov_b64 s[4:5], 0
87- ; GFX908-NEXT: s_branch .LBB0_1
75+ ; GFX908-NEXT: s_cbranch_execz .LBB0_4
76+ ; GFX908-NEXT: .LBB0_2: ; %bb
77+ ; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1
78+ ; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
79+ ; GFX908-NEXT: s_cbranch_vccz .LBB0_1
80+ ; GFX908-NEXT: ; %bb.3:
81+ ; GFX908-NEXT: ; implicit-def: $sgpr3
8882; GFX908-NEXT: .LBB0_4: ; %common.ret
8983; GFX908-NEXT: s_endpgm
84+ entry:
85+ br label %bb
86+
9087bb:
91- br label %bb2
88+ %i = phi { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } [ %i10 , %bb2 ], [ zeroinitializer , %entry ]
89+ %i1 = phi i32 [ %i5 , %bb2 ], [ 0 , %entry ]
90+ %c0 = icmp ne i32 %a0 , 0
91+ br i1 %c0 , label %bb2 , label %bb11
9292
9393bb2:
94- %i = phi { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } [ %i12 , %bb4 ], [ zeroinitializer , %bb ]
95- %i3 = phi i32 [ %i7 , %bb4 ], [ 0 , %bb ]
96- br i1 %arg , label %bb4 , label %bb13
97-
98- bb4:
99- %i5 = or i32 %i3 , 1
100- %i6 = icmp slt i32 %i3 , 0
101- %i7 = select i1 %i6 , i32 %i5 , i32 0
102- %i8 = extractvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } %i , 123
103- %i9 = insertelement <4 x float > zeroinitializer , float %i8 , i32 0
104- %i10 = call <4 x float > @llvm.amdgcn.mfma.f32.16x16x16f16 (<4 x half > zeroinitializer , <4 x half > zeroinitializer , <4 x float > %i9 , i32 0 , i32 0 , i32 0 )
105- %i11 = extractelement <4 x float > %i10 , i32 0
106- %i12 = insertvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } zeroinitializer , float %i11 , 123
107- br label %bb2
94+ %i3 = or i32 %i1 , 1
95+ %i4 = icmp slt i32 %i1 , 0
96+ %i5 = select i1 %i4 , i32 %i3 , i32 0
97+ %i6 = extractvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } %i , 123
98+ %i7 = insertelement <4 x float > zeroinitializer , float %i6 , i32 0
99+ %i8 = call <4 x float > @llvm.amdgcn.mfma.f32.16x16x16f16 (<4 x half > zeroinitializer , <4 x half > zeroinitializer , <4 x float > %i7 , i32 0 , i32 0 , i32 0 )
100+ %i9 = extractelement <4 x float > %i8 , i32 0
101+ %i10 = insertvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } zeroinitializer , float %i9 , 123
102+ br label %bb
108103
109- bb13:
110- br i1 %arg1 , label %bb14 , label %common.ret
104+ bb11:
105+ %c1 = icmp ne i32 %a1 , 0
106+ br i1 %c1 , label %bb12 , label %common.ret
111107
112108common .ret:
113109 ret void
114110
115- bb14:
116- %i15 = extractvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } %i , 0
117- %i16 = insertelement <4 x float > zeroinitializer , float %i15 , i32 0
111+ bb12:
112+ %i13 = extractvalue { float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float , float } %i , 0
113+ %i14 = insertelement <4 x float > zeroinitializer , float %i13 , i32 0
114+ %i15 = insertelement <4 x float > %i14 , float 0 .000000e+00 , i32 0
115+ %i16 = insertelement <4 x float > %i15 , float 0 .000000e+00 , i32 0
118116 br label %common.ret
119117}
120118
121119; Function Attrs: convergent nocallback nofree nosync nounwind willreturn memory(none)
122- declare <4 x float > @llvm.amdgcn.mfma.f32.16x16x16f16 (<4 x half >, <4 x half >, <4 x float >, i32 immarg, i32 immarg, i32 immarg) #0
123-
124- attributes #0 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }
120+ declare <4 x float > @llvm.amdgcn.mfma.f32.16x16x16f16 (<4 x half >, <4 x half >, <4 x float >, i32 immarg, i32 immarg, i32 immarg)
0 commit comments