1//===-- NVPTXLowerUnreachable.cpp - Lower unreachables to exit =====--===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// PTX does not have a notion of `unreachable`, which results in emitted basic
10// blocks having an edge to the next block:
11//
12// block1:
13// call @does_not_return();
14// // unreachable
15// block2:
16// // ptxas will create a CFG edge from block1 to block2
17//
18// This may result in significant changes to the control flow graph, e.g., when
19// LLVM moves unreachable blocks to the end of the function. That's a problem
20// in the context of divergent control flow, as `ptxas` uses the CFG to
21// determine divergent regions, and some intructions may not be executed
22// divergently.
23//
24// For example, `bar.sync` is not allowed to be executed divergently on Pascal
25// or earlier. If we start with the following:
26//
27// entry:
28// // start of divergent region
29// @%p0 bra cont;
30// @%p1 bra unlikely;
31// ...
32// bra.uni cont;
33// unlikely:
34// ...
35// // unreachable
36// cont:
37// // end of divergent region
38// bar.sync 0;
39// bra.uni exit;
40// exit:
41// ret;
42//
43// it is transformed by the branch-folder and block-placement passes to:
44//
45// entry:
46// // start of divergent region
47// @%p0 bra cont;
48// @%p1 bra unlikely;
49// ...
50// bra.uni cont;
51// cont:
52// bar.sync 0;
53// bra.uni exit;
54// unlikely:
55// ...
56// // unreachable
57// exit:
58// // end of divergent region
59// ret;
60//
61// After moving the `unlikely` block to the end of the function, it has an edge
62// to the `exit` block, which widens the divergent region and makes the
63// `bar.sync` instruction happen divergently.
64//
65// To work around this, we add an `exit` instruction before every `unreachable`,
66// as `ptxas` understands that exit terminates the CFG. We do only do this if
67// `unreachable` is not lowered to `trap`, which has the same effect (although
68// with current versions of `ptxas` only because it is emited as `trap; exit;`).
69//
70//===----------------------------------------------------------------------===//
71
72#include "NVPTX.h"
73#include "llvm/IR/Function.h"
74#include "llvm/IR/InlineAsm.h"
75#include "llvm/IR/Instructions.h"
76#include "llvm/IR/Type.h"
77#include "llvm/Pass.h"
78
79using namespace llvm;
80
81// =============================================================================
82// Returns whether a `trap` intrinsic would be emitted before I.
83//
84// This is a copy of the logic in SelectionDAGBuilder::visitUnreachable().
85// =============================================================================
86static bool isLoweredToTrap(const UnreachableInst &I, bool TrapUnreachable,
87 bool NoTrapAfterNoreturn) {
88 if (const auto *Call = dyn_cast_or_null<CallInst>(Val: I.getPrevNode())) {
89 // We've already emitted a non-continuable trap.
90 if (Call->isNonContinuableTrap())
91 return true;
92
93 // No traps are emitted for calls that do not return
94 // when this option is enabled.
95 if (NoTrapAfterNoreturn && Call->doesNotReturn())
96 return false;
97 }
98
99 // In all other cases, we will generate a trap if TrapUnreachable is set.
100 return TrapUnreachable;
101}
102
103// =============================================================================
104// Main function for this pass.
105// =============================================================================
106static bool lowerUnreachable(Function &F, bool TrapUnreachable,
107 bool NoTrapAfterNoreturn) {
108 // Early out iff isLoweredToTrap() always returns true.
109 if (TrapUnreachable && !NoTrapAfterNoreturn)
110 return false;
111
112 LLVMContext &C = F.getContext();
113 FunctionType *ExitFTy = FunctionType::get(Result: Type::getVoidTy(C), isVarArg: false);
114 InlineAsm *Exit = InlineAsm::get(Ty: ExitFTy, AsmString: "exit;", Constraints: "", hasSideEffects: true);
115
116 bool Changed = false;
117 for (auto &BB : F)
118 for (auto &I : BB) {
119 if (auto unreachableInst = dyn_cast<UnreachableInst>(Val: &I)) {
120 if (isLoweredToTrap(I: *unreachableInst, TrapUnreachable,
121 NoTrapAfterNoreturn))
122 continue; // trap is emitted as `trap; exit;`.
123 CallInst::Create(Ty: ExitFTy, F: Exit, NameStr: "", InsertBefore: unreachableInst->getIterator());
124 Changed = true;
125 }
126 }
127 return Changed;
128}
129
130namespace {
131class NVPTXLowerUnreachableLegacyPass : public FunctionPass {
132 StringRef getPassName() const override {
133 return "add an exit instruction before every unreachable";
134 }
135
136 bool runOnFunction(Function &F) override {
137 if (skipFunction(F))
138 return false;
139 return lowerUnreachable(F, TrapUnreachable, NoTrapAfterNoreturn);
140 }
141
142 void getAnalysisUsage(AnalysisUsage &AU) const override {
143 AU.setPreservesCFG();
144 }
145
146public:
147 static char ID; // Pass identification, replacement for typeid
148 NVPTXLowerUnreachableLegacyPass(bool TrapUnreachable,
149 bool NoTrapAfterNoreturn)
150 : FunctionPass(ID), TrapUnreachable(TrapUnreachable),
151 NoTrapAfterNoreturn(NoTrapAfterNoreturn) {}
152
153private:
154 bool TrapUnreachable;
155 bool NoTrapAfterNoreturn;
156};
157} // namespace
158
159char NVPTXLowerUnreachableLegacyPass::ID = 0;
160
161INITIALIZE_PASS(NVPTXLowerUnreachableLegacyPass, "nvptx-lower-unreachable",
162 "Lower Unreachable", false, false)
163
164FunctionPass *
165llvm::createNVPTXLowerUnreachableLegacyPass(bool TrapUnreachable,
166 bool NoTrapAfterNoreturn) {
167 return new NVPTXLowerUnreachableLegacyPass(TrapUnreachable,
168 NoTrapAfterNoreturn);
169}
170
171PreservedAnalyses NVPTXLowerUnreachablePass::run(Function &F,
172 FunctionAnalysisManager &FAM) {
173 if (!lowerUnreachable(F, TrapUnreachable, NoTrapAfterNoreturn))
174 return PreservedAnalyses::all();
175 return PreservedAnalyses::none().preserveSet<CFGAnalyses>();
176}
177