27#include "mlir/IR/Builders.h"
32#include "llvm/ADT/DenseMap.h"
33#include "llvm/ADT/SmallVector.h"
34#include "llvm/TargetParser/Triple.h"
40#define GEN_PASS_DEF_CIREHABILOWERING
41#include "clang/CIR/Dialect/Passes.h.inc"
52static cir::FuncOp getOrCreateRuntimeFuncDecl(mlir::ModuleOp mod,
55 cir::FuncType funcTy) {
56 if (
auto existing = mod.lookupSymbol<cir::FuncOp>(
name))
59 mlir::OpBuilder builder(mod.getContext());
60 builder.setInsertionPointToEnd(mod.getBody());
61 auto funcOp = cir::FuncOp::create(builder, loc,
name, funcTy);
62 funcOp.setLinkage(cir::GlobalLinkageKind::ExternalLinkage);
75 explicit EHABILowering(mlir::ModuleOp mod)
76 : mod(mod), ctx(mod.getContext()), builder(ctx) {}
77 virtual ~EHABILowering() =
default;
80 virtual mlir::LogicalResult
run() = 0;
84 mlir::MLIRContext *ctx;
85 mlir::OpBuilder builder;
99class ItaniumEHLowering :
public EHABILowering {
101 using EHABILowering::EHABILowering;
102 mlir::LogicalResult
run()
override;
107 using EhTokenMap = DenseMap<mlir::Value, std::pair<mlir::Value, mlir::Value>>;
109 cir::VoidType voidType;
110 cir::PointerType voidPtrType;
111 cir::PointerType u8PtrType;
112 cir::IntType u32Type;
116 cir::FuncOp personalityFunc;
117 cir::FuncOp beginCatchFunc;
118 cir::FuncOp endCatchFunc;
120 constexpr const static ::llvm::StringLiteral kGxxPersonality =
121 "__gxx_personality_v0";
123 void ensureRuntimeDecls(mlir::Location loc);
124 mlir::LogicalResult lowerFunc(cir::FuncOp funcOp);
125 void lowerEhInitiate(cir::EhInitiateOp initiateOp, EhTokenMap &ehTokenMap,
126 SmallVectorImpl<mlir::Operation *> &deadOps);
127 void lowerDispatch(cir::EhDispatchOp dispatch, mlir::Value exnPtr,
129 SmallVectorImpl<mlir::Operation *> &deadOps);
133mlir::LogicalResult ItaniumEHLowering::run() {
136 voidType = cir::VoidType::get(ctx);
137 voidPtrType = cir::PointerType::get(voidType);
138 auto u8Type = cir::IntType::get(ctx, 8,
false);
139 u8PtrType = cir::PointerType::get(u8Type);
140 u32Type = cir::IntType::get(ctx, 32,
false);
142 for (cir::FuncOp funcOp : mod.getOps<cir::FuncOp>()) {
143 if (mlir::failed(lowerFunc(funcOp)))
144 return mlir::failure();
146 return mlir::success();
151void ItaniumEHLowering::ensureRuntimeDecls(mlir::Location loc) {
154 if (!personalityFunc) {
155 auto s32Type = cir::IntType::get(ctx, 32,
true);
156 auto personalityFuncTy = cir::FuncType::get({}, s32Type,
true);
157 personalityFunc = getOrCreateRuntimeFuncDecl(mod, loc, kGxxPersonality,
161 if (!beginCatchFunc) {
162 auto beginCatchFuncTy =
163 cir::FuncType::get({voidPtrType}, u8PtrType,
false);
164 beginCatchFunc = getOrCreateRuntimeFuncDecl(mod, loc,
"__cxa_begin_catch",
169 auto endCatchFuncTy = cir::FuncType::get({}, voidType,
false);
171 getOrCreateRuntimeFuncDecl(mod, loc,
"__cxa_end_catch", endCatchFuncTy);
176mlir::LogicalResult ItaniumEHLowering::lowerFunc(cir::FuncOp funcOp) {
177 if (funcOp.isDeclaration())
178 return mlir::success();
184 SmallVector<cir::EhInitiateOp> initiateOps;
185 funcOp.walk([&](cir::EhInitiateOp op) { initiateOps.push_back(op); });
186 if (initiateOps.empty())
187 return mlir::success();
189 ensureRuntimeDecls(funcOp.getLoc());
196 if (!funcOp.getPersonality())
197 funcOp.setPersonality(kGxxPersonality);
204 EhTokenMap ehTokenMap;
205 SmallVector<mlir::Operation *> deadOps;
206 for (cir::EhInitiateOp initiateOp : initiateOps)
207 lowerEhInitiate(initiateOp, ehTokenMap, deadOps);
211 for (mlir::Operation *op : deadOps)
216 for (mlir::Block &block : funcOp.getBody()) {
217 for (
int i = block.getNumArguments() - 1; i >= 0; --i) {
218 if (mlir::isa<cir::EhTokenType>(block.getArgument(i).getType()))
219 block.eraseArgument(i);
223 return mlir::success();
252void ItaniumEHLowering::lowerEhInitiate(
253 cir::EhInitiateOp initiateOp, EhTokenMap &ehTokenMap,
254 SmallVectorImpl<mlir::Operation *> &deadOps) {
255 mlir::Value rootToken = initiateOp.getEhToken();
259 builder.setInsertionPoint(initiateOp);
260 auto inflightOp = cir::EhInflightOp::create(
261 builder, initiateOp.getLoc(), initiateOp.getCleanup(),
264 ehTokenMap[rootToken] = {inflightOp.getExceptionPtr(),
265 inflightOp.getTypeId()};
271 SmallVector<mlir::Value> worklist;
272 SmallPtrSet<mlir::Value, 8> visited;
273 worklist.push_back(rootToken);
275 while (!worklist.empty()) {
276 mlir::Value current = worklist.pop_back_val();
277 if (!visited.insert(current).second)
282 SmallVector<mlir::Operation *> users;
283 for (mlir::OpOperand &use : current.getUses())
284 users.push_back(use.getOwner());
288 for (mlir::Operation *user : users) {
294 for (
unsigned s = 0;
s < user->getNumSuccessors(); ++
s) {
295 mlir::Block *succ = user->getSuccessor(
s);
296 for (mlir::BlockArgument arg : succ->getArguments()) {
297 if (!mlir::isa<cir::EhTokenType>(
arg.getType()))
299 if (!ehTokenMap.count(arg)) {
300 mlir::Value ptrArg = succ->addArgument(voidPtrType,
arg.getLoc());
301 mlir::Value u32Arg = succ->addArgument(u32Type,
arg.getLoc());
302 ehTokenMap[
arg] = {ptrArg, u32Arg};
304 worklist.push_back(arg);
308 if (
auto op = mlir::dyn_cast<cir::BeginCleanupOp>(user)) {
311 for (
auto &tokenUsers :
312 llvm::make_early_inc_range(op.getCleanupToken().getUses())) {
314 mlir::dyn_cast<cir::EndCleanupOp>(tokenUsers.getOwner()))
318 }
else if (
auto op = mlir::dyn_cast<cir::BeginCatchOp>(user)) {
321 for (
auto &tokenUsers :
322 llvm::make_early_inc_range(op.getCatchToken().getUses())) {
324 mlir::dyn_cast<cir::EndCatchOp>(tokenUsers.getOwner())) {
325 builder.setInsertionPoint(endOp);
326 cir::CallOp::create(builder, endOp.getLoc(),
327 mlir::FlatSymbolRefAttr::get(endCatchFunc),
328 voidType, mlir::ValueRange{});
333 auto [exnPtr, typeId] = ehTokenMap.lookup(op.getEhToken());
334 builder.setInsertionPoint(op);
335 auto callOp = cir::CallOp::create(
336 builder, op.getLoc(), mlir::FlatSymbolRefAttr::get(beginCatchFunc),
337 u8PtrType, mlir::ValueRange{exnPtr});
338 mlir::Value castResult = callOp.getResult();
339 mlir::Type expectedPtrType = op.getExnPtr().getType();
340 if (castResult.getType() != expectedPtrType)
342 cir::CastOp::create(builder, op.getLoc(), expectedPtrType,
343 cir::CastKind::bitcast, callOp.getResult());
344 op.getExnPtr().replaceAllUsesWith(castResult);
346 }
else if (
auto op = mlir::dyn_cast<cir::EhDispatchOp>(user)) {
348 mlir::ArrayAttr catchTypes = op.getCatchTypesAttr();
349 if (catchTypes && catchTypes.size() > 0) {
350 SmallVector<mlir::Attribute> typeSymbols;
351 for (mlir::Attribute attr : catchTypes)
352 typeSymbols.push_back(
353 mlir::cast<cir::GlobalViewAttr>(attr).getSymbol());
354 inflightOp.setCatchTypeListAttr(builder.getArrayAttr(typeSymbols));
359 if (!llvm::is_contained(deadOps, op.getOperation())) {
360 auto [exnPtr, typeId] = ehTokenMap.lookup(op.getEhToken());
361 lowerDispatch(op, exnPtr, typeId, deadOps);
363 }
else if (
auto op = mlir::dyn_cast<cir::ResumeOp>(user)) {
364 auto [exnPtr, typeId] = ehTokenMap.lookup(op.getEhToken());
365 builder.setInsertionPoint(op);
366 cir::ResumeFlatOp::create(builder, op.getLoc(), exnPtr, typeId);
368 }
else if (
auto op = mlir::dyn_cast<cir::BrOp>(user)) {
370 SmallVector<mlir::Value> newOperands;
371 bool changed =
false;
372 for (mlir::Value operand : op.getDestOperands()) {
373 auto it = ehTokenMap.find(operand);
374 if (it != ehTokenMap.end()) {
375 newOperands.push_back(it->second.first);
376 newOperands.push_back(it->second.second);
379 newOperands.push_back(operand);
383 builder.setInsertionPoint(op);
384 cir::BrOp::create(builder, op.getLoc(), op.getDest(), newOperands);
397void ItaniumEHLowering::lowerDispatch(
398 cir::EhDispatchOp dispatch, mlir::Value exnPtr, mlir::Value typeId,
399 SmallVectorImpl<mlir::Operation *> &deadOps) {
400 mlir::Location dispLoc = dispatch.getLoc();
401 mlir::Block *defaultDest = dispatch.getDefaultDestination();
402 mlir::ArrayAttr catchTypes = dispatch.getCatchTypesAttr();
403 mlir::SuccessorRange catchDests = dispatch.getCatchDestinations();
404 mlir::Block *dispatchBlock = dispatch->getBlock();
409 if (!catchTypes || catchTypes.empty()) {
411 builder.setInsertionPoint(dispatch);
412 cir::BrOp::create(builder, dispLoc, defaultDest,
413 mlir::ValueRange{exnPtr, typeId});
415 unsigned numCatches = catchTypes.size();
421 mlir::Block *
insertBefore = dispatchBlock->getNextNode();
422 mlir::Block *falseDest = defaultDest;
423 mlir::Block *firstCmpBlock =
nullptr;
424 for (
int i = numCatches - 1; i >= 0; --i) {
425 auto *cmpBlock = builder.createBlock(insertBefore, {voidPtrType, u32Type},
428 mlir::Value cmpExnPtr = cmpBlock->getArgument(0);
429 mlir::Value cmpTypeId = cmpBlock->getArgument(1);
431 auto globalView = mlir::cast<cir::GlobalViewAttr>(catchTypes[i]);
433 cir::EhTypeIdOp::create(builder, dispLoc, globalView.getSymbol());
434 auto cmpOp = cir::CmpOp::create(builder, dispLoc, cir::CmpOpKind::eq,
435 cmpTypeId, ehTypeIdOp.getTypeId());
437 cir::BrCondOp::create(builder, dispLoc, cmpOp, catchDests[i], falseDest,
438 mlir::ValueRange{cmpExnPtr, cmpTypeId},
439 mlir::ValueRange{cmpExnPtr, cmpTypeId});
442 falseDest = cmpBlock;
443 firstCmpBlock = cmpBlock;
447 builder.setInsertionPoint(dispatch);
448 cir::BrOp::create(builder, dispLoc, firstCmpBlock,
449 mlir::ValueRange{exnPtr, typeId});
455 deadOps.push_back(dispatch);
462struct CIREHABILoweringPass
463 :
public impl::CIREHABILoweringBase<CIREHABILoweringPass> {
464 CIREHABILoweringPass() =
default;
465 void runOnOperation()
override;
468void CIREHABILoweringPass::runOnOperation() {
469 auto mod = mlir::cast<mlir::ModuleOp>(getOperation());
474 auto tripleAttr = mlir::dyn_cast_if_present<mlir::StringAttr>(
475 mod->getAttr(cir::CIRDialect::getTripleAttrName()));
477 mod.emitError(
"Module has no target triple");
484 llvm::Triple triple(tripleAttr.getValue());
485 std::unique_ptr<EHABILowering> lowering;
486 if (triple.isWindowsMSVCEnvironment()) {
488 "EH ABI lowering is not yet implemented for the Microsoft ABI");
489 return signalPassFailure();
491 lowering = std::make_unique<ItaniumEHLowering>(mod);
494 if (mlir::failed(lowering->run()))
495 return signalPassFailure();
501 return std::make_unique<CIREHABILoweringPass>();
__device__ __2f16 float __ockl_bool s
std::unique_ptr< Pass > createCIREHABILoweringPass()
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)