Skip to content

Commit d4ac1c2

Browse files
ronliebronlieb
committed
lit test regen/fix
1 parent fd0fd16 commit d4ac1c2

File tree

5 files changed

+47
-161
lines changed

5 files changed

+47
-161
lines changed

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl

Lines changed: 0 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -125,50 +125,6 @@ void test_s_barrier_signal_isfirst(int* a, int* b, int *c)
125125
__builtin_amdgcn_s_barrier_wait(1);
126126
}
127127

128-
<<<<<<< HEAD
129-
// CHECK-LABEL: @test_s_barrier_init(
130-
// CHECK-NEXT: entry:
131-
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
132-
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
133-
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8
134-
// CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4
135-
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8
136-
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
137-
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4
138-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) [[TMP1]], i32 [[TMP2]])
139-
// CHECK-NEXT: ret void
140-
//
141-
void test_s_barrier_init(void *bar, int a)
142-
{
143-
__builtin_amdgcn_s_barrier_init(bar, a);
144-
}
145-
146-
// CHECK-LABEL: @test_s_barrier_join(
147-
// CHECK-NEXT: entry:
148-
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
149-
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8
150-
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8
151-
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
152-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) [[TMP1]])
153-
// CHECK-NEXT: ret void
154-
//
155-
void test_s_barrier_join(void *bar)
156-
{
157-
__builtin_amdgcn_s_barrier_join(bar);
158-
}
159-
160-
// CHECK-LABEL: @test_s_barrier_leave(
161-
// CHECK-NEXT: entry:
162-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.leave(i16 1)
163-
// CHECK-NEXT: ret void
164-
//
165-
void test_s_barrier_leave()
166-
{
167-
__builtin_amdgcn_s_barrier_leave(1);
168-
}
169-
170-
=======
171-
>>>>>>> a614f2b489caa19001b2f44784514a6226f79cb7
172128
// CHECK-LABEL: @test_s_get_barrier_state(
173129
// CHECK-NEXT: entry:
174130
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)

‎llvm/test/CodeGen/AMDGPU/eliminate-frame-index-s-add-u32.mir

