Skip to content

Commit 4357852

Browse files
gandhi56gandhi56
committed
[Remarks] Emit optimization remarks for atomics generating CAS loop
Implements ORE in AtomicExpand pass to report atomics generating a compare and swap loop. Differential Revision: https://reviews.llvm.org/D106891
1 parent 530aa7e commit 4357852

File tree

11 files changed

+249
-7
lines changed

11 files changed

+249
-7
lines changed

‎clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
1+
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
2+
// RUN: -target-cpu gfx90a -Rpass=atomic-expand -S -o - 2>&1 | \
3+
// RUN: FileCheck %s --check-prefix=GFX90A-CAS
4+
5+
// REQUIRES: amdgpu-registered-target
6+
7+
#include "Inputs/cuda.h"
8+
#include <stdatomic.h>
9+
10+
// GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at system memory scope
11+
// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
12+
// GFX90A-CAS: flat_atomic_cmpswap v0, v[2:3], v[4:5] glc
13+
// GFX90A-CAS: s_cbranch_execnz
14+
__device__ float atomic_add_cas(float *p) {
15+
return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
16+
}

‎clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \
2+
// RUN: -Rpass=atomic-expand -S -o - 2>&1 | \
3+
// RUN: FileCheck %s --check-prefix=REMARK
4+
5+
// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \
6+
// RUN: -Rpass=atomic-expand -S -emit-llvm -o - 2>&1 | \
7+
// RUN: FileCheck %s --check-prefix=GFX90A-CAS
8+
9+
// REQUIRES: amdgpu-registered-target
10+
11+
typedef enum memory_order {
12+
memory_order_relaxed = __ATOMIC_RELAXED,
13+
memory_order_acquire = __ATOMIC_ACQUIRE,
14+
memory_order_release = __ATOMIC_RELEASE,
15+
memory_order_acq_rel = __ATOMIC_ACQ_REL,
16+
memory_order_seq_cst = __ATOMIC_SEQ_CST
17+
} memory_order;
18+
19+
typedef enum memory_scope {
20+
memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
21+
memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
22+
memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
23+
memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
24+
#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
25+
memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
26+
#endif
27+
} memory_scope;
28+
29+
// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at workgroup-one-as memory scope [-Rpass=atomic-expand]
30+
// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at agent-one-as memory scope [-Rpass=atomic-expand]
31+
// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at one-as memory scope [-Rpass=atomic-expand]
32+
// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at wavefront-one-as memory scope [-Rpass=atomic-expand]
33+
// GFX90A-CAS-LABEL: @atomic_cas
34+
// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
35+
// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("agent-one-as") monotonic
36+
// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("one-as") monotonic
37+
// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("wavefront-one-as") monotonic
38+
float atomic_cas(__global atomic_float *d, float a) {
39+
float ret1 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
40+
float ret2 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_device);
41+
float ret3 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_all_svm_devices);
42+
float ret4 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_sub_group);
43+
}
44+
45+
46+

‎llvm/lib/CodeGen/AtomicExpandPass.cpp

