10#include "mlir/IR/Attributes.h"
11#include "mlir/IR/BuiltinAttributeInterfaces.h"
12#include "mlir/IR/IRMapping.h"
13#include "mlir/IR/Location.h"
14#include "mlir/IR/Value.h"
32#include "llvm/ADT/StringRef.h"
33#include "llvm/ADT/TypeSwitch.h"
34#include "llvm/IR/Instructions.h"
35#include "llvm/Support/ErrorHandling.h"
36#include "llvm/Support/MemoryBuffer.h"
37#include "llvm/Support/Path.h"
38#include "llvm/Support/VirtualFileSystem.h"
47#define GEN_PASS_DEF_LOWERINGPREPARE
48#include "clang/CIR/Dialect/Passes.h.inc"
52 SmallString<128> fileName;
54 if (mlirModule.getSymName())
55 fileName = llvm::sys::path::filename(mlirModule.getSymName()->str());
60 for (
size_t i = 0; i < fileName.size(); ++i) {
72 mlir::SymbolRefAttr sym = llvm::dyn_cast_if_present<mlir::SymbolRefAttr>(
73 callOp.getCallableForCallee());
76 return dyn_cast_or_null<cir::FuncOp>(
77 mlir::SymbolTable::lookupNearestSymbolFrom(callOp, sym));
81struct LoweringPreparePass
82 :
public impl::LoweringPrepareBase<LoweringPreparePass> {
83 LoweringPreparePass() =
default;
84 void runOnOperation()
override;
86 void runOnOp(mlir::Operation *op, mlir::SymbolTableCollection &symbolTables);
87 void lowerCastOp(cir::CastOp op);
88 void lowerComplexDivOp(cir::ComplexDivOp op);
89 void lowerComplexMulOp(cir::ComplexMulOp op);
90 void lowerUnaryOp(cir::UnaryOpInterface op);
91 void lowerGlobalOp(cir::GlobalOp op);
92 void lowerThreeWayCmpOp(cir::CmpThreeWayOp op);
93 void lowerArrayDtor(cir::ArrayDtor op);
94 void lowerArrayCtor(cir::ArrayCtor op);
95 void lowerTrivialCopyCall(cir::CallOp op);
96 void lowerStoreOfConstAggregate(cir::StoreOp op,
97 mlir::SymbolTableCollection &symbolTables);
98 void lowerLocalInitOp(cir::LocalInitOp op,
99 mlir::SymbolTableCollection &symbolTables);
109 getOrCreateConstAggregateGlobal(CIRBaseBuilderTy &builder,
110 mlir::SymbolTableCollection &symbolTables,
111 mlir::Location loc, llvm::StringRef baseName,
112 mlir::Type ty, mlir::TypedAttr constant);
115 cir::FuncOp buildCXXGlobalVarDeclInitFunc(cir::GlobalOp op);
118 cir::FuncOp getOrCreateDtorFunc(CIRBaseBuilderTy &builder, cir::GlobalOp op,
119 mlir::Region &dtorRegion,
120 cir::CallOp &dtorCall);
123 void buildCXXGlobalInitFunc();
126 void buildGlobalCtorDtorList();
128 cir::FuncOp buildRuntimeFunction(
129 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
131 cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage);
133 cir::GlobalOp getOrCreateRuntimeVariable(
134 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
136 cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage,
137 cir::VisibilityKind visibility = cir::VisibilityKind::Default);
143 llvm::StringMap<FuncOp> cudaKernelMap;
147 void buildCUDAModuleCtor();
148 std::optional<FuncOp> buildCUDAModuleDtor();
149 std::optional<FuncOp> buildCUDARegisterGlobals();
150 void buildCUDARegisterGlobalFunctions(cir::CIRBaseBuilderTy &builder,
151 FuncOp regGlobalFunc);
154 void handleStaticLocal(cir::GlobalOp globalOp, cir::LocalInitOp localInitOp);
163 cir::GlobalOp createGuardGlobalOp(CIRBaseBuilderTy &builder,
164 mlir::Location loc, llvm::StringRef name,
165 cir::IntType guardTy,
166 cir::GlobalLinkageKind linkage);
169 cir::GlobalOp getStaticLocalDeclGuardAddress(llvm::StringRef globalSymName) {
170 auto it = staticLocalDeclGuardMap.find(globalSymName);
171 if (it != staticLocalDeclGuardMap.end())
177 void setStaticLocalDeclGuardAddress(llvm::StringRef globalSymName,
178 cir::GlobalOp guard) {
179 staticLocalDeclGuardMap[globalSymName] = guard;
183 cir::GlobalOp getOrCreateStaticLocalDeclGuardAddress(
184 CIRBaseBuilderTy &builder, cir::GlobalOp globalOp,
185 cir::ASTVarDeclInterface varDecl, cir::IntType guardTy,
186 clang::CharUnits guardAlignment) {
187 llvm::StringRef globalSymName = globalOp.getSymName();
188 cir::GlobalOp guard = getStaticLocalDeclGuardAddress(globalSymName);
191 llvm::StringRef guardName =
192 globalOp.getStaticLocalGuard()->getName().getValue();
195 guard = createGuardGlobalOp(builder, globalOp->getLoc(), guardName,
196 guardTy, globalOp.getLinkage());
197 guard.setInitialValueAttr(cir::IntAttr::get(guardTy, 0));
198 guard.setDSOLocal(globalOp.getDsoLocal());
199 guard.setAlignment(guardAlignment.
getAsAlign().value());
200 guard.setTlsModel(globalOp.getTlsModel());
206 bool hasComdat = globalOp.getComdat();
207 const llvm::Triple &triple = astCtx->getTargetInfo().getTriple();
208 if (!
varDecl.isLocalVarDecl() && hasComdat &&
209 (triple.isOSBinFormatELF() || triple.isOSBinFormatWasm())) {
210 globalOp->emitError(
"NYI: guard COMDAT for non-local variables");
212 }
else if (hasComdat && globalOp.isWeakForLinker()) {
213 guard.setComdat(
true);
216 setStaticLocalDeclGuardAddress(globalSymName, guard);
225 clang::ASTContext *astCtx;
228 mlir::ModuleOp mlirModule;
231 llvm::StringMap<uint32_t> dynamicInitializerNames;
232 llvm::SmallVector<cir::FuncOp> dynamicInitializers;
235 llvm::StringMap<cir::GlobalOp> staticLocalDeclGuardMap;
237 llvm::StringMap<llvm::SmallVector<cir::GlobalOp, 1>> constAggregateGlobals;
240 llvm::SmallVector<std::pair<std::string, uint32_t>, 4> globalCtorList;
242 llvm::SmallVector<std::pair<std::string, uint32_t>, 4> globalDtorList;
246 bool useARMGuardVarABI()
const {
247 switch (astCtx->getCXXABIKind()) {
248 case clang::TargetCXXABI::GenericARM:
249 case clang::TargetCXXABI::iOS:
250 case clang::TargetCXXABI::WatchOS:
251 case clang::TargetCXXABI::GenericAArch64:
252 case clang::TargetCXXABI::WebAssembly:
259 void emitGlobalGuardedDtorRegion(CIRBaseBuilderTy &builder,
260 cir::GlobalOp global,
261 mlir::Region &dtorRegion,
bool tls,
262 mlir::Block &entryBB) {
264 builder.setInsertionPointToStart(&mlirModule.getBodyRegion().front());
265 cir::GlobalOp handle = getOrCreateRuntimeVariable(
266 builder,
"__dso_handle", global.getLoc(), builder.getI8Type(),
267 cir::GlobalLinkageKind::ExternalLinkage, cir::VisibilityKind::Hidden);
273 cir::CallOp dtorCall;
274 cir::FuncOp dtorFunc =
275 getOrCreateDtorFunc(builder, global, dtorRegion, dtorCall);
280 cir::PointerType voidFnPtrTy = builder.
getVoidFnPtrTy({voidPtrTy});
281 cir::PointerType handlePtrTy = builder.
getPointerTo(handle.getSymType());
283 builder.
getVoidFnTy({voidFnPtrTy, voidPtrTy, handlePtrTy});
285 llvm::StringLiteral nameAtExit =
"__cxa_atexit";
287 nameAtExit = astCtx->getTargetInfo().getTriple().isOSDarwin()
288 ? llvm::StringLiteral(
"_tlv_atexit")
289 : llvm::StringLiteral(
"__cxa_thread_atexit");
291 cir::FuncOp fnAtExit = buildRuntimeFunction(builder, nameAtExit,
292 global.getLoc(), fnAtExitType);
296 builder.setInsertionPointAfter(dtorCall);
298 auto dtorPtrTy = cir::PointerType::get(dtorFunc.getFunctionType());
299 args[0] = cir::GetGlobalOp::create(builder, dtorCall.getLoc(), dtorPtrTy,
300 dtorFunc.getSymName());
301 args[0] = cir::CastOp::create(builder, dtorCall.getLoc(), voidFnPtrTy,
302 cir::CastKind::bitcast, args[0]);
304 cir::CastOp::create(builder, dtorCall.getLoc(), voidPtrTy,
305 cir::CastKind::bitcast, dtorCall.getArgOperand(0));
306 args[2] = cir::GetGlobalOp::create(builder, handle.getLoc(), handlePtrTy,
307 handle.getSymName());
308 builder.
createCallOp(dtorCall.getLoc(), fnAtExit, args);
310 mlir::Block &dtorBlock = dtorRegion.front();
311 entryBB.getOperations().splice(entryBB.end(), dtorBlock.getOperations(),
313 std::prev(dtorBlock.end()));
319 void emitCXXGuardedInitIf(CIRBaseBuilderTy &builder, cir::GlobalOp globalOp,
320 mlir::Region &ctorRegion, mlir::Region &dtorRegion,
321 cir::ASTVarDeclInterface varDecl,
322 mlir::Value guardPtr, cir::PointerType guardPtrTy,
324 auto loc = globalOp->getLoc();
344 mlir::Block *insertBlock = builder.getInsertionBlock();
345 if (!ctorRegion.empty()) {
346 assert(ctorRegion.hasOneBlock() &&
"Enforced by MaxSizedRegion<1>");
348 mlir::Block &block = ctorRegion.front();
349 insertBlock->getOperations().splice(
350 insertBlock->end(), block.getOperations(), block.begin(),
351 std::prev(block.end()));
354 if (!dtorRegion.empty()) {
355 assert(dtorRegion.hasOneBlock() &&
"Enforced by MaxSizedRegion<1>");
357 emitGlobalGuardedDtorRegion(builder, globalOp, dtorRegion, !threadsafe,
360 builder.setInsertionPointToEnd(insertBlock);
361 ctorRegion.getBlocks().clear();
369 mlir::Value acquireResult = acquireCall.getResult();
372 loc, mlir::cast<cir::IntType>(acquireResult.getType()), 0);
373 auto shouldInit = builder.
createCompare(loc, cir::CmpOpKind::ne,
374 acquireResult, acquireZero);
379 cir::IfOp::create(builder, loc, shouldInit,
false,
380 [](mlir::OpBuilder &, mlir::Location) {});
381 mlir::OpBuilder::InsertionGuard insertGuard(builder);
382 builder.setInsertionPointToStart(&ifOp.getThenRegion().front());
396 mlir::ValueRange{guardPtr});
399 }
else if (!
varDecl.isLocalVarDecl()) {
405 globalOp->emitError(
"NYI: non-threadsafe init for non-local variables");
420 void setASTContext(clang::ASTContext *
c) { astCtx =
c; }
425cir::GlobalOp LoweringPreparePass::getOrCreateRuntimeVariable(
426 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
427 mlir::Type type, cir::GlobalLinkageKind linkage,
428 cir::VisibilityKind visibility) {
429 cir::GlobalOp g = dyn_cast_or_null<cir::GlobalOp>(
430 mlir::SymbolTable::lookupNearestSymbolFrom(
431 mlirModule, mlir::StringAttr::get(mlirModule->getContext(), name)));
433 g = cir::GlobalOp::create(builder, loc, name, type);
435 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
436 mlir::SymbolTable::setSymbolVisibility(
437 g, mlir::SymbolTable::Visibility::Private);
438 g.setGlobalVisibility(visibility);
443cir::FuncOp LoweringPreparePass::buildRuntimeFunction(
444 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
445 cir::FuncType type, cir::GlobalLinkageKind linkage) {
446 cir::FuncOp f = dyn_cast_or_null<FuncOp>(SymbolTable::lookupNearestSymbolFrom(
447 mlirModule, StringAttr::get(mlirModule->getContext(), name)));
449 f = cir::FuncOp::create(builder, loc, name, type);
451 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
452 mlir::SymbolTable::setSymbolVisibility(
453 f, mlir::SymbolTable::Visibility::Private);
463 builder.setInsertionPoint(op);
465 mlir::Value src = op.getSrc();
466 mlir::Value imag = builder.
getNullValue(src.getType(), op.getLoc());
472 cir::CastKind elemToBoolKind) {
474 builder.setInsertionPoint(op);
476 mlir::Value src = op.getSrc();
477 if (!mlir::isa<cir::BoolType>(op.getType()))
484 cir::BoolType boolTy = builder.
getBoolTy();
485 mlir::Value srcRealToBool =
486 builder.
createCast(op.getLoc(), elemToBoolKind, srcReal, boolTy);
487 mlir::Value srcImagToBool =
488 builder.
createCast(op.getLoc(), elemToBoolKind, srcImag, boolTy);
489 return builder.
createLogicalOr(op.getLoc(), srcRealToBool, srcImagToBool);
494 cir::CastKind scalarCastKind) {
496 builder.setInsertionPoint(op);
498 mlir::Value src = op.getSrc();
499 auto dstComplexElemTy =
500 mlir::cast<cir::ComplexType>(op.getType()).getElementType();
505 mlir::Value dstReal = builder.
createCast(op.getLoc(), scalarCastKind, srcReal,
507 mlir::Value dstImag = builder.
createCast(op.getLoc(), scalarCastKind, srcImag,
512void LoweringPreparePass::lowerCastOp(cir::CastOp op) {
513 mlir::MLIRContext &ctx = getContext();
514 mlir::Value loweredValue = [&]() -> mlir::Value {
515 switch (op.getKind()) {
516 case cir::CastKind::float_to_complex:
517 case cir::CastKind::int_to_complex:
519 case cir::CastKind::float_complex_to_real:
520 case cir::CastKind::int_complex_to_real:
522 case cir::CastKind::float_complex_to_bool:
524 case cir::CastKind::int_complex_to_bool:
526 case cir::CastKind::float_complex:
528 case cir::CastKind::float_complex_to_int_complex:
530 case cir::CastKind::int_complex:
532 case cir::CastKind::int_complex_to_float_complex:
540 op.replaceAllUsesWith(loweredValue);
547 llvm::StringRef (*libFuncNameGetter)(llvm::APFloat::Semantics),
548 mlir::Location loc, cir::ComplexType ty, mlir::Value lhsReal,
549 mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag) {
550 cir::FPTypeInterface elementTy =
551 mlir::cast<cir::FPTypeInterface>(ty.getElementType());
553 llvm::StringRef libFuncName = libFuncNameGetter(
554 llvm::APFloat::SemanticsToEnum(elementTy.getFloatSemantics()));
557 cir::FuncType libFuncTy = cir::FuncType::get(libFuncInputTypes, ty);
563 mlir::OpBuilder::InsertionGuard ipGuard{builder};
564 builder.setInsertionPointToStart(pass.mlirModule.getBody());
565 libFunc = pass.buildRuntimeFunction(builder, libFuncName, loc, libFuncTy);
569 builder.
createCallOp(loc, libFunc, {lhsReal, lhsImag, rhsReal, rhsImag});
570 return call.getResult();
573static llvm::StringRef
576 case llvm::APFloat::S_IEEEhalf:
578 case llvm::APFloat::S_IEEEsingle:
580 case llvm::APFloat::S_IEEEdouble:
582 case llvm::APFloat::S_PPCDoubleDouble:
584 case llvm::APFloat::S_x87DoubleExtended:
586 case llvm::APFloat::S_IEEEquad:
589 llvm_unreachable(
"unsupported floating point type");
595 mlir::Value lhsReal, mlir::Value lhsImag,
596 mlir::Value rhsReal, mlir::Value rhsImag) {
598 mlir::Value &a = lhsReal;
599 mlir::Value &
b = lhsImag;
600 mlir::Value &
c = rhsReal;
601 mlir::Value &d = rhsImag;
603 mlir::Value ac = builder.
createMul(loc, a,
c);
604 mlir::Value bd = builder.
createMul(loc,
b, d);
606 mlir::Value dd = builder.
createMul(loc, d, d);
607 mlir::Value acbd = builder.
createAdd(loc, ac, bd);
608 mlir::Value ccdd = builder.
createAdd(loc, cc, dd);
609 mlir::Value resultReal = builder.
createDiv(loc, acbd, ccdd);
612 mlir::Value ad = builder.
createMul(loc, a, d);
613 mlir::Value bcad = builder.
createSub(loc, bc, ad);
614 mlir::Value resultImag = builder.
createDiv(loc, bcad, ccdd);
620 mlir::Value lhsReal, mlir::Value lhsImag,
621 mlir::Value rhsReal, mlir::Value rhsImag) {
642 mlir::Value &a = lhsReal;
643 mlir::Value &
b = lhsImag;
644 mlir::Value &
c = rhsReal;
645 mlir::Value &d = rhsImag;
647 auto trueBranchBuilder = [&](mlir::OpBuilder &, mlir::Location) {
649 mlir::Value rd = builder.
createMul(loc, r, d);
650 mlir::Value tmp = builder.
createAdd(loc,
c, rd);
652 mlir::Value br = builder.
createMul(loc,
b, r);
653 mlir::Value abr = builder.
createAdd(loc, a, br);
654 mlir::Value e = builder.
createDiv(loc, abr, tmp);
656 mlir::Value ar = builder.
createMul(loc, a, r);
657 mlir::Value bar = builder.
createSub(loc,
b, ar);
658 mlir::Value f = builder.
createDiv(loc, bar, tmp);
664 auto falseBranchBuilder = [&](mlir::OpBuilder &, mlir::Location) {
666 mlir::Value rc = builder.
createMul(loc, r,
c);
667 mlir::Value tmp = builder.
createAdd(loc, d, rc);
669 mlir::Value ar = builder.
createMul(loc, a, r);
670 mlir::Value arb = builder.
createAdd(loc, ar,
b);
671 mlir::Value e = builder.
createDiv(loc, arb, tmp);
673 mlir::Value br = builder.
createMul(loc,
b, r);
674 mlir::Value bra = builder.
createSub(loc, br, a);
675 mlir::Value f = builder.
createDiv(loc, bra, tmp);
681 auto cFabs = cir::FAbsOp::create(builder, loc,
c);
682 auto dFabs = cir::FAbsOp::create(builder, loc, d);
683 cir::CmpOp cmpResult =
684 builder.
createCompare(loc, cir::CmpOpKind::ge, cFabs, dFabs);
685 auto ternary = cir::TernaryOp::create(builder, loc, cmpResult,
686 trueBranchBuilder, falseBranchBuilder);
688 return ternary.getResult();
695 auto getHigherPrecisionFPType = [&context](mlir::Type type) -> mlir::Type {
696 if (mlir::isa<cir::FP16Type>(type))
697 return cir::SingleType::get(&context);
699 if (mlir::isa<cir::SingleType>(type) || mlir::isa<cir::BF16Type>(type))
700 return cir::DoubleType::get(&context);
702 if (mlir::isa<cir::DoubleType>(type))
703 return cir::LongDoubleType::get(&context, type);
708 auto getFloatTypeSemantics =
709 [&cc](mlir::Type type) ->
const llvm::fltSemantics & {
711 if (mlir::isa<cir::FP16Type>(type))
714 if (mlir::isa<cir::BF16Type>(type))
717 if (mlir::isa<cir::SingleType>(type))
720 if (mlir::isa<cir::DoubleType>(type))
723 if (mlir::isa<cir::LongDoubleType>(type)) {
725 llvm_unreachable(
"NYI Float type semantics with OpenMP");
729 if (mlir::isa<cir::FP128Type>(type)) {
731 llvm_unreachable(
"NYI Float type semantics with OpenMP");
735 llvm_unreachable(
"Unsupported float type semantics");
738 const mlir::Type higherElementType = getHigherPrecisionFPType(elementType);
739 const llvm::fltSemantics &elementTypeSemantics =
740 getFloatTypeSemantics(elementType);
741 const llvm::fltSemantics &higherElementTypeSemantics =
742 getFloatTypeSemantics(higherElementType);
751 if (llvm::APFloat::semanticsMaxExponent(elementTypeSemantics) * 2 + 1 <=
752 llvm::APFloat::semanticsMaxExponent(higherElementTypeSemantics)) {
753 return higherElementType;
763 mlir::Location loc, cir::ComplexDivOp op, mlir::Value lhsReal,
764 mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag,
766 cir::ComplexType complexTy = op.getType();
767 if (mlir::isa<cir::FPTypeInterface>(complexTy.getElementType())) {
768 cir::ComplexRangeKind range = op.getRange();
769 if (range == cir::ComplexRangeKind::Improved)
773 if (range == cir::ComplexRangeKind::Full)
775 loc, complexTy, lhsReal, lhsImag, rhsReal,
778 if (range == cir::ComplexRangeKind::Promoted) {
779 mlir::Type originalElementType = complexTy.getElementType();
780 mlir::Type higherPrecisionElementType =
782 originalElementType);
784 if (!higherPrecisionElementType)
788 cir::CastKind floatingCastKind = cir::CastKind::floating;
789 lhsReal = builder.
createCast(floatingCastKind, lhsReal,
790 higherPrecisionElementType);
791 lhsImag = builder.
createCast(floatingCastKind, lhsImag,
792 higherPrecisionElementType);
793 rhsReal = builder.
createCast(floatingCastKind, rhsReal,
794 higherPrecisionElementType);
795 rhsImag = builder.
createCast(floatingCastKind, rhsImag,
796 higherPrecisionElementType);
799 builder, loc, lhsReal, lhsImag, rhsReal, rhsImag);
804 mlir::Value finalReal =
805 builder.
createCast(floatingCastKind, resultReal, originalElementType);
806 mlir::Value finalImag =
807 builder.
createCast(floatingCastKind, resultImag, originalElementType);
816void LoweringPreparePass::lowerComplexDivOp(cir::ComplexDivOp op) {
817 cir::CIRBaseBuilderTy builder(getContext());
818 builder.setInsertionPointAfter(op);
819 mlir::Location loc = op.getLoc();
820 mlir::TypedValue<cir::ComplexType> lhs = op.getLhs();
821 mlir::TypedValue<cir::ComplexType> rhs = op.getRhs();
827 mlir::Value loweredResult =
829 rhsImag, getContext(), *astCtx);
830 op.replaceAllUsesWith(loweredResult);
834static llvm::StringRef
837 case llvm::APFloat::S_IEEEhalf:
839 case llvm::APFloat::S_IEEEsingle:
841 case llvm::APFloat::S_IEEEdouble:
843 case llvm::APFloat::S_PPCDoubleDouble:
845 case llvm::APFloat::S_x87DoubleExtended:
847 case llvm::APFloat::S_IEEEquad:
850 llvm_unreachable(
"unsupported floating point type");
856 mlir::Location loc, cir::ComplexMulOp op,
857 mlir::Value lhsReal, mlir::Value lhsImag,
858 mlir::Value rhsReal, mlir::Value rhsImag) {
860 mlir::Value resultRealLhs = builder.
createMul(loc, lhsReal, rhsReal);
861 mlir::Value resultRealRhs = builder.
createMul(loc, lhsImag, rhsImag);
862 mlir::Value resultImagLhs = builder.
createMul(loc, lhsReal, rhsImag);
863 mlir::Value resultImagRhs = builder.
createMul(loc, lhsImag, rhsReal);
864 mlir::Value resultReal = builder.
createSub(loc, resultRealLhs, resultRealRhs);
865 mlir::Value resultImag = builder.
createAdd(loc, resultImagLhs, resultImagRhs);
866 mlir::Value algebraicResult =
869 cir::ComplexType complexTy = op.getType();
870 cir::ComplexRangeKind rangeKind = op.getRange();
871 if (mlir::isa<cir::IntType>(complexTy.getElementType()) ||
872 rangeKind == cir::ComplexRangeKind::Basic ||
873 rangeKind == cir::ComplexRangeKind::Improved ||
874 rangeKind == cir::ComplexRangeKind::Promoted)
875 return algebraicResult;
882 mlir::Value resultRealIsNaN = builder.
createIsNaN(loc, resultReal);
883 mlir::Value resultImagIsNaN = builder.
createIsNaN(loc, resultImag);
884 mlir::Value resultRealAndImagAreNaN =
887 return cir::TernaryOp::create(
888 builder, loc, resultRealAndImagAreNaN,
889 [&](mlir::OpBuilder &, mlir::Location) {
892 lhsReal, lhsImag, rhsReal, rhsImag);
895 [&](mlir::OpBuilder &, mlir::Location) {
901void LoweringPreparePass::lowerComplexMulOp(cir::ComplexMulOp op) {
902 cir::CIRBaseBuilderTy builder(getContext());
903 builder.setInsertionPointAfter(op);
904 mlir::Location loc = op.getLoc();
905 mlir::TypedValue<cir::ComplexType> lhs = op.getLhs();
906 mlir::TypedValue<cir::ComplexType> rhs = op.getRhs();
911 mlir::Value loweredResult =
lowerComplexMul(*
this, builder, loc, op, lhsReal,
912 lhsImag, rhsReal, rhsImag);
913 op.replaceAllUsesWith(loweredResult);
917void LoweringPreparePass::lowerUnaryOp(cir::UnaryOpInterface op) {
918 if (!mlir::isa<cir::ComplexType>(op.getResult().getType()))
921 mlir::Location loc = op->getLoc();
922 CIRBaseBuilderTy builder(getContext());
923 builder.setInsertionPointAfter(op);
925 mlir::Value operand = op.getInput();
929 mlir::Value resultReal = operandReal;
930 mlir::Value resultImag = operandImag;
932 llvm::TypeSwitch<mlir::Operation *>(op)
934 [&](
auto) { resultReal = builder.
createInc(loc, operandReal); })
936 [&](
auto) { resultReal = builder.
createDec(loc, operandReal); })
937 .Case<cir::MinusOp>([&](
auto) {
938 resultReal = builder.
createMinus(loc, operandReal);
939 resultImag = builder.
createMinus(loc, operandImag);
942 [&](
auto) { resultImag = builder.
createMinus(loc, operandImag); })
943 .
Default([](
auto) { llvm_unreachable(
"unhandled unary complex op"); });
946 op->replaceAllUsesWith(mlir::ValueRange{result});
950cir::FuncOp LoweringPreparePass::getOrCreateDtorFunc(CIRBaseBuilderTy &builder,
952 mlir::Region &dtorRegion,
953 cir::CallOp &dtorCall) {
954 mlir::OpBuilder::InsertionGuard guard(builder);
958 cir::VoidType voidTy = builder.
getVoidTy();
959 auto voidPtrTy = cir::PointerType::get(voidTy);
962 mlir::Block &dtorBlock = dtorRegion.front();
966 auto opIt = dtorBlock.getOperations().begin();
967 cir::GetGlobalOp ggop = mlir::cast<cir::GetGlobalOp>(*opIt);
978 if (dtorBlock.getOperations().size() == 3) {
979 auto callOp = mlir::dyn_cast<cir::CallOp>(&*(++opIt));
980 auto yieldOp = mlir::dyn_cast<cir::YieldOp>(&*(++opIt));
981 if (yieldOp && callOp && callOp.getNumOperands() == 1 &&
982 callOp.getArgOperand(0) == ggop) {
991 builder.setInsertionPointAfter(op);
992 SmallString<256> fnName(
"__cxx_global_array_dtor");
993 uint32_t cnt = dynamicInitializerNames[fnName]++;
995 fnName +=
"." + std::to_string(cnt);
998 auto fnType = cir::FuncType::get({voidPtrTy}, voidTy);
999 cir::FuncOp dtorFunc =
1000 buildRuntimeFunction(builder, fnName, op.getLoc(), fnType,
1001 cir::GlobalLinkageKind::InternalLinkage);
1003 SmallVector<mlir::NamedAttribute> paramAttrs;
1004 paramAttrs.push_back(
1005 builder.getNamedAttr(
"llvm.noundef", builder.getUnitAttr()));
1006 SmallVector<mlir::Attribute> argAttrDicts;
1007 argAttrDicts.push_back(
1008 mlir::DictionaryAttr::get(builder.getContext(), paramAttrs));
1009 dtorFunc.setArgAttrsAttr(
1010 mlir::ArrayAttr::get(builder.getContext(), argAttrDicts));
1012 mlir::Block *entryBB = dtorFunc.addEntryBlock();
1015 entryBB->getOperations().splice(entryBB->begin(), dtorBlock.getOperations(),
1016 dtorBlock.begin(), dtorBlock.end());
1019 cir::GetGlobalOp dtorGGop =
1020 mlir::cast<cir::GetGlobalOp>(entryBB->getOperations().front());
1021 builder.setInsertionPointToStart(&dtorBlock);
1022 builder.clone(*dtorGGop.getOperation());
1026 mlir::Value dtorArg = entryBB->getArgument(0);
1027 dtorGGop.replaceAllUsesWith(dtorArg);
1031 mlir::Block &finalBlock = dtorFunc.getBody().back();
1032 auto yieldOp = cast<cir::YieldOp>(finalBlock.getTerminator());
1033 builder.setInsertionPoint(yieldOp);
1034 cir::ReturnOp::create(builder, yieldOp->getLoc());
1039 cir::GetGlobalOp origGGop =
1040 mlir::cast<cir::GetGlobalOp>(dtorBlock.getOperations().front());
1041 builder.setInsertionPointAfter(origGGop);
1042 mlir::Value ggopResult = origGGop.getResult();
1043 dtorCall = builder.
createCallOp(op.getLoc(), dtorFunc, ggopResult);
1046 auto finalYield = cir::YieldOp::create(builder, op.getLoc());
1049 dtorBlock.getOperations().erase(std::next(mlir::Block::iterator(finalYield)),
1051 dtorRegion.getBlocks().erase(std::next(dtorRegion.begin()), dtorRegion.end());
1057LoweringPreparePass::buildCXXGlobalVarDeclInitFunc(cir::GlobalOp op) {
1060 SmallString<256> fnName(
"__cxx_global_var_init");
1062 uint32_t cnt = dynamicInitializerNames[fnName]++;
1064 fnName +=
"." + std::to_string(cnt);
1067 CIRBaseBuilderTy builder(getContext());
1068 builder.setInsertionPointAfter(op);
1069 cir::VoidType voidTy = builder.
getVoidTy();
1070 auto fnType = cir::FuncType::get({}, voidTy);
1071 FuncOp f = buildRuntimeFunction(builder, fnName, op.getLoc(), fnType,
1072 cir::GlobalLinkageKind::InternalLinkage);
1080 mlir::Block *entryBB = f.addEntryBlock();
1081 if (!op.getCtorRegion().empty()) {
1082 mlir::Block &block = op.getCtorRegion().front();
1083 entryBB->getOperations().splice(entryBB->begin(), block.getOperations(),
1084 block.begin(), std::prev(block.end()));
1088 mlir::Region &dtorRegion = op.getDtorRegion();
1089 if (!dtorRegion.empty()) {
1093 emitGlobalGuardedDtorRegion(builder, op, dtorRegion,
1094 op.getTlsModel().has_value(), *entryBB);
1098 builder.setInsertionPointToEnd(entryBB);
1099 mlir::Operation *yieldOp =
nullptr;
1100 if (!op.getCtorRegion().empty()) {
1101 mlir::Block &block = op.getCtorRegion().front();
1102 yieldOp = &block.getOperations().back();
1104 assert(!dtorRegion.empty());
1105 mlir::Block &block = dtorRegion.front();
1106 yieldOp = &block.getOperations().back();
1109 assert(isa<cir::YieldOp>(*yieldOp));
1110 cir::ReturnOp::create(builder, yieldOp->getLoc());
1115LoweringPreparePass::getGuardAcquireFn(cir::PointerType guardPtrTy) {
1117 CIRBaseBuilderTy builder(getContext());
1118 mlir::OpBuilder::InsertionGuard ipGuard{builder};
1119 builder.setInsertionPointToStart(mlirModule.getBody());
1120 mlir::Location loc = mlirModule.getLoc();
1121 cir::IntType intTy = cir::IntType::get(&getContext(), 32,
true);
1122 auto fnType = cir::FuncType::get({guardPtrTy}, intTy);
1123 return buildRuntimeFunction(builder,
"__cxa_guard_acquire", loc, fnType);
1127LoweringPreparePass::getGuardReleaseFn(cir::PointerType guardPtrTy) {
1129 CIRBaseBuilderTy builder(getContext());
1130 mlir::OpBuilder::InsertionGuard ipGuard{builder};
1131 builder.setInsertionPointToStart(mlirModule.getBody());
1132 mlir::Location loc = mlirModule.getLoc();
1133 cir::VoidType voidTy = cir::VoidType::get(&getContext());
1134 auto fnType = cir::FuncType::get({guardPtrTy}, voidTy);
1135 return buildRuntimeFunction(builder,
"__cxa_guard_release", loc, fnType);
1138cir::GlobalOp LoweringPreparePass::createGuardGlobalOp(
1139 CIRBaseBuilderTy &builder, mlir::Location loc, llvm::StringRef name,
1140 cir::IntType guardTy, cir::GlobalLinkageKind linkage) {
1141 mlir::OpBuilder::InsertionGuard guard(builder);
1142 builder.setInsertionPointToStart(mlirModule.getBody());
1143 cir::GlobalOp g = cir::GlobalOp::create(builder, loc, name, guardTy);
1145 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
1146 mlir::SymbolTable::setSymbolVisibility(
1147 g, mlir::SymbolTable::Visibility::Private);
1151void LoweringPreparePass::handleStaticLocal(cir::GlobalOp globalOp,
1152 cir::LocalInitOp localInitOp) {
1153 CIRBaseBuilderTy builder(getContext());
1155 std::optional<cir::ASTVarDeclInterface> astOption = globalOp.getAst();
1156 assert(astOption.has_value());
1157 cir::ASTVarDeclInterface
varDecl = astOption.value();
1159 builder.setInsertionPointAfter(localInitOp);
1160 mlir::Block *localInitBlock = builder.getInsertionBlock();
1163 mlir::Operation *ret = localInitBlock->getTerminator();
1167 builder.setInsertionPointAfter(localInitOp);
1171 bool nonTemplateInline =
1177 if (nonTemplateInline) {
1178 globalOp->emitError(
1179 "NYI: guarded initialization for inline namespace-scope variables");
1186 bool threadsafe = astCtx->
getLangOpts().ThreadsafeStatics &&
1187 (
varDecl.isLocalVarDecl() || nonTemplateInline) &&
1192 bool useInt8GuardVariable = !threadsafe && globalOp.hasInternalLinkage();
1193 cir::CIRDataLayout dataLayout(mlirModule);
1194 cir::IntType guardTy;
1195 clang::CharUnits guardAlignment;
1198 if (useInt8GuardVariable) {
1199 guardTy = cir::IntType::get(&getContext(), 8,
true);
1201 }
else if (useARMGuardVarABI()) {
1203 const unsigned sizeTypeSize =
1205 guardTy = cir::IntType::get(&getContext(), sizeTypeSize,
true);
1209 guardTy = cir::IntType::get(&getContext(), 64,
true);
1213 assert(guardTy && guardAlignment.
getQuantity() != 0);
1215 auto guardPtrTy = cir::PointerType::get(guardTy);
1218 cir::GlobalOp guard = getOrCreateStaticLocalDeclGuardAddress(
1219 builder, globalOp, varDecl, guardTy, guardAlignment);
1222 localInitBlock->push_back(ret);
1226 mlir::Value guardPtr = builder.
createGetGlobal(guard, localInitOp.getTls());
1248 unsigned maxInlineWidthInBits =
1251 if (!threadsafe || maxInlineWidthInBits) {
1253 auto bytePtrTy = cir::PointerType::get(builder.
getSIntNTy(8));
1254 mlir::Value bytePtr = builder.
createBitcast(guardPtr, bytePtrTy);
1256 localInitOp.getLoc(), bytePtr, guardAlignment.
getAsAlign().value());
1265 auto loadOp = mlir::cast<cir::LoadOp>(guardLoad.getDefiningOp());
1266 loadOp.setMemOrder(cir::MemOrder::Acquire);
1267 loadOp.setSyncScope(cir::SyncScopeKind::System);
1290 if (useARMGuardVarABI() && !useInt8GuardVariable) {
1292 localInitOp.getLoc(), mlir::cast<cir::IntType>(guardLoad.getType()),
1294 guardLoad = builder.
createAnd(localInitOp.getLoc(), guardLoad, one);
1299 localInitOp.getLoc(), mlir::cast<cir::IntType>(guardLoad.getType()), 0);
1300 auto needsInit = builder.
createCompare(localInitOp.getLoc(),
1301 cir::CmpOpKind::eq, guardLoad, zero);
1305 builder, globalOp.getLoc(), needsInit,
1306 false, [&](mlir::OpBuilder &, mlir::Location) {
1307 emitCXXGuardedInitIf(builder, globalOp, localInitOp.getCtorRegion(),
1308 localInitOp.getDtorRegion(), varDecl, guardPtr,
1309 guardPtrTy, threadsafe);
1314 globalOp->emitError(
"NYI: guarded init without inline atomics support");
1319 builder.getInsertionBlock()->push_back(ret);
1322void LoweringPreparePass::lowerLocalInitOp(
1323 cir::LocalInitOp initOp, mlir::SymbolTableCollection &symbolTables) {
1326 if (initOp.getCtorRegion().empty() && initOp.getDtorRegion().empty()) {
1331 cir::GlobalOp globalOp = initOp.getReferencedGlobal(symbolTables);
1332 assert(globalOp &&
"No global-op found");
1334 handleStaticLocal(globalOp, initOp);
1340void LoweringPreparePass::lowerGlobalOp(GlobalOp op) {
1342 if (op.getStaticLocalGuard())
1345 mlir::Region &ctorRegion = op.getCtorRegion();
1346 mlir::Region &dtorRegion = op.getDtorRegion();
1348 if (!ctorRegion.empty() || !dtorRegion.empty()) {
1351 cir::FuncOp f = buildCXXGlobalVarDeclInitFunc(op);
1354 ctorRegion.getBlocks().clear();
1355 dtorRegion.getBlocks().clear();
1358 dynamicInitializers.push_back(f);
1364void LoweringPreparePass::lowerThreeWayCmpOp(CmpThreeWayOp op) {
1365 CIRBaseBuilderTy builder(getContext());
1366 builder.setInsertionPointAfter(op);
1368 mlir::Location loc = op->getLoc();
1369 cir::CmpThreeWayInfoAttr cmpInfo = op.getInfo();
1378 mlir::Value transformedResult;
1379 if (cmpInfo.getOrdering() != CmpOrdering::Partial) {
1382 builder.
createCompare(loc, CmpOpKind::lt, op.getLhs(), op.getRhs());
1383 mlir::Value selectOnLt = builder.
createSelect(loc, lt, ltRes, gtRes);
1385 builder.
createCompare(loc, CmpOpKind::eq, op.getLhs(), op.getRhs());
1386 transformedResult = builder.
createSelect(loc, eq, eqRes, selectOnLt);
1390 loc, op.getType(), cmpInfo.getUnordered().value());
1393 builder.
createCompare(loc, CmpOpKind::eq, op.getLhs(), op.getRhs());
1394 mlir::Value selectOnEq = builder.
createSelect(loc, eq, eqRes, unorderedRes);
1396 builder.
createCompare(loc, CmpOpKind::gt, op.getLhs(), op.getRhs());
1397 mlir::Value selectOnGt = builder.
createSelect(loc, gt, gtRes, selectOnEq);
1399 builder.
createCompare(loc, CmpOpKind::lt, op.getLhs(), op.getRhs());
1400 transformedResult = builder.
createSelect(loc, lt, ltRes, selectOnGt);
1403 op.replaceAllUsesWith(transformedResult);
1407template <
typename AttributeTy>
1408static llvm::SmallVector<mlir::Attribute>
1412 for (
const auto &[name, priority] : list)
1413 attrs.push_back(AttributeTy::get(context, name, priority));
1417void LoweringPreparePass::buildGlobalCtorDtorList() {
1418 if (!globalCtorList.empty()) {
1419 llvm::SmallVector<mlir::Attribute> globalCtors =
1423 mlirModule->setAttr(cir::CIRDialect::getGlobalCtorsAttrName(),
1424 mlir::ArrayAttr::get(&getContext(), globalCtors));
1427 if (!globalDtorList.empty()) {
1428 llvm::SmallVector<mlir::Attribute> globalDtors =
1431 mlirModule->setAttr(cir::CIRDialect::getGlobalDtorsAttrName(),
1432 mlir::ArrayAttr::get(&getContext(), globalDtors));
1436void LoweringPreparePass::buildCXXGlobalInitFunc() {
1437 if (dynamicInitializers.empty())
1444 SmallString<256> fnName;
1452 llvm::raw_svector_ostream
out(fnName);
1453 std::unique_ptr<clang::MangleContext> mangleCtx(
1455 cast<clang::ItaniumMangleContext>(*mangleCtx)
1458 fnName +=
"_GLOBAL__sub_I_";
1462 CIRBaseBuilderTy builder(getContext());
1463 builder.setInsertionPointToEnd(&mlirModule.getBodyRegion().back());
1464 auto fnType = cir::FuncType::get({}, builder.
getVoidTy());
1466 buildRuntimeFunction(builder, fnName, mlirModule.getLoc(), fnType,
1467 cir::GlobalLinkageKind::ExternalLinkage);
1468 builder.setInsertionPointToStart(f.addEntryBlock());
1469 for (cir::FuncOp &f : dynamicInitializers)
1473 globalCtorList.emplace_back(fnName,
1474 cir::GlobalCtorAttr::getDefaultPriority());
1476 cir::ReturnOp::create(builder, f.getLoc());
1485 mlir::Operation *op, mlir::Type eltTy,
1487 mlir::Value numElements,
1488 uint64_t arrayLen,
bool isCtor) {
1489 mlir::Location loc = op->getLoc();
1490 bool isDynamic = numElements !=
nullptr;
1494 const unsigned sizeTypeSize =
1500 mlir::Value begin, end;
1503 end = cir::PtrStrideOp::create(builder, loc, eltTy, begin, numElements);
1505 mlir::Value endOffsetVal =
1507 begin = cir::CastOp::create(builder, loc, eltTy,
1508 cir::CastKind::array_to_ptrdecay, addr);
1509 end = cir::PtrStrideOp::create(builder, loc, eltTy, begin, endOffsetVal);
1512 mlir::Value start = isCtor ? begin : end;
1513 mlir::Value stop = isCtor ? end : begin;
1519 mlir::Value guardCond;
1522 guardCond = cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne,
1528 cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne, start, stop);
1530 ifOp = cir::IfOp::create(builder, loc, guardCond,
1532 [&](mlir::OpBuilder &, mlir::Location) {});
1533 builder.setInsertionPointToStart(&ifOp.getThenRegion().front());
1541 mlir::Block *bodyBlock = &op->getRegion(0).front();
1546 auto cloneRegionBodyInto = [&](mlir::Block *srcBlock,
1547 mlir::Value replacement) {
1548 mlir::IRMapping map;
1549 map.map(srcBlock->getArgument(0), replacement);
1550 for (mlir::Operation ®ionOp : *srcBlock) {
1551 if (!mlir::isa<cir::YieldOp>(®ionOp))
1552 builder.clone(regionOp, map);
1556 mlir::Block *partialDtorBlock =
nullptr;
1557 if (
auto arrayCtor = mlir::dyn_cast<cir::ArrayCtor>(op)) {
1558 mlir::Region &partialDtor = arrayCtor.getPartialDtor();
1559 if (!partialDtor.empty())
1560 partialDtorBlock = &partialDtor.front();
1561 }
else if (
auto arrayDtor = mlir::dyn_cast<cir::ArrayDtor>(op)) {
1570 if (arrayDtor.getDtorMayThrow())
1571 partialDtorBlock = bodyBlock;
1574 auto emitCtorDtorLoop = [&]() {
1578 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1579 auto currentElement = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1580 auto cmp = cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne,
1581 currentElement, stop);
1585 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1586 auto currentElement = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1588 cloneRegionBodyInto(bodyBlock, currentElement);
1589 mlir::Value stride = builder.
getUnsignedInt(loc, 1, sizeTypeSize);
1590 auto nextElement = cir::PtrStrideOp::create(builder, loc, eltTy,
1591 currentElement, stride);
1594 mlir::Value stride = builder.
getSignedInt(loc, -1, sizeTypeSize);
1595 auto prevElement = cir::PtrStrideOp::create(builder, loc, eltTy,
1596 currentElement, stride);
1598 cloneRegionBodyInto(bodyBlock, prevElement);
1601 cir::YieldOp::create(
b, loc);
1605 if (partialDtorBlock) {
1606 cir::CleanupScopeOp::create(
1607 builder, loc, cir::CleanupKind::EH,
1609 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1611 cir::YieldOp::create(
b, loc);
1614 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1615 auto cur = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1617 cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne, cur, begin);
1619 builder, loc, cmp,
false,
1620 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1624 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1625 auto el = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1626 auto neq = cir::CmpOp::create(
1627 builder, loc, cir::CmpOpKind::ne, el, begin);
1631 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1632 auto el = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1633 mlir::Value negOne =
1635 auto prev = cir::PtrStrideOp::create(builder, loc, eltTy,
1638 cloneRegionBodyInto(partialDtorBlock, prev);
1641 cir::YieldOp::create(builder, loc);
1643 cir::YieldOp::create(
b, loc);
1650 cir::YieldOp::create(builder, loc);
1655void LoweringPreparePass::lowerArrayDtor(cir::ArrayDtor op) {
1656 CIRBaseBuilderTy builder(getContext());
1657 builder.setInsertionPointAfter(op.getOperation());
1659 mlir::Type eltTy = op->getRegion(0).getArgument(0).getType();
1661 if (op.getNumElements()) {
1663 op.getNumElements(), 0,
1669 mlir::cast<cir::ArrayType>(op.getAddr().getType().getPointee()).getSize();
1675void LoweringPreparePass::lowerArrayCtor(cir::ArrayCtor op) {
1676 cir::CIRBaseBuilderTy builder(getContext());
1677 builder.setInsertionPointAfter(op.getOperation());
1679 mlir::Type eltTy = op->getRegion(0).getArgument(0).getType();
1681 if (op.getNumElements()) {
1683 op.getNumElements(), 0,
1689 mlir::cast<cir::ArrayType>(op.getAddr().getType().getPointee()).getSize();
1695void LoweringPreparePass::lowerTrivialCopyCall(cir::CallOp op) {
1700 std::optional<cir::CtorKind> ctorKind = funcOp.getCxxConstructorKind();
1701 if (ctorKind && *ctorKind == cir::CtorKind::Copy &&
1702 funcOp.isCxxTrivialMemberFunction()) {
1704 CIRBaseBuilderTy builder(getContext());
1705 mlir::ValueRange operands = op.getOperands();
1706 mlir::Value dest = operands[0];
1707 mlir::Value src = operands[1];
1708 builder.setInsertionPoint(op);
1714cir::GlobalOp LoweringPreparePass::getOrCreateConstAggregateGlobal(
1715 CIRBaseBuilderTy &builder, mlir::SymbolTableCollection &symbolTables,
1716 mlir::Location loc, llvm::StringRef baseName, mlir::Type ty,
1717 mlir::TypedAttr constant) {
1719 llvm::SmallVector<cir::GlobalOp, 1> &versions =
1720 constAggregateGlobals[baseName];
1723 for (cir::GlobalOp gv : versions) {
1724 if (gv.getSymType() == ty && gv.getInitialValue() == constant)
1732 llvm::SmallString<128>
name(baseName);
1733 size_t baseLen =
name.size();
1734 unsigned version = versions.size();
1736 name.resize(baseLen);
1738 name.push_back(
'.');
1739 llvm::Twine(version).toVector(name);
1741 auto existingGv = symbolTables.lookupSymbolIn<cir::GlobalOp>(
1742 mlirModule, mlir::StringAttr::get(&getContext(), name));
1745 versions.push_back(existingGv);
1746 if (existingGv.getSymType() == ty &&
1747 existingGv.getInitialValue() == constant)
1753 mlir::OpBuilder::InsertionGuard guard(builder);
1754 builder.setInsertionPointToStart(mlirModule.getBody());
1756 cir::GlobalOp::create(builder, loc, name, ty,
1758 cir::LangAddressSpaceAttr::get(
1759 &getContext(), cir::LangAddressSpace::Default),
1760 cir::GlobalLinkageKind::PrivateLinkage);
1761 mlir::SymbolTable::setSymbolVisibility(
1762 gv, mlir::SymbolTable::Visibility::Private);
1763 gv.setInitialValueAttr(constant);
1767 symbolTables.getSymbolTable(mlirModule).insert(gv);
1769 versions.push_back(gv);
1773void LoweringPreparePass::lowerStoreOfConstAggregate(
1774 cir::StoreOp op, mlir::SymbolTableCollection &symbolTables) {
1776 auto constOp = op.getValue().getDefiningOp<cir::ConstantOp>();
1780 mlir::Type ty = constOp.getType();
1781 if (!mlir::isa<cir::ArrayType, cir::RecordType>(ty))
1787 auto alloca = op.getAddr().getDefiningOp<cir::AllocaOp>();
1791 mlir::TypedAttr constant = constOp.getValue();
1802 auto func = op->getParentOfType<cir::FuncOp>();
1805 llvm::StringRef funcName = func.getSymName();
1808 llvm::StringRef varName = alloca.getName();
1811 std::string baseName = (
"__const." + funcName +
"." + varName).str();
1812 CIRBaseBuilderTy builder(getContext());
1816 cir::GlobalOp gv = getOrCreateConstAggregateGlobal(
1817 builder, symbolTables, op.getLoc(), baseName, ty, constant);
1820 builder.setInsertionPoint(op);
1822 auto ptrTy = cir::PointerType::get(ty);
1823 mlir::Value globalPtr =
1824 cir::GetGlobalOp::create(builder, op.getLoc(), ptrTy, gv.getSymName());
1833 if (constOp.use_empty())
1837void LoweringPreparePass::runOnOp(mlir::Operation *op,
1838 mlir::SymbolTableCollection &symbolTables) {
1839 if (
auto arrayCtor = dyn_cast<cir::ArrayCtor>(op)) {
1840 lowerArrayCtor(arrayCtor);
1841 }
else if (
auto arrayDtor = dyn_cast<cir::ArrayDtor>(op)) {
1842 lowerArrayDtor(arrayDtor);
1843 }
else if (
auto cast = mlir::dyn_cast<cir::CastOp>(op)) {
1845 }
else if (
auto complexDiv = mlir::dyn_cast<cir::ComplexDivOp>(op)) {
1846 lowerComplexDivOp(complexDiv);
1847 }
else if (
auto complexMul = mlir::dyn_cast<cir::ComplexMulOp>(op)) {
1848 lowerComplexMulOp(complexMul);
1849 }
else if (
auto glob = mlir::dyn_cast<cir::GlobalOp>(op)) {
1850 lowerGlobalOp(glob);
1851 }
else if (
auto unaryOp = mlir::dyn_cast<cir::UnaryOpInterface>(op)) {
1852 lowerUnaryOp(unaryOp);
1853 }
else if (
auto callOp = dyn_cast<cir::CallOp>(op)) {
1854 lowerTrivialCopyCall(callOp);
1855 }
else if (
auto storeOp = dyn_cast<cir::StoreOp>(op)) {
1856 lowerStoreOfConstAggregate(storeOp, symbolTables);
1857 }
else if (
auto fnOp = dyn_cast<cir::FuncOp>(op)) {
1858 if (
auto globalCtor = fnOp.getGlobalCtorPriority())
1859 globalCtorList.emplace_back(fnOp.getName(), globalCtor.value());
1860 else if (
auto globalDtor = fnOp.getGlobalDtorPriority())
1861 globalDtorList.emplace_back(fnOp.getName(), globalDtor.value());
1863 if (mlir::Attribute attr =
1864 fnOp->getAttr(cir::CUDAKernelNameAttr::getMnemonic())) {
1865 auto kernelNameAttr = dyn_cast<CUDAKernelNameAttr>(attr);
1866 llvm::StringRef kernelName = kernelNameAttr.getKernelName();
1867 cudaKernelMap[kernelName] = fnOp;
1869 }
else if (
auto threeWayCmp = dyn_cast<cir::CmpThreeWayOp>(op)) {
1870 lowerThreeWayCmpOp(threeWayCmp);
1871 }
else if (
auto initOp = dyn_cast<cir::LocalInitOp>(op)) {
1872 lowerLocalInitOp(initOp, symbolTables);
1883 llvm::StringRef name) {
1884 return (
"__" + prefix + name).str();
1906void LoweringPreparePass::buildCUDAModuleCtor() {
1911 if (astCtx->
getLangOpts().GPURelocatableDeviceCode)
1912 llvm_unreachable(
"GPU RDC NYI");
1916 if (cudaKernelMap.empty())
1921 mlir::Attribute cudaBinaryHandleAttr =
1922 mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName());
1923 if (!cudaBinaryHandleAttr) {
1929 llvm::StringRef cudaGPUBinaryName =
1930 mlir::cast<CUDABinaryHandleAttr>(cudaBinaryHandleAttr)
1934 llvm::vfs::FileSystem &vfs =
1936 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> gpuBinaryOrErr =
1937 vfs.getBufferForFile(cudaGPUBinaryName);
1938 if (std::error_code ec = gpuBinaryOrErr.getError()) {
1939 mlirModule->emitError(
"cannot open GPU binary file: " + cudaGPUBinaryName +
1940 ": " + ec.message());
1943 std::unique_ptr<llvm::MemoryBuffer> gpuBinary =
1944 std::move(gpuBinaryOrErr.get());
1948 mlir::Location loc = mlirModule->getLoc();
1949 CIRBaseBuilderTy builder(getContext());
1950 builder.setInsertionPointToStart(mlirModule.getBody());
1954 PointerType voidPtrPtrTy = builder.
getPointerTo(voidPtrTy);
1956 IntType charTy = cir::IntType::get(&getContext(), astCtx->
getCharWidth(),
1962 llvm::StringRef fatbinConstName =
1963 astCtx->
getLangOpts().HIP ?
".hip_fatbin" :
".nv_fatbin";
1965 llvm::StringRef fatbinSectionName =
1966 astCtx->
getLangOpts().HIP ?
".hipFatBinSegment" :
".nvFatBinSegment";
1970 ArrayType::get(&getContext(), charTy, gpuBinary->getBuffer().size());
1972 GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType,
1974 GlobalLinkageKind::PrivateLinkage);
1975 fatbinStr.setAlignment(8);
1976 fatbinStr.setInitialValueAttr(cir::ConstArrayAttr::get(
1977 fatbinType, StringAttr::get(gpuBinary->getBuffer(), fatbinType)));
1978 fatbinStr.setSection(fatbinConstName);
1979 fatbinStr.setPrivate();
1983 auto fatbinWrapperType = RecordType::get(
1984 &getContext(), {intTy, intTy, voidPtrTy, voidPtrTy},
1985 false,
false, RecordType::RecordKind::Struct);
1986 std::string fatbinWrapperName =
1988 GlobalOp fatbinWrapper = GlobalOp::create(
1989 builder, loc, fatbinWrapperName, fatbinWrapperType,
1990 true, {}, GlobalLinkageKind::PrivateLinkage);
1991 fatbinWrapper.setSection(fatbinSectionName);
1993 constexpr unsigned cudaFatMagic = 0x466243b1;
1994 constexpr unsigned hipFatMagic = 0x48495046;
1995 unsigned fatMagic =
isHIP ? hipFatMagic : cudaFatMagic;
1997 auto magicInit = IntAttr::get(intTy, fatMagic);
1998 auto versionInit = IntAttr::get(intTy, 1);
1999 auto fatbinStrSymbol =
2000 mlir::FlatSymbolRefAttr::get(fatbinStr.getSymNameAttr());
2001 auto fatbinInit = GlobalViewAttr::get(voidPtrTy, fatbinStrSymbol);
2003 fatbinWrapper.setInitialValueAttr(cir::ConstRecordAttr::get(
2005 mlir::ArrayAttr::get(&getContext(),
2006 {magicInit, versionInit, fatbinInit, unusedInit})));
2009 std::string gpubinHandleName =
2012 GlobalOp gpuBinHandle = GlobalOp::create(
2013 builder, loc, gpubinHandleName, voidPtrPtrTy,
2014 false, {}, cir::GlobalLinkageKind::InternalLinkage);
2016 gpuBinHandle.setPrivate();
2021 std::string regFuncName =
2023 FuncType regFuncType = FuncType::get({voidPtrTy}, voidPtrPtrTy);
2024 cir::FuncOp regFunc =
2025 buildRuntimeFunction(builder, regFuncName, loc, regFuncType);
2028 cir::FuncOp moduleCtor = buildRuntimeFunction(
2029 builder, moduleCtorName, loc, FuncType::get({}, voidTy),
2030 GlobalLinkageKind::InternalLinkage);
2032 globalCtorList.emplace_back(moduleCtorName,
2033 cir::GlobalCtorAttr::getDefaultPriority());
2034 builder.setInsertionPointToStart(moduleCtor.addEntryBlock());
2037 llvm_unreachable(
"HIP Module Constructor Support");
2038 }
else if (!astCtx->
getLangOpts().GPURelocatableDeviceCode) {
2046 mlir::Value fatbinVoidPtr = builder.
createBitcast(wrapper, voidPtrTy);
2047 cir::CallOp gpuBinaryHandleCall =
2049 mlir::Value gpuBinaryHandle = gpuBinaryHandleCall.getResult();
2051 mlir::Value gpuBinaryHandleGlobal = builder.
createGetGlobal(gpuBinHandle);
2052 builder.
createStore(loc, gpuBinaryHandle, gpuBinaryHandleGlobal);
2055 if (std::optional<FuncOp> regGlobal = buildCUDARegisterGlobals()) {
2056 builder.
createCallOp(loc, *regGlobal, gpuBinaryHandle);
2065 cir::CIRBaseBuilderTy globalBuilder(getContext());
2066 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
2068 buildRuntimeFunction(globalBuilder,
"__cudaRegisterFatBinaryEnd", loc,
2069 FuncType::get({voidPtrPtrTy}, voidTy));
2073 llvm_unreachable(
"GPU RDC NYI");
2078 if (std::optional<FuncOp> dtor = buildCUDAModuleDtor()) {
2081 cir::CIRBaseBuilderTy globalBuilder(getContext());
2082 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
2083 FuncOp atexit = buildRuntimeFunction(
2084 globalBuilder,
"atexit", loc,
2085 FuncType::get(PointerType::get(dtor->getFunctionType()), intTy));
2086 mlir::Value dtorFunc = GetGlobalOp::create(
2087 builder, loc, PointerType::get(dtor->getFunctionType()),
2088 mlir::FlatSymbolRefAttr::get(dtor->getSymNameAttr()));
2091 cir::ReturnOp::create(builder, loc);
2094std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() {
2095 if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName()))
2100 VoidType voidTy = VoidType::get(&getContext());
2101 PointerType voidPtrPtrTy = PointerType::get(PointerType::get(voidTy));
2103 mlir::Location loc = mlirModule.getLoc();
2105 cir::CIRBaseBuilderTy builder(getContext());
2106 builder.setInsertionPointToStart(mlirModule.getBody());
2109 std::string unregisterFuncName =
2111 FuncOp unregisterFunc = buildRuntimeFunction(
2112 builder, unregisterFuncName, loc, FuncType::get({voidPtrPtrTy}, voidTy));
2121 buildRuntimeFunction(builder, dtorName, loc, FuncType::get({}, voidTy),
2122 GlobalLinkageKind::InternalLinkage);
2124 builder.setInsertionPointToStart(dtor.addEntryBlock());
2130 GlobalOp gpubinGlobal = cast<GlobalOp>(mlirModule.lookupSymbol(gpubinName));
2132 mlir::Value gpubin = builder.
createLoad(loc, gpubinAddress);
2134 ReturnOp::create(builder, loc);
2139std::optional<FuncOp> LoweringPreparePass::buildCUDARegisterGlobals() {
2141 if (cudaKernelMap.empty())
2144 cir::CIRBaseBuilderTy builder(getContext());
2145 builder.setInsertionPointToStart(mlirModule.getBody());
2147 mlir::Location loc = mlirModule.getLoc();
2150 auto voidTy = VoidType::get(&getContext());
2151 auto voidPtrTy = PointerType::get(voidTy);
2152 auto voidPtrPtrTy = PointerType::get(voidPtrTy);
2156 std::string regGlobalFuncName =
2158 auto regGlobalFuncTy = FuncType::get({voidPtrPtrTy}, voidTy);
2159 FuncOp regGlobalFunc =
2160 buildRuntimeFunction(builder, regGlobalFuncName, loc, regGlobalFuncTy,
2161 GlobalLinkageKind::InternalLinkage);
2162 builder.setInsertionPointToStart(regGlobalFunc.addEntryBlock());
2164 buildCUDARegisterGlobalFunctions(builder, regGlobalFunc);
2168 ReturnOp::create(builder, loc);
2169 return regGlobalFunc;
2172void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
2173 cir::CIRBaseBuilderTy &builder, FuncOp regGlobalFunc) {
2174 mlir::Location loc = mlirModule.getLoc();
2176 cir::CIRDataLayout dataLayout(mlirModule);
2178 auto voidTy = VoidType::get(&getContext());
2179 auto voidPtrTy = PointerType::get(voidTy);
2180 auto voidPtrPtrTy = PointerType::get(voidPtrTy);
2182 IntType charTy = cir::IntType::get(&getContext(), astCtx->
getCharWidth(),
2186 mlir::Value fatbinHandle = *regGlobalFunc.args_begin();
2188 cir::CIRBaseBuilderTy globalBuilder(getContext());
2189 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
2203 FuncOp cudaRegisterFunction = buildRuntimeFunction(
2205 FuncType::get({voidPtrPtrTy, voidPtrTy, voidPtrTy, voidPtrTy, intTy,
2206 voidPtrTy, voidPtrTy, voidPtrTy, voidPtrTy, voidPtrTy},
2209 auto makeConstantString = [&](llvm::StringRef str) -> GlobalOp {
2210 auto strType = ArrayType::get(&getContext(), charTy, 1 + str.size());
2211 auto tmpString = cir::GlobalOp::create(
2212 globalBuilder, loc, (
".str" + str).str(), strType,
2214 cir::GlobalLinkageKind::PrivateLinkage);
2217 tmpString.setInitialValueAttr(
2218 ConstArrayAttr::get(strType, StringAttr::get(str +
"\0", strType)));
2219 tmpString.setPrivate();
2223 cir::ConstantOp cirNullPtr = builder.
getNullPtr(voidPtrTy, loc);
2225 for (
auto kernelName : cudaKernelMap.keys()) {
2226 FuncOp deviceStub = cudaKernelMap[kernelName];
2227 GlobalOp deviceFuncStr = makeConstantString(kernelName);
2232 llvm_unreachable(
"HIP kernel registration NYI");
2235 GetGlobalOp::create(
2236 builder, loc, PointerType::get(deviceStub.getFunctionType()),
2237 mlir::FlatSymbolRefAttr::get(deviceStub.getSymNameAttr())),
2240 loc, cudaRegisterFunction,
2241 {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
2242 ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)),
2243 cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
2248void LoweringPreparePass::runOnOperation() {
2249 mlir::Operation *op = getOperation();
2250 if (isa<::mlir::ModuleOp>(op))
2251 mlirModule = cast<::mlir::ModuleOp>(op);
2253 llvm::SmallVector<mlir::Operation *> opsToTransform;
2254 mlir::SymbolTableCollection symbolTables;
2256 op->walk([&](mlir::Operation *op) {
2257 if (mlir::isa<cir::ArrayCtor, cir::ArrayDtor, cir::CastOp,
2258 cir::ComplexMulOp, cir::ComplexDivOp, cir::DynamicCastOp,
2259 cir::FuncOp, cir::CallOp, cir::GlobalOp, cir::StoreOp,
2260 cir::CmpThreeWayOp, cir::IncOp, cir::DecOp, cir::MinusOp,
2261 cir::NotOp, cir::LocalInitOp>(op))
2262 opsToTransform.push_back(op);
2265 for (mlir::Operation *o : opsToTransform)
2266 runOnOp(o, symbolTables);
2268 buildCXXGlobalInitFunc();
2270 buildCUDAModuleCtor();
2272 buildGlobalCtorDtorList();
2276 return std::make_unique<LoweringPreparePass>();
2279std::unique_ptr<Pass>
2281 auto pass = std::make_unique<LoweringPreparePass>();
2282 pass->setASTContext(astCtx);
2283 return std::move(pass);
Defines the clang::ASTContext interface.
static void emitBody(CodeGenFunction &CGF, const Stmt *S, const Stmt *NextLoop, int MaxLevel, int Level=0)
static llvm::FunctionCallee getGuardReleaseFn(CodeGenModule &CGM, llvm::PointerType *GuardPtrTy)
static llvm::FunctionCallee getGuardAcquireFn(CodeGenModule &CGM, llvm::PointerType *GuardPtrTy)
static mlir::Value buildRangeReductionComplexDiv(CIRBaseBuilderTy &builder, mlir::Location loc, mlir::Value lhsReal, mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag)
static llvm::StringRef getComplexDivLibCallName(llvm::APFloat::Semantics semantics)
static llvm::SmallVector< mlir::Attribute > prepareCtorDtorAttrList(mlir::MLIRContext *context, llvm::ArrayRef< std::pair< std::string, uint32_t > > list)
static llvm::StringRef getComplexMulLibCallName(llvm::APFloat::Semantics semantics)
static mlir::Value buildComplexBinOpLibCall(LoweringPreparePass &pass, CIRBaseBuilderTy &builder, llvm::StringRef(*libFuncNameGetter)(llvm::APFloat::Semantics), mlir::Location loc, cir::ComplexType ty, mlir::Value lhsReal, mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag)
static mlir::Value lowerComplexMul(LoweringPreparePass &pass, CIRBaseBuilderTy &builder, mlir::Location loc, cir::ComplexMulOp op, mlir::Value lhsReal, mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag)
static std::string addUnderscoredPrefix(llvm::StringRef prefix, llvm::StringRef name)
static SmallString< 128 > getTransformedFileName(mlir::ModuleOp mlirModule)
static mlir::Value lowerComplexToComplexCast(mlir::MLIRContext &ctx, cir::CastOp op, cir::CastKind scalarCastKind)
static void lowerArrayDtorCtorIntoLoop(cir::CIRBaseBuilderTy &builder, clang::ASTContext *astCtx, mlir::Operation *op, mlir::Type eltTy, mlir::Value addr, mlir::Value numElements, uint64_t arrayLen, bool isCtor)
Lower a cir.array.ctor or cir.array.dtor into a do-while loop that iterates over every element.
static mlir::Value lowerComplexToScalarCast(mlir::MLIRContext &ctx, cir::CastOp op, cir::CastKind elemToBoolKind)
static mlir::Value buildAlgebraicComplexDiv(CIRBaseBuilderTy &builder, mlir::Location loc, mlir::Value lhsReal, mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag)
static llvm::StringRef getCUDAPrefix(clang::ASTContext *astCtx)
static cir::FuncOp getCalledFunction(cir::CallOp callOp)
Return the FuncOp called by callOp.
static mlir::Type higherPrecisionElementTypeForComplexArithmetic(mlir::MLIRContext &context, clang::ASTContext &cc, CIRBaseBuilderTy &builder, mlir::Type elementType)
static mlir::Value lowerScalarToComplexCast(mlir::MLIRContext &ctx, cir::CastOp op)
static mlir::Value lowerComplexDiv(LoweringPreparePass &pass, CIRBaseBuilderTy &builder, mlir::Location loc, cir::ComplexDivOp op, mlir::Value lhsReal, mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag, mlir::MLIRContext &mlirCx, clang::ASTContext &cc)
Defines the clang::Module class, which describes a module in the source code.
Defines the SourceManager interface.
Defines various enumerations that describe declaration and type specifiers.
Defines the TargetCXXABI class, which abstracts details of the C++ ABI that we're targeting.
__device__ __2f16 float c
mlir::Value createDiv(mlir::Location loc, mlir::Value lhs, mlir::Value rhs)
mlir::TypedAttr getConstNullPtrAttr(mlir::Type t)
mlir::Value createDec(mlir::Location loc, mlir::Value input, bool nsw=false)
mlir::Value createLogicalOr(mlir::Location loc, mlir::Value lhs, mlir::Value rhs)
mlir::Value createSub(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, OverflowBehavior ob=OverflowBehavior::None)
cir::ConditionOp createCondition(mlir::Value condition)
Create a loop condition.
mlir::Value createInc(mlir::Location loc, mlir::Value input, bool nsw=false)
cir::CopyOp createCopy(mlir::Value dst, mlir::Value src, bool isVolatile=false, bool skipTailPadding=false)
Create a copy with inferred length.
cir::VoidType getVoidTy()
cir::ConstantOp getNullValue(mlir::Type ty, mlir::Location loc)
mlir::Value createCast(mlir::Location loc, cir::CastKind kind, mlir::Value src, mlir::Type newTy)
cir::PointerType getVoidFnPtrTy(mlir::TypeRange argTypes={})
Returns void (*)(T...) as a cir::PointerType.
mlir::Value createAdd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, OverflowBehavior ob=OverflowBehavior::None)
cir::PointerType getPointerTo(mlir::Type ty)
mlir::Value createComplexImag(mlir::Location loc, mlir::Value operand)
cir::ConstantOp getNullPtr(mlir::Type ty, mlir::Location loc)
cir::DoWhileOp createDoWhile(mlir::Location loc, llvm::function_ref< void(mlir::OpBuilder &, mlir::Location)> condBuilder, llvm::function_ref< void(mlir::OpBuilder &, mlir::Location)> bodyBuilder)
Create a do-while operation.
cir::GetGlobalOp createGetGlobal(mlir::Location loc, cir::GlobalOp global, bool threadLocal=false)
mlir::Value getSignedInt(mlir::Location loc, int64_t val, unsigned numBits)
mlir::Value createAnd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs)
mlir::Value createBitcast(mlir::Value src, mlir::Type newTy)
cir::FuncType getVoidFnTy(mlir::TypeRange argTypes={})
Returns void (T...) as a cir::FuncType.
cir::CmpOp createCompare(mlir::Location loc, cir::CmpOpKind kind, mlir::Value lhs, mlir::Value rhs)
mlir::IntegerAttr getAlignmentAttr(clang::CharUnits alignment)
mlir::Value createSelect(mlir::Location loc, mlir::Value condition, mlir::Value trueValue, mlir::Value falseValue)
mlir::Value createMul(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, OverflowBehavior ob=OverflowBehavior::None)
cir::LoadOp createLoad(mlir::Location loc, mlir::Value ptr, bool isVolatile=false, uint64_t alignment=0)
mlir::Value createMinus(mlir::Location loc, mlir::Value input, bool nsw=false)
cir::ConstantOp getConstantInt(mlir::Location loc, mlir::Type ty, int64_t value)
mlir::Value createComplexCreate(mlir::Location loc, mlir::Value real, mlir::Value imag)
cir::PointerType getVoidPtrTy(clang::LangAS langAS=clang::LangAS::Default)
mlir::Value createIsNaN(mlir::Location loc, mlir::Value operand)
cir::IntType getSIntNTy(int n)
mlir::Value createAlignedLoad(mlir::Location loc, mlir::Value ptr, uint64_t alignment)
cir::CallOp createCallOp(mlir::Location loc, mlir::SymbolRefAttr callee, mlir::Type returnType, mlir::ValueRange operands, llvm::ArrayRef< mlir::NamedAttribute > attrs={}, llvm::ArrayRef< mlir::NamedAttrList > argAttrs={}, llvm::ArrayRef< mlir::NamedAttribute > resAttrs={})
cir::StoreOp createStore(mlir::Location loc, mlir::Value val, mlir::Value dst, bool isVolatile=false, mlir::IntegerAttr align={}, cir::SyncScopeKindAttr scope={}, cir::MemOrderAttr order={})
cir::YieldOp createYield(mlir::Location loc, mlir::ValueRange value={})
Create a yield operation.
mlir::Value createLogicalAnd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs)
mlir::Value createAlloca(mlir::Location loc, cir::PointerType addrType, mlir::Type type, llvm::StringRef name, mlir::IntegerAttr alignment, mlir::Value dynAllocSize)
cir::BoolType getBoolTy()
mlir::Value getUnsignedInt(mlir::Location loc, uint64_t val, unsigned numBits)
mlir::Value createComplexReal(mlir::Location loc, mlir::Value operand)
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
SourceManager & getSourceManager()
MangleContext * createMangleContext(const TargetInfo *T=nullptr)
If T is null pointer, assume the target in ASTContext.
const LangOptions & getLangOpts() const
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
const TargetInfo & getTargetInfo() const
QualType getSignedSizeType() const
Return the unique signed counterpart of the integer type corresponding to size_t.
Module * getCurrentNamedModule() const
Get module under construction, nullptr if this is not a C++20 module.
uint64_t getCharWidth() const
Return the size of the character type, in bits.
llvm::Align getAsAlign() const
getAsAlign - Returns Quantity as a valid llvm::Align, Beware llvm::Align assumes power of two 8-bit b...
QuantityType getQuantity() const
getQuantity - Get the raw integer representation of this quantity.
static CharUnits One()
One - Construct a CharUnits quantity of one.
static CharUnits fromQuantity(QuantityType Quantity)
fromQuantity - Construct a CharUnits quantity from a raw integer type.
llvm::vfs::FileSystem & getVirtualFileSystem() const
bool isModuleImplementation() const
Is this a module implementation.
FileManager & getFileManager() const
Exposes information about the current target.
unsigned getMaxAtomicInlineWidth() const
Return the maximum width lock-free atomic operation which can be inlined given the supported features...
const llvm::fltSemantics & getDoubleFormat() const
const llvm::fltSemantics & getHalfFormat() const
const llvm::fltSemantics & getBFloat16Format() const
const llvm::fltSemantics & getLongDoubleFormat() const
const llvm::fltSemantics & getFloatFormat() const
const llvm::fltSemantics & getFloat128Format() const
const llvm::VersionTuple & getSDKVersion() const
Defines the clang::TargetInfo interface.
const internal::VariadicDynCastAllOfMatcher< Decl, VarDecl > varDecl
Matches variable declarations.
bool isHIP(ID Id)
isHIP - Is this a HIP input.
bool isTemplateInstantiation(TemplateSpecializationKind Kind)
Determine whether this template specialization kind refers to an instantiation of an entity (as oppos...
bool CudaFeatureEnabled(llvm::VersionTuple, CudaFeature)
LLVM_READONLY bool isPreprocessingNumberBody(unsigned char c)
Return true if this is the body character of a C preprocessing number, which is [a-zA-Z0-9_.
@ CUDA_USES_FATBIN_REGISTER_END
std::unique_ptr< Pass > createLoweringPreparePass()
static bool opGlobalThreadLocal()
static bool hipModuleCtor()
static bool guardAbortOnException()
static bool opGlobalAnnotations()
static bool opGlobalCtorPriority()
static bool shouldSplitConstantStore()
static bool shouldUseMemSetToInitialize()
static bool opFuncExtraAttrs()
static bool shouldUseBZeroPlusStoresToInitialize()
static bool globalRegistration()
static bool fastMathFlags()
static bool astVarDeclInterface()