Lines changed: 36 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -21,21 +21,33 @@ machineFunctionInfo:
2121
body: |
2222
bb.0:
2323
; MUBUFW64-LABEL: name: s_add_u32__inline_imm__fi_offset0
24-
; MUBUFW64: renamable $sgpr4 = S_LSHR_B32 $sgpr32, 6, implicit-def dead $scc
24+
; MUBUFW64: frame-setup CFI_INSTRUCTION llvm_def_aspace_cfa $sgpr32, 0, 6
25+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION llvm_register_pair $pc_reg, $sgpr30, 32, $sgpr31, 32
26+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr7
27+
; MUBUFW64-NEXT: renamable $sgpr4 = S_LSHR_B32 $sgpr32, 6, implicit-def dead $scc
2528
; MUBUFW64-NEXT: renamable $sgpr7 = S_ADD_U32 12, $sgpr4, implicit-def dead $scc
2629
; MUBUFW64-NEXT: SI_RETURN implicit $sgpr7
2730
;
2831
; MUBUFW32-LABEL: name: s_add_u32__inline_imm__fi_offset0
29-
; MUBUFW32: renamable $sgpr4 = S_LSHR_B32 $sgpr32, 5, implicit-def dead $scc
32+
; MUBUFW32: frame-setup CFI_INSTRUCTION llvm_def_aspace_cfa $sgpr32, 0, 6
33+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION llvm_register_pair $pc_reg, $sgpr30, 32, $sgpr31, 32
34+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr7
35+
; MUBUFW32-NEXT: renamable $sgpr4 = S_LSHR_B32 $sgpr32, 5, implicit-def dead $scc
3036
; MUBUFW32-NEXT: renamable $sgpr7 = S_ADD_U32 12, $sgpr4, implicit-def dead $scc
3137
; MUBUFW32-NEXT: SI_RETURN implicit $sgpr7
3238
;
3339
; FLATSCRW64-LABEL: name: s_add_u32__inline_imm__fi_offset0
34-
; FLATSCRW64: renamable $sgpr7 = S_ADD_U32 12, $sgpr32, implicit-def dead $scc
40+
; FLATSCRW64: frame-setup CFI_INSTRUCTION escape 0x0f, 0x09, 0x90, 0x40, 0x94, 0x04, 0x36, 0x24, 0x36, 0xe9, 0x02
41+
; FLATSCRW64-NEXT: frame-setup CFI_INSTRUCTION llvm_register_pair $pc_reg, $sgpr30, 32, $sgpr31, 32
42+
; FLATSCRW64-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr7
43+
; FLATSCRW64-NEXT: renamable $sgpr7 = S_ADD_U32 12, $sgpr32, implicit-def dead $scc
3544
; FLATSCRW64-NEXT: SI_RETURN implicit $sgpr7
3645
;
3746
; FLATSCRW32-LABEL: name: s_add_u32__inline_imm__fi_offset0
38-
; FLATSCRW32: renamable $sgpr7 = S_ADD_U32 12, $sgpr32, implicit-def dead $scc
47+
; FLATSCRW32: frame-setup CFI_INSTRUCTION escape 0x0f, 0x09, 0x90, 0x40, 0x94, 0x04, 0x35, 0x24, 0x36, 0xe9, 0x02
48+
; FLATSCRW32-NEXT: frame-setup CFI_INSTRUCTION llvm_register_pair $pc_reg, $sgpr30, 32, $sgpr31, 32
49+
; FLATSCRW32-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr7
50+
; FLATSCRW32-NEXT: renamable $sgpr7 = S_ADD_U32 12, $sgpr32, implicit-def dead $scc
3951
; FLATSCRW32-NEXT: SI_RETURN implicit $sgpr7
4052
renamable $sgpr7 = S_ADD_U32 12, %stack.0, implicit-def dead $scc
4153
SI_RETURN implicit $sgpr7
@@ -58,6 +70,8 @@ body: |
5870
; MUBUFW64-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal
5971
; MUBUFW64: liveins: $sgpr0_sgpr1_sgpr2_sgpr3
6072
; MUBUFW64-NEXT: {{ $}}
73+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
74+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
6175
; MUBUFW64-NEXT: $sgpr0 = S_ADD_U32 $sgpr0, $noreg, implicit-def $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
6276
; MUBUFW64-NEXT: $sgpr1 = S_ADDC_U32 $sgpr1, 0, implicit-def dead $scc, implicit $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
6377
; MUBUFW64-NEXT: renamable $sgpr7 = S_MOV_B32 164
@@ -66,17 +80,23 @@ body: |
6680
; MUBUFW32-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal
6781
; MUBUFW32: liveins: $sgpr0_sgpr1_sgpr2_sgpr3
6882
; MUBUFW32-NEXT: {{ $}}
83+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
84+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
6985
; MUBUFW32-NEXT: $sgpr0 = S_ADD_U32 $sgpr0, $noreg, implicit-def $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
7086
; MUBUFW32-NEXT: $sgpr1 = S_ADDC_U32 $sgpr1, 0, implicit-def dead $scc, implicit $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
7187
; MUBUFW32-NEXT: renamable $sgpr7 = S_MOV_B32 164
7288
; MUBUFW32-NEXT: SI_RETURN implicit $sgpr7
7389
;
7490
; FLATSCRW64-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal
75-
; FLATSCRW64: renamable $sgpr7 = S_MOV_B32 164
91+
; FLATSCRW64: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
92+
; FLATSCRW64-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
93+
; FLATSCRW64-NEXT: renamable $sgpr7 = S_MOV_B32 164
7694
; FLATSCRW64-NEXT: SI_RETURN implicit $sgpr7
7795
;
7896
; FLATSCRW32-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal
79-
; FLATSCRW32: renamable $sgpr7 = S_MOV_B32 164
97+
; FLATSCRW32: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
98+
; FLATSCRW32-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
99+
; FLATSCRW32-NEXT: renamable $sgpr7 = S_MOV_B32 164
80100
; FLATSCRW32-NEXT: SI_RETURN implicit $sgpr7
81101
renamable $sgpr7 = S_ADD_U32 68, %stack.1, implicit-def dead $scc
82102
SI_RETURN implicit $sgpr7
@@ -98,6 +118,8 @@ body: |
98118
; MUBUFW64-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal_live_scc
99119
; MUBUFW64: liveins: $sgpr0_sgpr1_sgpr2_sgpr3
100120
; MUBUFW64-NEXT: {{ $}}
121+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
122+
; MUBUFW64-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
101123
; MUBUFW64-NEXT: $sgpr0 = S_ADD_U32 $sgpr0, $noreg, implicit-def $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
102124
; MUBUFW64-NEXT: $sgpr1 = S_ADDC_U32 $sgpr1, 0, implicit-def dead $scc, implicit $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
103125
; MUBUFW64-NEXT: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
@@ -106,17 +128,23 @@ body: |
106128
; MUBUFW32-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal_live_scc
107129
; MUBUFW32: liveins: $sgpr0_sgpr1_sgpr2_sgpr3
108130
; MUBUFW32-NEXT: {{ $}}
131+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
132+
; MUBUFW32-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
109133
; MUBUFW32-NEXT: $sgpr0 = S_ADD_U32 $sgpr0, $noreg, implicit-def $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
110134
; MUBUFW32-NEXT: $sgpr1 = S_ADDC_U32 $sgpr1, 0, implicit-def dead $scc, implicit $scc, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3
111135
; MUBUFW32-NEXT: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
112136
; MUBUFW32-NEXT: SI_RETURN implicit $sgpr7, implicit $scc
113137
;
114138
; FLATSCRW64-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal_live_scc
115-
; FLATSCRW64: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
139+
; FLATSCRW64: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
140+
; FLATSCRW64-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
141+
; FLATSCRW64-NEXT: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
116142
; FLATSCRW64-NEXT: SI_RETURN implicit $sgpr7, implicit $scc
117143
;
118144
; FLATSCRW32-LABEL: name: s_add_u32__kernel__literal__fi_offset96__offset_literal_live_scc
119-
; FLATSCRW32: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
145+
; FLATSCRW32: frame-setup CFI_INSTRUCTION escape 0x0f, 0x04, 0x30, 0x36, 0xe9, 0x02
146+
; FLATSCRW32-NEXT: frame-setup CFI_INSTRUCTION undefined $pc_reg
147+
; FLATSCRW32-NEXT: renamable $sgpr7 = S_ADD_U32 164, 0, implicit-def $scc
120148
; FLATSCRW32-NEXT: SI_RETURN implicit $sgpr7, implicit $scc
121149
renamable $sgpr7 = S_ADD_U32 68, %stack.1, implicit-def $scc
122150
SI_RETURN implicit $sgpr7, implicit $scc

