//===-- NVPTXLowerUnreachable.cpp - Lower unreachables to exit =====--===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // // PTX does not have a notion of `unreachable`, which results in emitted basic // blocks having an edge to the next block: // // block1: // call @does_not_return(); // // unreachable // block2: // // ptxas will create a CFG edge from block1 to block2 // // This may result in significant changes to the control flow graph, e.g., when // LLVM moves unreachable blocks to the end of the function. That's a problem // in the context of divergent control flow, as `ptxas` uses the CFG to // determine divergent regions, and some intructions may not be executed // divergently. // // For example, `bar.sync` is not allowed to be executed divergently on Pascal // or earlier. If we start with the following: // // entry: // // start of divergent region // @%p0 bra cont; // @%p1 bra unlikely; // ... // bra.uni cont; // unlikely: // ... // // unreachable // cont: // // end of divergent region // bar.sync 0; // bra.uni exit; // exit: // ret; // // it is transformed by the branch-folder and block-placement passes to: // // entry: // // start of divergent region // @%p0 bra cont; // @%p1 bra unlikely; // ... // bra.uni cont; // cont: // bar.sync 0; // bra.uni exit; // unlikely: // ... // // unreachable // exit: // // end of divergent region // ret; // // After moving the `unlikely` block to the end of the function, it has an edge // to the `exit` block, which widens the divergent region and makes the // `bar.sync` instruction happen divergently. // // To work around this, we add an `exit` instruction before every `unreachable`, // as `ptxas` understands that exit terminates the CFG. We do only do this if // `unreachable` is not lowered to `trap`, which has the same effect (although // with current versions of `ptxas` only because it is emited as `trap; exit;`). // //===----------------------------------------------------------------------===// #include "NVPTX.h" #include "llvm/IR/Function.h" #include "llvm/IR/InlineAsm.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/Type.h" #include "llvm/Pass.h" usingnamespacellvm; namespace llvm { void initializeNVPTXLowerUnreachablePass(PassRegistry &); } namespace { class NVPTXLowerUnreachable : public FunctionPass { … }; } // namespace char NVPTXLowerUnreachable::ID = …; INITIALIZE_PASS(…) StringRef NVPTXLowerUnreachable::getPassName() const { … } // ============================================================================= // Returns whether a `trap` intrinsic should be emitted before I. // // This is a copy of the logic in SelectionDAGBuilder::visitUnreachable(). // ============================================================================= bool NVPTXLowerUnreachable::isLoweredToTrap(const UnreachableInst &I) const { … } // ============================================================================= // Main function for this pass. // ============================================================================= bool NVPTXLowerUnreachable::runOnFunction(Function &F) { … } FunctionPass *llvm::createNVPTXLowerUnreachablePass(bool TrapUnreachable, bool NoTrapAfterNoreturn) { … }