Lines changed: 22 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,7 @@
1717
#include "llvm/ADT/ArrayRef.h"
1818
#include "llvm/ADT/STLExtras.h"
1919
#include "llvm/ADT/SmallVector.h"
20+
#include "llvm/Analysis/OptimizationRemarkEmitter.h"
2021
#include "llvm/CodeGen/AtomicExpandUtils.h"
2122
#include "llvm/CodeGen/RuntimeLibcalls.h"
2223
#include "llvm/CodeGen/TargetLowering.h"
@@ -58,6 +59,7 @@ namespace {
5859

5960
class AtomicExpand: public FunctionPass {
6061
const TargetLowering *TLI = nullptr;
62+
OptimizationRemarkEmitter *ORE;
6163

6264
public:
6365
static char ID; // Pass identification, replacement for typeid
@@ -69,6 +71,7 @@ namespace {
6971
bool runOnFunction(Function &F) override;
7072

7173
private:
74+
void getAnalysisUsage(AnalysisUsage &AU) const override;
7275
bool bracketInstWithFences(Instruction *I, AtomicOrdering Order);
7376
IntegerType *getCorrespondingIntegerType(Type *T, const DataLayout &DL);
7477
LoadInst *convertAtomicLoadToIntegerType(LoadInst *LI);
@@ -165,11 +168,16 @@ static bool atomicSizeSupported(const TargetLowering *TLI, Inst *I) {
165168
Size <= TLI->getMaxAtomicSizeInBitsSupported() / 8;
166169
}
167170

171+
void AtomicExpand::getAnalysisUsage(AnalysisUsage &AU) const {
172+
AU.addRequired<OptimizationRemarkEmitterWrapperPass>();
173+
}
174+
168175
bool AtomicExpand::runOnFunction(Function &F) {
169176
auto *TPC = getAnalysisIfAvailable<TargetPassConfig>();
170177
if (!TPC)
171178
return false;
172179

180+
ORE = &getAnalysis<OptimizationRemarkEmitterWrapperPass>().getORE();
173181
auto &TM = TPC->getTM<TargetMachine>();
174182
if (!TM.getSubtargetImpl(F)->enableAtomicExpand())
175183
return false;
@@ -570,7 +578,9 @@ static Value *performAtomicOp(AtomicRMWInst::BinOp Op, IRBuilder<> &Builder,
570578
}
571579

572580
bool AtomicExpand::tryExpandAtomicRMW(AtomicRMWInst *AI) {
573-
switch (TLI->shouldExpandAtomicRMWInIR(AI)) {
581+
LLVMContext &Ctx = AI->getModule()->getContext();
582+
TargetLowering::AtomicExpansionKind Kind = TLI->shouldExpandAtomicRMWInIR(AI);
583+
switch (Kind) {
574584
case TargetLoweringBase::AtomicExpansionKind::None:
575585
return false;
576586
case TargetLoweringBase::AtomicExpansionKind::LLSC: {
@@ -600,6 +610,17 @@ bool AtomicExpand::tryExpandAtomicRMW(AtomicRMWInst *AI) {
600610
expandPartwordAtomicRMW(AI,
601611
TargetLoweringBase::AtomicExpansionKind::CmpXChg);
602612
} else {
613+
SmallVector<StringRef> SSNs;
614+
Ctx.getSyncScopeNames(SSNs);
615+
auto MemScope = SSNs[AI->getSyncScopeID()].empty()
616+
? "system"
617+
: SSNs[AI->getSyncScopeID()];
618+
ORE->emit([&]() {
619+
return OptimizationRemark(DEBUG_TYPE, "Passed", AI->getFunction())
620+
<< "A compare and swap loop was generated for an atomic "
621+
<< AI->getOperationName(AI->getOperation()) << " operation at "
622+
<< MemScope << " memory scope";
623+
});
603624
expandAtomicRMWToCmpXchg(AI, createCmpXchgInstFun);
604625
}
605626
return true;

‎llvm/test/CodeGen/AArch64/O0-pipeline.ll

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,13 +8,18 @@
88
; CHECK-NEXT: Target Pass Configuration
99
; CHECK-NEXT: Machine Module Information
1010
; CHECK-NEXT: Target Transform Information
11+
; CHECK-NEXT: Profile summary info
1112
; CHECK-NEXT: Create Garbage Collector Module Metadata
1213
; CHECK-NEXT: Assumption Cache Tracker
13-
; CHECK-NEXT: Profile summary info
1414
; CHECK-NEXT: Machine Branch Probability Analysis
1515
; CHECK-NEXT: ModulePass Manager
1616
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
1717
; CHECK-NEXT: FunctionPass Manager
18+
; CHECK-NEXT: Dominator Tree Construction
19+
; CHECK-NEXT: Natural Loop Information
20+
; CHECK-NEXT: Lazy Branch Probability Analysis
21+
; CHECK-NEXT: Lazy Block Frequency Analysis
22+
; CHECK-NEXT: Optimization Remark Emitter
1823
; CHECK-NEXT: Expand Atomic instructions
1924
; CHECK-NEXT: Module Verifier
2025
; CHECK-NEXT: Lower Garbage Collection Instructions

‎llvm/test/CodeGen/AArch64/O3-pipeline.ll

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,15 +8,20 @@
88
; CHECK-NEXT: Target Pass Configuration
99
; CHECK-NEXT: Machine Module Information
1010
; CHECK-NEXT: Target Transform Information
11-
; CHECK-NEXT: Assumption Cache Tracker
1211
; CHECK-NEXT: Profile summary info
12+
; CHECK-NEXT: Assumption Cache Tracker
1313
; CHECK-NEXT: Type-Based Alias Analysis
1414
; CHECK-NEXT: Scoped NoAlias Alias Analysis
1515
; CHECK-NEXT: Create Garbage Collector Module Metadata
1616
; CHECK-NEXT: Machine Branch Probability Analysis
1717
; CHECK-NEXT: ModulePass Manager
1818
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
1919
; CHECK-NEXT: FunctionPass Manager
20+
; CHECK-NEXT: Dominator Tree Construction
21+
; CHECK-NEXT: Natural Loop Information
22+
; CHECK-NEXT: Lazy Branch Probability Analysis
23+
; CHECK-NEXT: Lazy Block Frequency Analysis
24+
; CHECK-NEXT: Optimization Remark Emitter
2025
; CHECK-NEXT: Expand Atomic instructions
2126
; CHECK-NEXT: SVE intrinsics optimizations
2227
; CHECK-NEXT: FunctionPass Manager

‎llvm/test/CodeGen/AMDGPU/atomics-remarks-gfx90a.ll

Lines changed: 103 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,103 @@
1+
; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs --pass-remarks=atomic-expand \
2+
; RUN: %s -o - 2>&1 | FileCheck %s --check-prefix=GFX90A-CAS
3+
4+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at system memory scope
5+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at agent memory scope
6+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at workgroup memory scope
7+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at wavefront memory scope
8+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at singlethread memory scope
9+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at one-as memory scope
10+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at agent-one-as memory scope
11+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at workgroup-one-as memory scope
12+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at wavefront-one-as memory scope
13+
; GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at singlethread-one-as memory scope
14+
15+
; GFX90A-CAS-LABEL: atomic_add_cas:
16+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
17+
; GFX90A-CAS: s_cbranch_execnz
18+
define dso_local void @atomic_add_cas(float* %p, float %q) {
19+
entry:
20+
%ret = atomicrmw fadd float* %p, float %q monotonic, align 4
21+
ret void
22+
}
23+
24+
; GFX90A-CAS-LABEL: atomic_add_cas_agent:
25+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
26+
; GFX90A-CAS: s_cbranch_execnz
27+
define dso_local void @atomic_add_cas_agent(float* %p, float %q) {
28+
entry:
29+
%ret = atomicrmw fadd float* %p, float %q syncscope("agent") monotonic, align 4
30+
ret void
31+
}
32+
33+
; GFX90A-CAS-LABEL: atomic_add_cas_workgroup:
34+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
35+
; GFX90A-CAS: s_cbranch_execnz
36+
define dso_local void @atomic_add_cas_workgroup(float* %p, float %q) {
37+
entry:
38+
%ret = atomicrmw fadd float* %p, float %q syncscope("workgroup") monotonic, align 4
39+
ret void
40+
}
41+
42+
; GFX90A-CAS-LABEL: atomic_add_cas_wavefront:
43+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
44+
; GFX90A-CAS: s_cbranch_execnz
45+
define dso_local void @atomic_add_cas_wavefront(float* %p, float %q) {
46+
entry:
47+
%ret = atomicrmw fadd float* %p, float %q syncscope("wavefront") monotonic, align 4
48+
ret void
49+
}
50+
51+
; GFX90A-CAS-LABEL: atomic_add_cas_singlethread:
52+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
53+
; GFX90A-CAS: s_cbranch_execnz
54+
define dso_local void @atomic_add_cas_singlethread(float* %p, float %q) {
55+
entry:
56+
%ret = atomicrmw fadd float* %p, float %q syncscope("singlethread") monotonic, align 4
57+
ret void
58+
}
59+
60+
; GFX90A-CAS-LABEL: atomic_add_cas_one_as:
61+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
62+
; GFX90A-CAS: s_cbranch_execnz
63+
define dso_local void @atomic_add_cas_one_as(float* %p, float %q) {
64+
entry:
65+
%ret = atomicrmw fadd float* %p, float %q syncscope("one-as") monotonic, align 4
66+
ret void
67+
}
68+
69+
; GFX90A-CAS-LABEL: atomic_add_cas_agent_one_as:
70+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
71+
; GFX90A-CAS: s_cbranch_execnz
72+
define dso_local void @atomic_add_cas_agent_one_as(float* %p, float %q) {
73+
entry:
74+
%ret = atomicrmw fadd float* %p, float %q syncscope("agent-one-as") monotonic, align 4
75+
ret void
76+
}
77+
78+
; GFX90A-CAS-LABEL: atomic_add_cas_workgroup_one_as:
79+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
80+
; GFX90A-CAS: s_cbranch_execnz
81+
define dso_local void @atomic_add_cas_workgroup_one_as(float* %p, float %q) {
82+
entry:
83+
%ret = atomicrmw fadd float* %p, float %q syncscope("workgroup-one-as") monotonic, align 4
84+
ret void
85+
}
86+
87+
; GFX90A-CAS-LABEL: atomic_add_cas_wavefront_one_as:
88+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
89+
; GFX90A-CAS: s_cbranch_execnz
90+
define dso_local void @atomic_add_cas_wavefront_one_as(float* %p, float %q) {
91+
entry:
92+
%ret = atomicrmw fadd float* %p, float %q syncscope("wavefront-one-as") monotonic, align 4
93+
ret void
94+
}
95+
96+
; GFX90A-CAS-LABEL: atomic_add_cas_singlethread_one_as:
97+
; GFX90A-CAS: flat_atomic_cmpswap v3, v[0:1], v[4:5] glc
98+
; GFX90A-CAS: s_cbranch_execnz
99+
define dso_local void @atomic_add_cas_singlethread_one_as(float* %p, float %q) {
100+
entry:
101+
%ret = atomicrmw fadd float* %p, float %q syncscope("singlethread-one-as") monotonic, align 4
102+
ret void
103+
}

‎llvm/test/CodeGen/AMDGPU/llc-pipeline.ll

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,11 @@
4444
; GCN-O0-NEXT: Lower OpenCL enqueued blocks
4545
; GCN-O0-NEXT: Lower uses of LDS variables from non-kernel functions
4646
; GCN-O0-NEXT: FunctionPass Manager
47+
; GCN-O0-NEXT: Dominator Tree Construction
48+
; GCN-O0-NEXT: Natural Loop Information
49+
; GCN-O0-NEXT: Lazy Branch Probability Analysis
50+
; GCN-O0-NEXT: Lazy Block Frequency Analysis
51+
; GCN-O0-NEXT: Optimization Remark Emitter
4752
; GCN-O0-NEXT: Expand Atomic instructions
4853
; GCN-O0-NEXT: Lower constant intrinsics
4954
; GCN-O0-NEXT: Remove unreachable blocks from the CFG
@@ -180,6 +185,11 @@
180185
; GCN-O1-NEXT: Lower uses of LDS variables from non-kernel functions
181186
; GCN-O1-NEXT: FunctionPass Manager
182187
; GCN-O1-NEXT: Infer address spaces
188+
; GCN-O1-NEXT: Dominator Tree Construction
189+
; GCN-O1-NEXT: Natural Loop Information
190+
; GCN-O1-NEXT: Lazy Branch Probability Analysis
191+
; GCN-O1-NEXT: Lazy Block Frequency Analysis
192+
; GCN-O1-NEXT: Optimization Remark Emitter
183193
; GCN-O1-NEXT: Expand Atomic instructions
184194
; GCN-O1-NEXT: AMDGPU Promote Alloca
185195
; GCN-O1-NEXT: Dominator Tree Construction
@@ -431,6 +441,11 @@
431441
; GCN-O1-OPTS-NEXT: Lower uses of LDS variables from non-kernel functions
432442
; GCN-O1-OPTS-NEXT: FunctionPass Manager
433443
; GCN-O1-OPTS-NEXT: Infer address spaces
444+
; GCN-O1-OPTS-NEXT: Dominator Tree Construction
445+
; GCN-O1-OPTS-NEXT: Natural Loop Information
446+
; GCN-O1-OPTS-NEXT: Lazy Branch Probability Analysis
447+
; GCN-O1-OPTS-NEXT: Lazy Block Frequency Analysis
448+
; GCN-O1-OPTS-NEXT: Optimization Remark Emitter
434449
; GCN-O1-OPTS-NEXT: Expand Atomic instructions
435450
; GCN-O1-OPTS-NEXT: AMDGPU Promote Alloca
436451
; GCN-O1-OPTS-NEXT: Dominator Tree Construction
@@ -715,6 +730,11 @@
715730
; GCN-O2-NEXT: Lower uses of LDS variables from non-kernel functions
716731
; GCN-O2-NEXT: FunctionPass Manager
717732
; GCN-O2-NEXT: Infer address spaces
733+
; GCN-O2-NEXT: Dominator Tree Construction
734+
; GCN-O2-NEXT: Natural Loop Information
735+
; GCN-O2-NEXT: Lazy Branch Probability Analysis
736+
; GCN-O2-NEXT: Lazy Block Frequency Analysis
737+
; GCN-O2-NEXT: Optimization Remark Emitter
718738
; GCN-O2-NEXT: Expand Atomic instructions
719739
; GCN-O2-NEXT: AMDGPU Promote Alloca
720740
; GCN-O2-NEXT: Dominator Tree Construction
@@ -1001,6 +1021,11 @@
10011021
; GCN-O3-NEXT: Lower uses of LDS variables from non-kernel functions
10021022
; GCN-O3-NEXT: FunctionPass Manager
10031023
; GCN-O3-NEXT: Infer address spaces
1024+
; GCN-O3-NEXT: Dominator Tree Construction
1025+
; GCN-O3-NEXT: Natural Loop Information
1026+
; GCN-O3-NEXT: Lazy Branch Probability Analysis
1027+
; GCN-O3-NEXT: Lazy Block Frequency Analysis
1028+
; GCN-O3-NEXT: Optimization Remark Emitter
10041029
; GCN-O3-NEXT: Expand Atomic instructions
10051030
; GCN-O3-NEXT: AMDGPU Promote Alloca
10061031
; GCN-O3-NEXT: Dominator Tree Construction

‎llvm/test/CodeGen/ARM/O3-pipeline.ll

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,11 @@
55
; CHECK: ModulePass Manager
66
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
77
; CHECK-NEXT: FunctionPass Manager
8+
; CHECK-NEXT: Dominator Tree Construction
9+
; CHECK-NEXT: Natural Loop Information
10+
; CHECK-NEXT: Lazy Branch Probability Analysis
11+
; CHECK-NEXT: Lazy Block Frequency Analysis
12+
; CHECK-NEXT: Optimization Remark Emitter
813
; CHECK-NEXT: Expand Atomic instructions
914
; CHECK-NEXT: Simplify the CFG
1015
; CHECK-NEXT: Dominator Tree Construction

‎llvm/test/CodeGen/PowerPC/O3-pipeline.ll

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,16 +8,21 @@
88
; CHECK-NEXT: Target Pass Configuration
99
; CHECK-NEXT: Machine Module Information
1010
; CHECK-NEXT: Target Transform Information
11+
; CHECK-NEXT: Profile summary info
1112
; CHECK-NEXT: Assumption Cache Tracker
1213
; CHECK-NEXT: Type-Based Alias Analysis
1314
; CHECK-NEXT: Scoped NoAlias Alias Analysis
14-
; CHECK-NEXT: Profile summary info
1515
; CHECK-NEXT: Create Garbage Collector Module Metadata
1616
; CHECK-NEXT: Machine Branch Probability Analysis
1717
; CHECK-NEXT: ModulePass Manager
1818
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
1919
; CHECK-NEXT: FunctionPass Manager
2020
; CHECK-NEXT: Convert i1 constants to i32/i64 if they are returned
21+
; CHECK-NEXT: Dominator Tree Construction
22+
; CHECK-NEXT: Natural Loop Information
23+
; CHECK-NEXT: Lazy Branch Probability Analysis
24+
; CHECK-NEXT: Lazy Block Frequency Analysis
25+
; CHECK-NEXT: Optimization Remark Emitter
2126
; CHECK-NEXT: Expand Atomic instructions
2227
; CHECK-NEXT: PPC Lower MASS Entries
2328
; CHECK-NEXT: FunctionPass Manager
@@ -206,4 +211,5 @@
206211

207212
define void @f() {
208213
ret void
209-
}
214+
}
215+

0 commit comments

Comments
 (0)