‎llvm/test/CodeGen/AMDGPU/frame-index-elimination.ll

Lines changed: 0 additions & 70 deletions
Original file line numberDiff line numberDiff line change
@@ -346,74 +346,4 @@ entry:
346346
ret void
347347
}
348348

349-
; Check for "SOP2/SOPC instruction requires too many immediate
350-
; constants" verifier error. Frame index would fold into low half of
351-
; the lowered flat pointer add, and use s_add_u32 instead of
352-
; s_add_i32.
353-
354-
; GCN-LABEL: {{^}}fi_sop2_s_add_u32_literal_error:
355-
; GCN: s_add_u32 [[ADD_LO:s[0-9]+]], 0, 0x2010
356-
; GCN: s_addc_u32 [[ADD_HI:s[0-9]+]], s{{[0-9]+}}, 0
357-
define amdgpu_kernel void @fi_sop2_s_add_u32_literal_error() #0 {
358-
entry:
359-
%.omp.reduction.element.i.i.i.i = alloca [1024 x i32], align 4, addrspace(5)
360-
%Total3.i.i = alloca [1024 x i32], align 16, addrspace(5)
361-
%Total3.ascast.i.i = addrspacecast ptr addrspace(5) %Total3.i.i to ptr
362-
%gep = getelementptr i8, ptr %Total3.ascast.i.i, i64 4096
363-
%p2i = ptrtoint ptr %gep to i64
364-
br label %.shuffle.then.i.i.i.i
365-
366-
.shuffle.then.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i, %entry
367-
store i64 0, ptr addrspace(5) null, align 4
368-
%icmp = icmp ugt i64 %p2i, 1
369-
br i1 %icmp, label %.shuffle.then.i.i.i.i, label %vector.body.i.i.i.i
370-
371-
vector.body.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i
372-
%wide.load9.i.i.i.i = load <2 x i32>, ptr addrspace(5) %.omp.reduction.element.i.i.i.i, align 4
373-
store <2 x i32> %wide.load9.i.i.i.i, ptr addrspace(5) null, align 4
374-
ret void
375-
}
376-
377-
; GCN-LABEL: {{^}}fi_sop2_and_literal_error:
378-
; GCN: s_and_b32 s{{[0-9]+}}, s{{[0-9]+}}, 0x1fe00
379-
define amdgpu_kernel void @fi_sop2_and_literal_error() #0 {
380-
entry:
381-
%.omp.reduction.element.i.i.i.i = alloca [1024 x i32], align 4, addrspace(5)
382-
%Total3.i.i = alloca [1024 x i32], align 16, addrspace(5)
383-
%p2i = ptrtoint ptr addrspace(5) %Total3.i.i to i32
384-
br label %.shuffle.then.i.i.i.i
385-
386-
.shuffle.then.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i, %entry
387-
store i64 0, ptr addrspace(5) null, align 4
388-
%or = and i32 %p2i, -512
389-
%icmp = icmp ugt i32 %or, 9999999
390-
br i1 %icmp, label %.shuffle.then.i.i.i.i, label %vector.body.i.i.i.i
391-
392-
vector.body.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i
393-
%wide.load9.i.i.i.i = load <2 x i32>, ptr addrspace(5) %.omp.reduction.element.i.i.i.i, align 4
394-
store <2 x i32> %wide.load9.i.i.i.i, ptr addrspace(5) null, align 4
395-
ret void
396-
}
397-
398-
; GCN-LABEL: {{^}}fi_sop2_or_literal_error:
399-
; GCN: s_or_b32 s{{[0-9]+}}, s{{[0-9]+}}, 0x3039
400-
define amdgpu_kernel void @fi_sop2_or_literal_error() #0 {
401-
entry:
402-
%.omp.reduction.element.i.i.i.i = alloca [1024 x i32], align 4, addrspace(5)
403-
%Total3.i.i = alloca [1024 x i32], align 16, addrspace(5)
404-
%p2i = ptrtoint ptr addrspace(5) %Total3.i.i to i32
405-
br label %.shuffle.then.i.i.i.i
406-
407-
.shuffle.then.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i, %entry
408-
store i64 0, ptr addrspace(5) null, align 4
409-
%or = or i32 %p2i, 12345
410-
%icmp = icmp ugt i32 %or, 9999999
411-
br i1 %icmp, label %.shuffle.then.i.i.i.i, label %vector.body.i.i.i.i
412-
413-
vector.body.i.i.i.i: ; preds = %.shuffle.then.i.i.i.i
414-
%wide.load9.i.i.i.i = load <2 x i32>, ptr addrspace(5) %.omp.reduction.element.i.i.i.i, align 4
415-
store <2 x i32> %wide.load9.i.i.i.i, ptr addrspace(5) null, align 4
416-
ret void
417-
}
418-
419349
attributes #0 = { nounwind }

