Skip to content

Commit dc606e8

Browse files
arsenmarsenm
committed
Revert the iteration change
Only keep the refactor of the main loop body into a function. Keep the worklist for now, later patch would have just un-inverted it later.
1 parent ac04cf0 commit dc606e8

File tree

6 files changed

+107
-103
lines changed

6 files changed

+107
-103
lines changed

‎llvm/lib/CodeGen/AtomicExpandPass.cpp

Lines changed: 11 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -351,13 +351,17 @@ bool AtomicExpandImpl::run(Function &F, const TargetMachine *TM) {
351351

352352
bool MadeChange = false;
353353

354-
for (BasicBlock &BB : make_early_inc_range(F)) {
355-
for (Instruction &I : make_early_inc_range(reverse(BB))) {
356-
// We do this iteration backwards because the control flow introducing
357-
// transforms split the block at the end.
358-
if (processAtomicInstr(&I))
359-
MadeChange = true;
360-
}
354+
SmallVector<Instruction *, 1> AtomicInsts;
355+
356+
// Changing control-flow while iterating through it is a bad idea, so gather a
357+
// list of all atomic instructions before we start.
358+
for (Instruction &I : instructions(F))
359+
if (I.isAtomic() && !isa<FenceInst>(&I))
360+
AtomicInsts.push_back(&I);
361+
362+
for (auto *I : AtomicInsts) {
363+
if (processAtomicInstr(I))
364+
MadeChange = true;
361365
}
362366

363367
return MadeChange;

‎llvm/test/CodeGen/NVPTX/atomics-sm70.ll

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -61,7 +61,7 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
6161
; CHECKPTX62-NEXT: shl.b32 %r27, %r26, %r2;
6262
; CHECKPTX62-NEXT: not.b32 %r3, %r27;
6363
; CHECKPTX62-NEXT: ld.u32 %r54, [%r1];
64-
; CHECKPTX62-NEXT: $L__BB0_1: // %atomicrmw.start45
64+
; CHECKPTX62-NEXT: $L__BB0_1: // %atomicrmw.start
6565
; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1
6666
; CHECKPTX62-NEXT: shr.u32 %r28, %r54, %r2;
6767
; CHECKPTX62-NEXT: cvt.u16.u32 %rs2, %r28;
@@ -74,9 +74,9 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
7474
; CHECKPTX62-NEXT: setp.ne.s32 %p1, %r6, %r54;
7575
; CHECKPTX62-NEXT: mov.u32 %r54, %r6;
7676
; CHECKPTX62-NEXT: @%p1 bra $L__BB0_1;
77-
; CHECKPTX62-NEXT: // %bb.2: // %atomicrmw.end44
77+
; CHECKPTX62-NEXT: // %bb.2: // %atomicrmw.end
7878
; CHECKPTX62-NEXT: ld.u32 %r55, [%r1];
79-
; CHECKPTX62-NEXT: $L__BB0_3: // %atomicrmw.start27
79+
; CHECKPTX62-NEXT: $L__BB0_3: // %atomicrmw.start9
8080
; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1
8181
; CHECKPTX62-NEXT: shr.u32 %r33, %r55, %r2;
8282
; CHECKPTX62-NEXT: cvt.u16.u32 %rs6, %r33;
@@ -90,14 +90,14 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
9090
; CHECKPTX62-NEXT: setp.ne.s32 %p2, %r9, %r55;
9191
; CHECKPTX62-NEXT: mov.u32 %r55, %r9;
9292
; CHECKPTX62-NEXT: @%p2 bra $L__BB0_3;
93-
; CHECKPTX62-NEXT: // %bb.4: // %atomicrmw.end26
93+
; CHECKPTX62-NEXT: // %bb.4: // %atomicrmw.end8
9494
; CHECKPTX62-NEXT: and.b32 %r10, %r22, -4;
9595
; CHECKPTX62-NEXT: shl.b32 %r38, %r22, 3;
9696
; CHECKPTX62-NEXT: and.b32 %r11, %r38, 24;
9797
; CHECKPTX62-NEXT: shl.b32 %r40, %r26, %r11;
9898
; CHECKPTX62-NEXT: not.b32 %r12, %r40;
9999
; CHECKPTX62-NEXT: ld.global.u32 %r56, [%r10];
100-
; CHECKPTX62-NEXT: $L__BB0_5: // %atomicrmw.start9
100+
; CHECKPTX62-NEXT: $L__BB0_5: // %atomicrmw.start27
101101
; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1
102102
; CHECKPTX62-NEXT: shr.u32 %r41, %r56, %r11;
103103
; CHECKPTX62-NEXT: cvt.u16.u32 %rs11, %r41;
@@ -110,14 +110,14 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
110110
; CHECKPTX62-NEXT: setp.ne.s32 %p3, %r15, %r56;
111111
; CHECKPTX62-NEXT: mov.u32 %r56, %r15;
112112
; CHECKPTX62-NEXT: @%p3 bra $L__BB0_5;
113-
; CHECKPTX62-NEXT: // %bb.6: // %atomicrmw.end8
113+
; CHECKPTX62-NEXT: // %bb.6: // %atomicrmw.end26
114114
; CHECKPTX62-NEXT: and.b32 %r16, %r23, -4;
115115
; CHECKPTX62-NEXT: shl.b32 %r46, %r23, 3;
116116
; CHECKPTX62-NEXT: and.b32 %r17, %r46, 24;
117117
; CHECKPTX62-NEXT: shl.b32 %r48, %r26, %r17;
118118
; CHECKPTX62-NEXT: not.b32 %r18, %r48;
119119
; CHECKPTX62-NEXT: ld.shared.u32 %r57, [%r16];
120-
; CHECKPTX62-NEXT: $L__BB0_7: // %atomicrmw.start
120+
; CHECKPTX62-NEXT: $L__BB0_7: // %atomicrmw.start45
121121
; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1
122122
; CHECKPTX62-NEXT: shr.u32 %r49, %r57, %r17;
123123
; CHECKPTX62-NEXT: cvt.u16.u32 %rs15, %r49;
@@ -130,7 +130,7 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
130130
; CHECKPTX62-NEXT: setp.ne.s32 %p4, %r21, %r57;
131131
; CHECKPTX62-NEXT: mov.u32 %r57, %r21;
132132
; CHECKPTX62-NEXT: @%p4 bra $L__BB0_7;
133-
; CHECKPTX62-NEXT: // %bb.8: // %atomicrmw.end
133+
; CHECKPTX62-NEXT: // %bb.8: // %atomicrmw.end44
134134
; CHECKPTX62-NEXT: ret;
135135
%r1 = atomicrmw fadd ptr %dp0, half %val seq_cst
136136
%r2 = atomicrmw fadd ptr %dp0, half 1.0 seq_cst

‎llvm/test/CodeGen/NVPTX/atomics-sm90.ll

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat
6363
; CHECKPTX71-NEXT: not.b32 %r3, %r27;
6464
; CHECKPTX71-NEXT: ld.u32 %r54, [%r1];
6565
; CHECKPTX71-NEXT: cvt.f32.bf16 %f2, %rs1;
66-
; CHECKPTX71-NEXT: $L__BB0_1: // %atomicrmw.start45
66+
; CHECKPTX71-NEXT: $L__BB0_1: // %atomicrmw.start
6767
; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1
6868
; CHECKPTX71-NEXT: shr.u32 %r28, %r54, %r2;
6969
; CHECKPTX71-NEXT: cvt.u16.u32 %rs2, %r28;
@@ -78,9 +78,9 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat
7878
; CHECKPTX71-NEXT: setp.ne.s32 %p1, %r6, %r54;
7979
; CHECKPTX71-NEXT: mov.u32 %r54, %r6;
8080
; CHECKPTX71-NEXT: @%p1 bra $L__BB0_1;
81-
; CHECKPTX71-NEXT: // %bb.2: // %atomicrmw.end44
81+
; CHECKPTX71-NEXT: // %bb.2: // %atomicrmw.end
8282
; CHECKPTX71-NEXT: ld.u32 %r55, [%r1];
83-
; CHECKPTX71-NEXT: $L__BB0_3: // %atomicrmw.start27
83+
; CHECKPTX71-NEXT: $L__BB0_3: // %atomicrmw.start9
8484
; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1
8585
; CHECKPTX71-NEXT: shr.u32 %r33, %r55, %r2;
8686
; CHECKPTX71-NEXT: cvt.u16.u32 %rs6, %r33;
@@ -95,14 +95,14 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat
9595
; CHECKPTX71-NEXT: setp.ne.s32 %p2, %r9, %r55;
9696
; CHECKPTX71-NEXT: mov.u32 %r55, %r9;
9797
; CHECKPTX71-NEXT: @%p2 bra $L__BB0_3;
98-
; CHECKPTX71-NEXT: // %bb.4: // %atomicrmw.end26
98+
; CHECKPTX71-NEXT: // %bb.4: // %atomicrmw.end8
9999
; CHECKPTX71-NEXT: and.b32 %r10, %r22, -4;
100100
; CHECKPTX71-NEXT: shl.b32 %r38, %r22, 3;
101101
; CHECKPTX71-NEXT: and.b32 %r11, %r38, 24;
102102
; CHECKPTX71-NEXT: shl.b32 %r40, %r26, %r11;
103103
; CHECKPTX71-NEXT: not.b32 %r12, %r40;
104104
; CHECKPTX71-NEXT: ld.global.u32 %r56, [%r10];
105-
; CHECKPTX71-NEXT: $L__BB0_5: // %atomicrmw.start9
105+
; CHECKPTX71-NEXT: $L__BB0_5: // %atomicrmw.start27
106106
; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1
107107
; CHECKPTX71-NEXT: shr.u32 %r41, %r56, %r11;
108108
; CHECKPTX71-NEXT: cvt.u16.u32 %rs10, %r41;
@@ -117,14 +117,14 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat
117117
; CHECKPTX71-NEXT: setp.ne.s32 %p3, %r15, %r56;
118118
; CHECKPTX71-NEXT: mov.u32 %r56, %r15;
119119
; CHECKPTX71-NEXT: @%p3 bra $L__BB0_5;
120-
; CHECKPTX71-NEXT: // %bb.6: // %atomicrmw.end8
120+
; CHECKPTX71-NEXT: // %bb.6: // %atomicrmw.end26
121121
; CHECKPTX71-NEXT: and.b32 %r16, %r23, -4;
122122
; CHECKPTX71-NEXT: shl.b32 %r46, %r23, 3;
123123
; CHECKPTX71-NEXT: and.b32 %r17, %r46, 24;
124124
; CHECKPTX71-NEXT: shl.b32 %r48, %r26, %r17;
125125
; CHECKPTX71-NEXT: not.b32 %r18, %r48;
126126
; CHECKPTX71-NEXT: ld.shared.u32 %r57, [%r16];
127-
; CHECKPTX71-NEXT: $L__BB0_7: // %atomicrmw.start
127+
; CHECKPTX71-NEXT: $L__BB0_7: // %atomicrmw.start45
128128
; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1
129129
; CHECKPTX71-NEXT: shr.u32 %r49, %r57, %r17;
130130
; CHECKPTX71-NEXT: cvt.u16.u32 %rs14, %r49;
@@ -139,7 +139,7 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat
139139
; CHECKPTX71-NEXT: setp.ne.s32 %p4, %r21, %r57;
140140
; CHECKPTX71-NEXT: mov.u32 %r57, %r21;
141141
; CHECKPTX71-NEXT: @%p4 bra $L__BB0_7;
142-
; CHECKPTX71-NEXT: // %bb.8: // %atomicrmw.end
142+
; CHECKPTX71-NEXT: // %bb.8: // %atomicrmw.end44
143143
; CHECKPTX71-NEXT: ret;
144144
%r1 = atomicrmw fadd ptr %dp0, bfloat %val seq_cst
145145
%r2 = atomicrmw fadd ptr %dp0, bfloat 1.0 seq_cst

0 commit comments

Comments
 (0)