‎llvm/test/CodeGen/AMDGPU/frame-index.mir

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -98,19 +98,14 @@ body: |
9898
; GCN-LABEL: name: func_add_constant_to_fi_uniform_SCC_clobber_i32
9999
; GCN: liveins: $sgpr30_sgpr31
100100
; GCN-NEXT: {{ $}}
101-
<<<<<<< HEAD
102101
; GCN-NEXT: frame-setup CFI_INSTRUCTION llvm_def_aspace_cfa $sgpr32, 0, 6
103102
; GCN-NEXT: frame-setup CFI_INSTRUCTION llvm_register_pair $pc_reg, $sgpr30, 32, $sgpr31, 32
104103
; GCN-NEXT: frame-setup CFI_INSTRUCTION undefined $vgpr0
105104
; GCN-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr4
106105
; GCN-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr5
107106
; GCN-NEXT: frame-setup CFI_INSTRUCTION undefined $sgpr6
108-
; GCN-NEXT: $sgpr0 = S_LSHR_B32 $sgpr32, 6, implicit-def dead $scc
109-
; GCN-NEXT: renamable $sgpr4 = nuw S_ADD_U32 killed $sgpr0, 4, implicit-def $scc
110-
=======
111107
; GCN-NEXT: renamable $sgpr0 = S_LSHR_B32 $sgpr32, 6, implicit-def dead $scc
112108
; GCN-NEXT: renamable $sgpr4 = nuw S_ADD_U32 $sgpr0, 4, implicit-def $scc
113-
>>>>>>> a614f2b489caa19001b2f44784514a6226f79cb7
114109
; GCN-NEXT: renamable $sgpr5 = S_ADDC_U32 $sgpr4, 1234567, implicit-def $scc, implicit $scc
115110
; GCN-NEXT: $sgpr0 = S_LSHR_B32 $sgpr32, 6, implicit-def $scc
116111
; GCN-NEXT: $sgpr0 = S_ADD_I32 killed $sgpr0, 8, implicit-def $scc

‎llvm/test/CodeGen/AMDGPU/local-stack-alloc-block-sp-reference.ll

Lines changed: 11 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -228,58 +228,35 @@ define amdgpu_kernel void @local_stack_offset_uses_sp_flat(ptr addrspace(1) %out
228228
; MUBUF-NEXT: s_waitcnt vmcnt(0)
229229
; MUBUF-NEXT: s_cbranch_scc1 .LBB2_1
230230
; MUBUF-NEXT: ; %bb.2: ; %split
231+
; MUBUF-NEXT: s_movk_i32 s5, 0x12d4
231232
; MUBUF-NEXT: v_mov_b32_e32 v1, 0x4000
232-
<<<<<<< HEAD
233-
; MUBUF-NEXT: s_movk_i32 s4, 0x12d4
234-
; MUBUF-NEXT: v_mov_b32_e32 v2, 0x4000
235-
; MUBUF-NEXT: v_or_b32_e32 v0, 0x12c0, v1
236-
; MUBUF-NEXT: v_or_b32_e32 v1, s4, v2
237-
; MUBUF-NEXT: s_movk_i32 s4, 0x12d0
238-
; MUBUF-NEXT: v_mov_b32_e32 v2, 0x4000
239-
; MUBUF-NEXT: buffer_load_dword v5, v1, s[0:3], 0 offen glc
240-
; MUBUF-NEXT: s_waitcnt vmcnt(0)
241-
; MUBUF-NEXT: v_or_b32_e32 v1, s4, v2
242-
; MUBUF-NEXT: s_movk_i32 s4, 0x12c4
243-
; MUBUF-NEXT: v_mov_b32_e32 v2, 0x4000
244-
; MUBUF-NEXT: buffer_load_dword v4, v1, s[0:3], 0 offen glc
245-
; MUBUF-NEXT: s_waitcnt vmcnt(0)
246-
; MUBUF-NEXT: v_or_b32_e32 v1, s4, v2
247-
; MUBUF-NEXT: buffer_load_dword v6, v1, s[0:3], 0 offen glc
248-
; MUBUF-NEXT: s_waitcnt vmcnt(0)
249-
; MUBUF-NEXT: buffer_load_dword v7, v0, s[0:3], 0 offen glc
250-
; MUBUF-NEXT: s_waitcnt vmcnt(0)
251-
; MUBUF-NEXT: s_movk_i32 s4, 0x12cc
252-
; MUBUF-NEXT: v_mov_b32_e32 v1, 0x4000
253-
; MUBUF-NEXT: v_or_b32_e32 v0, s4, v1
254-
; MUBUF-NEXT: s_movk_i32 s4, 0x12c8
255-
; MUBUF-NEXT: v_mov_b32_e32 v2, 0x4000
256-
; MUBUF-NEXT: v_or_b32_e32 v1, s4, v2
257-
; MUBUF-NEXT: v_mov_b32_e32 v2, 0x4000
258-
=======
259-
; MUBUF-NEXT: v_or_b32_e32 v0, 0x12d4, v1
233+
; MUBUF-NEXT: v_or_b32_e32 v0, s5, v1
234+
; MUBUF-NEXT: s_movk_i32 s5, 0x12d0
260235
; MUBUF-NEXT: v_mov_b32_e32 v1, 0x4000
261236
; MUBUF-NEXT: s_movk_i32 s4, 0x4000
262237
; MUBUF-NEXT: buffer_load_dword v5, v0, s[0:3], 0 offen glc
263238
; MUBUF-NEXT: s_waitcnt vmcnt(0)
264-
; MUBUF-NEXT: v_or_b32_e32 v0, 0x12d0, v1
239+
; MUBUF-NEXT: v_or_b32_e32 v0, s5, v1
240+
; MUBUF-NEXT: s_movk_i32 s5, 0x12c4
265241
; MUBUF-NEXT: v_mov_b32_e32 v1, 0x4000
266242
; MUBUF-NEXT: s_or_b32 s4, s4, 0x12c0
267243
; MUBUF-NEXT: buffer_load_dword v4, v0, s[0:3], 0 offen glc
268244
; MUBUF-NEXT: s_waitcnt vmcnt(0)
269-
; MUBUF-NEXT: v_or_b32_e32 v0, 0x12c4, v1
270-
; MUBUF-NEXT: v_mov_b32_e32 v3, 0x4000
245+
; MUBUF-NEXT: v_or_b32_e32 v0, s5, v1
271246
; MUBUF-NEXT: buffer_load_dword v1, v0, s[0:3], 0 offen glc
272247
; MUBUF-NEXT: s_waitcnt vmcnt(0)
273248
; MUBUF-NEXT: v_mov_b32_e32 v0, s4
274-
; MUBUF-NEXT: v_or_b32_e32 v2, 0x12cc, v3
249+
; MUBUF-NEXT: s_movk_i32 s4, 0x12cc
250+
; MUBUF-NEXT: v_mov_b32_e32 v3, 0x4000
251+
; MUBUF-NEXT: v_or_b32_e32 v2, s4, v3
252+
; MUBUF-NEXT: s_movk_i32 s4, 0x12c8
275253
; MUBUF-NEXT: v_mov_b32_e32 v6, 0x4000
276-
>>>>>>> a614f2b489caa19001b2f44784514a6226f79cb7
277254
; MUBUF-NEXT: buffer_load_dword v0, v0, s[0:3], 0 offen glc
278255
; MUBUF-NEXT: s_waitcnt vmcnt(0)
279256
; MUBUF-NEXT: v_mov_b32_e32 v7, 0x4000
280257
; MUBUF-NEXT: buffer_load_dword v3, v2, s[0:3], 0 offen glc
281258
; MUBUF-NEXT: s_waitcnt vmcnt(0)
282-
; MUBUF-NEXT: v_or_b32_e32 v2, 0x12c8, v6
259+
; MUBUF-NEXT: v_or_b32_e32 v2, s4, v6
283260
; MUBUF-NEXT: v_mov_b32_e32 v8, 0x4000
284261
; MUBUF-NEXT: v_mov_b32_e32 v9, 0x4000
285262
; MUBUF-NEXT: buffer_load_dword v2, v2, s[0:3], 0 offen glc

0 commit comments

Comments
 (0)