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);
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);
99 cir::FuncOp buildCXXGlobalVarDeclInitFunc(cir::GlobalOp op);
102 cir::FuncOp getOrCreateDtorFunc(CIRBaseBuilderTy &builder, cir::GlobalOp op,
103 mlir::Region &dtorRegion,
104 cir::CallOp &dtorCall);
107 void buildCXXGlobalInitFunc();
110 void buildGlobalCtorDtorList();
112 cir::FuncOp buildRuntimeFunction(
113 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
115 cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage);
117 cir::GlobalOp buildRuntimeVariable(
118 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
120 cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage,
121 cir::VisibilityKind visibility = cir::VisibilityKind::Default);
127 llvm::StringMap<FuncOp> cudaKernelMap;
131 void buildCUDAModuleCtor();
132 std::optional<FuncOp> buildCUDAModuleDtor();
133 std::optional<FuncOp> buildCUDARegisterGlobals();
134 void buildCUDARegisterGlobalFunctions(cir::CIRBaseBuilderTy &builder,
135 FuncOp regGlobalFunc);
138 void handleStaticLocal(cir::GlobalOp globalOp, cir::GetGlobalOp getGlobalOp);
147 cir::GlobalOp createGuardGlobalOp(CIRBaseBuilderTy &builder,
148 mlir::Location loc, llvm::StringRef name,
149 cir::IntType guardTy,
150 cir::GlobalLinkageKind linkage);
153 cir::GlobalOp getStaticLocalDeclGuardAddress(llvm::StringRef globalSymName) {
154 auto it = staticLocalDeclGuardMap.find(globalSymName);
155 if (it != staticLocalDeclGuardMap.end())
161 void setStaticLocalDeclGuardAddress(llvm::StringRef globalSymName,
162 cir::GlobalOp guard) {
163 staticLocalDeclGuardMap[globalSymName] = guard;
167 cir::GlobalOp getOrCreateStaticLocalDeclGuardAddress(
168 CIRBaseBuilderTy &builder, cir::GlobalOp globalOp,
169 cir::ASTVarDeclInterface varDecl, cir::IntType guardTy,
170 clang::CharUnits guardAlignment) {
171 llvm::StringRef globalSymName = globalOp.getSymName();
172 cir::GlobalOp guard = getStaticLocalDeclGuardAddress(globalSymName);
175 llvm::StringRef guardName =
176 globalOp.getStaticLocalGuard()->getName().getValue();
179 guard = createGuardGlobalOp(builder, globalOp->getLoc(), guardName,
180 guardTy, globalOp.getLinkage());
181 guard.setInitialValueAttr(cir::IntAttr::get(guardTy, 0));
182 guard.setDSOLocal(globalOp.getDsoLocal());
183 guard.setAlignment(guardAlignment.
getAsAlign().value());
189 bool hasComdat = globalOp.getComdat();
190 const llvm::Triple &triple = astCtx->getTargetInfo().getTriple();
191 if (!
varDecl.isLocalVarDecl() && hasComdat &&
192 (triple.isOSBinFormatELF() || triple.isOSBinFormatWasm())) {
193 globalOp->emitError(
"NYI: guard COMDAT for non-local variables");
195 }
else if (hasComdat && globalOp.isWeakForLinker()) {
196 globalOp->emitError(
"NYI: guard COMDAT for weak linkage");
200 setStaticLocalDeclGuardAddress(globalSymName, guard);
209 clang::ASTContext *astCtx;
212 mlir::ModuleOp mlirModule;
215 llvm::StringMap<uint32_t> dynamicInitializerNames;
216 llvm::SmallVector<cir::FuncOp> dynamicInitializers;
219 llvm::StringMap<cir::GlobalOp> staticLocalDeclGuardMap;
222 llvm::SmallVector<std::pair<std::string, uint32_t>, 4> globalCtorList;
224 llvm::SmallVector<std::pair<std::string, uint32_t>, 4> globalDtorList;
228 bool useARMGuardVarABI()
const {
229 switch (astCtx->getCXXABIKind()) {
230 case clang::TargetCXXABI::GenericARM:
231 case clang::TargetCXXABI::iOS:
232 case clang::TargetCXXABI::WatchOS:
233 case clang::TargetCXXABI::GenericAArch64:
234 case clang::TargetCXXABI::WebAssembly:
244 void emitCXXGuardedInitIf(CIRBaseBuilderTy &builder, cir::GlobalOp globalOp,
245 cir::ASTVarDeclInterface varDecl,
246 mlir::Value guardPtr, cir::PointerType guardPtrTy,
248 auto loc = globalOp->getLoc();
271 mlir::Value acquireResult = acquireCall.getResult();
274 loc, mlir::cast<cir::IntType>(acquireResult.getType()), 0);
275 auto shouldInit = builder.
createCompare(loc, cir::CmpOpKind::ne,
276 acquireResult, acquireZero);
281 cir::IfOp::create(builder, loc, shouldInit,
false,
282 [](mlir::OpBuilder &, mlir::Location) {});
283 mlir::OpBuilder::InsertionGuard insertGuard(builder);
284 builder.setInsertionPointToStart(&ifOp.getThenRegion().front());
291 auto &ctorRegion = globalOp.getCtorRegion();
292 assert(!ctorRegion.empty() &&
"This should never be empty here.");
293 if (!ctorRegion.hasOneBlock())
294 llvm_unreachable(
"Multiple blocks NYI");
295 mlir::Block &block = ctorRegion.front();
296 mlir::Block *insertBlock = builder.getInsertionBlock();
297 insertBlock->getOperations().splice(insertBlock->end(),
298 block.getOperations(), block.begin(),
299 std::prev(block.end()));
300 builder.setInsertionPointToEnd(insertBlock);
301 ctorRegion.getBlocks().clear();
309 mlir::ValueRange{guardPtr});
312 }
else if (!
varDecl.isLocalVarDecl()) {
318 globalOp->emitError(
"NYI: non-threadsafe init for non-local variables");
324 globalOp->emitError(
"NYI: non-threadsafe init for local variables");
331 void setASTContext(clang::ASTContext *
c) { astCtx =
c; }
336cir::GlobalOp LoweringPreparePass::buildRuntimeVariable(
337 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
338 mlir::Type type, cir::GlobalLinkageKind linkage,
339 cir::VisibilityKind visibility) {
340 cir::GlobalOp g = dyn_cast_or_null<cir::GlobalOp>(
341 mlir::SymbolTable::lookupNearestSymbolFrom(
342 mlirModule, mlir::StringAttr::get(mlirModule->getContext(), name)));
344 g = cir::GlobalOp::create(builder, loc, name, type);
346 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
347 mlir::SymbolTable::setSymbolVisibility(
348 g, mlir::SymbolTable::Visibility::Private);
349 g.setGlobalVisibility(visibility);
354cir::FuncOp LoweringPreparePass::buildRuntimeFunction(
355 mlir::OpBuilder &builder, llvm::StringRef name, mlir::Location loc,
356 cir::FuncType type, cir::GlobalLinkageKind linkage) {
357 cir::FuncOp f = dyn_cast_or_null<FuncOp>(SymbolTable::lookupNearestSymbolFrom(
358 mlirModule, StringAttr::get(mlirModule->getContext(), name)));
360 f = cir::FuncOp::create(builder, loc, name, type);
362 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
363 mlir::SymbolTable::setSymbolVisibility(
364 f, mlir::SymbolTable::Visibility::Private);
374 builder.setInsertionPoint(op);
376 mlir::Value src = op.getSrc();
377 mlir::Value imag = builder.
getNullValue(src.getType(), op.getLoc());
383 cir::CastKind elemToBoolKind) {
385 builder.setInsertionPoint(op);
387 mlir::Value src = op.getSrc();
388 if (!mlir::isa<cir::BoolType>(op.getType()))
395 cir::BoolType boolTy = builder.
getBoolTy();
396 mlir::Value srcRealToBool =
397 builder.
createCast(op.getLoc(), elemToBoolKind, srcReal, boolTy);
398 mlir::Value srcImagToBool =
399 builder.
createCast(op.getLoc(), elemToBoolKind, srcImag, boolTy);
400 return builder.
createLogicalOr(op.getLoc(), srcRealToBool, srcImagToBool);
405 cir::CastKind scalarCastKind) {
407 builder.setInsertionPoint(op);
409 mlir::Value src = op.getSrc();
410 auto dstComplexElemTy =
411 mlir::cast<cir::ComplexType>(op.getType()).getElementType();
416 mlir::Value dstReal = builder.
createCast(op.getLoc(), scalarCastKind, srcReal,
418 mlir::Value dstImag = builder.
createCast(op.getLoc(), scalarCastKind, srcImag,
423void LoweringPreparePass::lowerCastOp(cir::CastOp op) {
424 mlir::MLIRContext &ctx = getContext();
425 mlir::Value loweredValue = [&]() -> mlir::Value {
426 switch (op.getKind()) {
427 case cir::CastKind::float_to_complex:
428 case cir::CastKind::int_to_complex:
430 case cir::CastKind::float_complex_to_real:
431 case cir::CastKind::int_complex_to_real:
433 case cir::CastKind::float_complex_to_bool:
435 case cir::CastKind::int_complex_to_bool:
437 case cir::CastKind::float_complex:
439 case cir::CastKind::float_complex_to_int_complex:
441 case cir::CastKind::int_complex:
443 case cir::CastKind::int_complex_to_float_complex:
451 op.replaceAllUsesWith(loweredValue);
458 llvm::StringRef (*libFuncNameGetter)(llvm::APFloat::Semantics),
459 mlir::Location loc, cir::ComplexType ty, mlir::Value lhsReal,
460 mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag) {
461 cir::FPTypeInterface elementTy =
462 mlir::cast<cir::FPTypeInterface>(ty.getElementType());
464 llvm::StringRef libFuncName = libFuncNameGetter(
465 llvm::APFloat::SemanticsToEnum(elementTy.getFloatSemantics()));
468 cir::FuncType libFuncTy = cir::FuncType::get(libFuncInputTypes, ty);
474 mlir::OpBuilder::InsertionGuard ipGuard{builder};
475 builder.setInsertionPointToStart(pass.mlirModule.getBody());
476 libFunc = pass.buildRuntimeFunction(builder, libFuncName, loc, libFuncTy);
480 builder.
createCallOp(loc, libFunc, {lhsReal, lhsImag, rhsReal, rhsImag});
481 return call.getResult();
484static llvm::StringRef
487 case llvm::APFloat::S_IEEEhalf:
489 case llvm::APFloat::S_IEEEsingle:
491 case llvm::APFloat::S_IEEEdouble:
493 case llvm::APFloat::S_PPCDoubleDouble:
495 case llvm::APFloat::S_x87DoubleExtended:
497 case llvm::APFloat::S_IEEEquad:
500 llvm_unreachable(
"unsupported floating point type");
506 mlir::Value lhsReal, mlir::Value lhsImag,
507 mlir::Value rhsReal, mlir::Value rhsImag) {
509 mlir::Value &a = lhsReal;
510 mlir::Value &
b = lhsImag;
511 mlir::Value &
c = rhsReal;
512 mlir::Value &d = rhsImag;
514 mlir::Value ac = builder.
createMul(loc, a,
c);
515 mlir::Value bd = builder.
createMul(loc,
b, d);
517 mlir::Value dd = builder.
createMul(loc, d, d);
518 mlir::Value acbd = builder.
createAdd(loc, ac, bd);
519 mlir::Value ccdd = builder.
createAdd(loc, cc, dd);
520 mlir::Value resultReal = builder.
createDiv(loc, acbd, ccdd);
523 mlir::Value ad = builder.
createMul(loc, a, d);
524 mlir::Value bcad = builder.
createSub(loc, bc, ad);
525 mlir::Value resultImag = builder.
createDiv(loc, bcad, ccdd);
531 mlir::Value lhsReal, mlir::Value lhsImag,
532 mlir::Value rhsReal, mlir::Value rhsImag) {
553 mlir::Value &a = lhsReal;
554 mlir::Value &
b = lhsImag;
555 mlir::Value &
c = rhsReal;
556 mlir::Value &d = rhsImag;
558 auto trueBranchBuilder = [&](mlir::OpBuilder &, mlir::Location) {
560 mlir::Value rd = builder.
createMul(loc, r, d);
561 mlir::Value tmp = builder.
createAdd(loc,
c, rd);
563 mlir::Value br = builder.
createMul(loc,
b, r);
564 mlir::Value abr = builder.
createAdd(loc, a, br);
565 mlir::Value e = builder.
createDiv(loc, abr, tmp);
567 mlir::Value ar = builder.
createMul(loc, a, r);
568 mlir::Value bar = builder.
createSub(loc,
b, ar);
569 mlir::Value f = builder.
createDiv(loc, bar, tmp);
575 auto falseBranchBuilder = [&](mlir::OpBuilder &, mlir::Location) {
577 mlir::Value rc = builder.
createMul(loc, r,
c);
578 mlir::Value tmp = builder.
createAdd(loc, d, rc);
580 mlir::Value ar = builder.
createMul(loc, a, r);
581 mlir::Value arb = builder.
createAdd(loc, ar,
b);
582 mlir::Value e = builder.
createDiv(loc, arb, tmp);
584 mlir::Value br = builder.
createMul(loc,
b, r);
585 mlir::Value bra = builder.
createSub(loc, br, a);
586 mlir::Value f = builder.
createDiv(loc, bra, tmp);
592 auto cFabs = cir::FAbsOp::create(builder, loc,
c);
593 auto dFabs = cir::FAbsOp::create(builder, loc, d);
594 cir::CmpOp cmpResult =
595 builder.
createCompare(loc, cir::CmpOpKind::ge, cFabs, dFabs);
596 auto ternary = cir::TernaryOp::create(builder, loc, cmpResult,
597 trueBranchBuilder, falseBranchBuilder);
599 return ternary.getResult();
606 auto getHigherPrecisionFPType = [&context](mlir::Type type) -> mlir::Type {
607 if (mlir::isa<cir::FP16Type>(type))
608 return cir::SingleType::get(&context);
610 if (mlir::isa<cir::SingleType>(type) || mlir::isa<cir::BF16Type>(type))
611 return cir::DoubleType::get(&context);
613 if (mlir::isa<cir::DoubleType>(type))
614 return cir::LongDoubleType::get(&context, type);
619 auto getFloatTypeSemantics =
620 [&cc](mlir::Type type) ->
const llvm::fltSemantics & {
622 if (mlir::isa<cir::FP16Type>(type))
625 if (mlir::isa<cir::BF16Type>(type))
628 if (mlir::isa<cir::SingleType>(type))
631 if (mlir::isa<cir::DoubleType>(type))
634 if (mlir::isa<cir::LongDoubleType>(type)) {
636 llvm_unreachable(
"NYI Float type semantics with OpenMP");
640 if (mlir::isa<cir::FP128Type>(type)) {
642 llvm_unreachable(
"NYI Float type semantics with OpenMP");
646 llvm_unreachable(
"Unsupported float type semantics");
649 const mlir::Type higherElementType = getHigherPrecisionFPType(elementType);
650 const llvm::fltSemantics &elementTypeSemantics =
651 getFloatTypeSemantics(elementType);
652 const llvm::fltSemantics &higherElementTypeSemantics =
653 getFloatTypeSemantics(higherElementType);
662 if (llvm::APFloat::semanticsMaxExponent(elementTypeSemantics) * 2 + 1 <=
663 llvm::APFloat::semanticsMaxExponent(higherElementTypeSemantics)) {
664 return higherElementType;
674 mlir::Location loc, cir::ComplexDivOp op, mlir::Value lhsReal,
675 mlir::Value lhsImag, mlir::Value rhsReal, mlir::Value rhsImag,
677 cir::ComplexType complexTy = op.getType();
678 if (mlir::isa<cir::FPTypeInterface>(complexTy.getElementType())) {
679 cir::ComplexRangeKind range = op.getRange();
680 if (range == cir::ComplexRangeKind::Improved)
684 if (range == cir::ComplexRangeKind::Full)
686 loc, complexTy, lhsReal, lhsImag, rhsReal,
689 if (range == cir::ComplexRangeKind::Promoted) {
690 mlir::Type originalElementType = complexTy.getElementType();
691 mlir::Type higherPrecisionElementType =
693 originalElementType);
695 if (!higherPrecisionElementType)
699 cir::CastKind floatingCastKind = cir::CastKind::floating;
700 lhsReal = builder.
createCast(floatingCastKind, lhsReal,
701 higherPrecisionElementType);
702 lhsImag = builder.
createCast(floatingCastKind, lhsImag,
703 higherPrecisionElementType);
704 rhsReal = builder.
createCast(floatingCastKind, rhsReal,
705 higherPrecisionElementType);
706 rhsImag = builder.
createCast(floatingCastKind, rhsImag,
707 higherPrecisionElementType);
710 builder, loc, lhsReal, lhsImag, rhsReal, rhsImag);
715 mlir::Value finalReal =
716 builder.
createCast(floatingCastKind, resultReal, originalElementType);
717 mlir::Value finalImag =
718 builder.
createCast(floatingCastKind, resultImag, originalElementType);
727void LoweringPreparePass::lowerComplexDivOp(cir::ComplexDivOp op) {
728 cir::CIRBaseBuilderTy builder(getContext());
729 builder.setInsertionPointAfter(op);
730 mlir::Location loc = op.getLoc();
731 mlir::TypedValue<cir::ComplexType> lhs = op.getLhs();
732 mlir::TypedValue<cir::ComplexType> rhs = op.getRhs();
738 mlir::Value loweredResult =
740 rhsImag, getContext(), *astCtx);
741 op.replaceAllUsesWith(loweredResult);
745static llvm::StringRef
748 case llvm::APFloat::S_IEEEhalf:
750 case llvm::APFloat::S_IEEEsingle:
752 case llvm::APFloat::S_IEEEdouble:
754 case llvm::APFloat::S_PPCDoubleDouble:
756 case llvm::APFloat::S_x87DoubleExtended:
758 case llvm::APFloat::S_IEEEquad:
761 llvm_unreachable(
"unsupported floating point type");
767 mlir::Location loc, cir::ComplexMulOp op,
768 mlir::Value lhsReal, mlir::Value lhsImag,
769 mlir::Value rhsReal, mlir::Value rhsImag) {
771 mlir::Value resultRealLhs = builder.
createMul(loc, lhsReal, rhsReal);
772 mlir::Value resultRealRhs = builder.
createMul(loc, lhsImag, rhsImag);
773 mlir::Value resultImagLhs = builder.
createMul(loc, lhsReal, rhsImag);
774 mlir::Value resultImagRhs = builder.
createMul(loc, lhsImag, rhsReal);
775 mlir::Value resultReal = builder.
createSub(loc, resultRealLhs, resultRealRhs);
776 mlir::Value resultImag = builder.
createAdd(loc, resultImagLhs, resultImagRhs);
777 mlir::Value algebraicResult =
780 cir::ComplexType complexTy = op.getType();
781 cir::ComplexRangeKind rangeKind = op.getRange();
782 if (mlir::isa<cir::IntType>(complexTy.getElementType()) ||
783 rangeKind == cir::ComplexRangeKind::Basic ||
784 rangeKind == cir::ComplexRangeKind::Improved ||
785 rangeKind == cir::ComplexRangeKind::Promoted)
786 return algebraicResult;
793 mlir::Value resultRealIsNaN = builder.
createIsNaN(loc, resultReal);
794 mlir::Value resultImagIsNaN = builder.
createIsNaN(loc, resultImag);
795 mlir::Value resultRealAndImagAreNaN =
798 return cir::TernaryOp::create(
799 builder, loc, resultRealAndImagAreNaN,
800 [&](mlir::OpBuilder &, mlir::Location) {
803 lhsReal, lhsImag, rhsReal, rhsImag);
806 [&](mlir::OpBuilder &, mlir::Location) {
812void LoweringPreparePass::lowerComplexMulOp(cir::ComplexMulOp op) {
813 cir::CIRBaseBuilderTy builder(getContext());
814 builder.setInsertionPointAfter(op);
815 mlir::Location loc = op.getLoc();
816 mlir::TypedValue<cir::ComplexType> lhs = op.getLhs();
817 mlir::TypedValue<cir::ComplexType> rhs = op.getRhs();
822 mlir::Value loweredResult =
lowerComplexMul(*
this, builder, loc, op, lhsReal,
823 lhsImag, rhsReal, rhsImag);
824 op.replaceAllUsesWith(loweredResult);
828void LoweringPreparePass::lowerUnaryOp(cir::UnaryOpInterface op) {
829 if (!mlir::isa<cir::ComplexType>(op.getResult().getType()))
832 mlir::Location loc = op->getLoc();
833 CIRBaseBuilderTy builder(getContext());
834 builder.setInsertionPointAfter(op);
836 mlir::Value operand = op.getInput();
840 mlir::Value resultReal = operandReal;
841 mlir::Value resultImag = operandImag;
843 llvm::TypeSwitch<mlir::Operation *>(op)
845 [&](
auto) { resultReal = builder.
createInc(loc, operandReal); })
847 [&](
auto) { resultReal = builder.
createDec(loc, operandReal); })
848 .Case<cir::MinusOp>([&](
auto) {
849 resultReal = builder.
createMinus(loc, operandReal);
850 resultImag = builder.
createMinus(loc, operandImag);
853 [&](
auto) { resultImag = builder.
createMinus(loc, operandImag); })
854 .
Default([](
auto) { llvm_unreachable(
"unhandled unary complex op"); });
857 op->replaceAllUsesWith(mlir::ValueRange{result});
861cir::FuncOp LoweringPreparePass::getOrCreateDtorFunc(CIRBaseBuilderTy &builder,
863 mlir::Region &dtorRegion,
864 cir::CallOp &dtorCall) {
865 mlir::OpBuilder::InsertionGuard guard(builder);
869 cir::VoidType voidTy = builder.
getVoidTy();
870 auto voidPtrTy = cir::PointerType::get(voidTy);
873 mlir::Block &dtorBlock = dtorRegion.front();
877 auto opIt = dtorBlock.getOperations().begin();
878 cir::GetGlobalOp ggop = mlir::cast<cir::GetGlobalOp>(*opIt);
889 if (dtorBlock.getOperations().size() == 3) {
890 auto callOp = mlir::dyn_cast<cir::CallOp>(&*(++opIt));
891 auto yieldOp = mlir::dyn_cast<cir::YieldOp>(&*(++opIt));
892 if (yieldOp && callOp && callOp.getNumOperands() == 1 &&
893 callOp.getArgOperand(0) == ggop) {
902 builder.setInsertionPointAfter(op);
903 SmallString<256> fnName(
"__cxx_global_array_dtor");
904 uint32_t cnt = dynamicInitializerNames[fnName]++;
906 fnName +=
"." + std::to_string(cnt);
909 auto fnType = cir::FuncType::get({voidPtrTy}, voidTy);
910 cir::FuncOp dtorFunc =
911 buildRuntimeFunction(builder, fnName, op.getLoc(), fnType,
912 cir::GlobalLinkageKind::InternalLinkage);
914 SmallVector<mlir::NamedAttribute> paramAttrs;
915 paramAttrs.push_back(
916 builder.getNamedAttr(
"llvm.noundef", builder.getUnitAttr()));
917 SmallVector<mlir::Attribute> argAttrDicts;
918 argAttrDicts.push_back(
919 mlir::DictionaryAttr::get(builder.getContext(), paramAttrs));
920 dtorFunc.setArgAttrsAttr(
921 mlir::ArrayAttr::get(builder.getContext(), argAttrDicts));
923 mlir::Block *entryBB = dtorFunc.addEntryBlock();
926 entryBB->getOperations().splice(entryBB->begin(), dtorBlock.getOperations(),
927 dtorBlock.begin(), dtorBlock.end());
930 cir::GetGlobalOp dtorGGop =
931 mlir::cast<cir::GetGlobalOp>(entryBB->getOperations().front());
932 builder.setInsertionPointToStart(&dtorBlock);
933 builder.clone(*dtorGGop.getOperation());
937 mlir::Value dtorArg = entryBB->getArgument(0);
938 dtorGGop.replaceAllUsesWith(dtorArg);
942 mlir::Block &finalBlock = dtorFunc.getBody().back();
943 auto yieldOp = cast<cir::YieldOp>(finalBlock.getTerminator());
944 builder.setInsertionPoint(yieldOp);
945 cir::ReturnOp::create(builder, yieldOp->getLoc());
950 cir::GetGlobalOp origGGop =
951 mlir::cast<cir::GetGlobalOp>(dtorBlock.getOperations().front());
952 builder.setInsertionPointAfter(origGGop);
953 mlir::Value ggopResult = origGGop.getResult();
954 dtorCall = builder.
createCallOp(op.getLoc(), dtorFunc, ggopResult);
957 auto finalYield = cir::YieldOp::create(builder, op.getLoc());
960 dtorBlock.getOperations().erase(std::next(mlir::Block::iterator(finalYield)),
962 dtorRegion.getBlocks().erase(std::next(dtorRegion.begin()), dtorRegion.end());
968LoweringPreparePass::buildCXXGlobalVarDeclInitFunc(cir::GlobalOp op) {
971 SmallString<256> fnName(
"__cxx_global_var_init");
973 uint32_t cnt = dynamicInitializerNames[fnName]++;
975 fnName +=
"." + std::to_string(cnt);
978 CIRBaseBuilderTy builder(getContext());
979 builder.setInsertionPointAfter(op);
980 cir::VoidType voidTy = builder.
getVoidTy();
981 auto fnType = cir::FuncType::get({}, voidTy);
982 FuncOp f = buildRuntimeFunction(builder, fnName, op.getLoc(), fnType,
983 cir::GlobalLinkageKind::InternalLinkage);
986 mlir::Block *entryBB = f.addEntryBlock();
987 if (!op.getCtorRegion().empty()) {
988 mlir::Block &block = op.getCtorRegion().front();
989 entryBB->getOperations().splice(entryBB->begin(), block.getOperations(),
990 block.begin(), std::prev(block.end()));
994 mlir::Region &dtorRegion = op.getDtorRegion();
995 if (!dtorRegion.empty()) {
1000 builder.setInsertionPointToStart(&mlirModule.getBodyRegion().front());
1001 cir::GlobalOp handle = buildRuntimeVariable(
1002 builder,
"__dso_handle", op.getLoc(), builder.getI8Type(),
1003 cir::GlobalLinkageKind::ExternalLinkage, cir::VisibilityKind::Hidden);
1009 cir::CallOp dtorCall;
1010 cir::FuncOp dtorFunc =
1011 getOrCreateDtorFunc(builder, op, dtorRegion, dtorCall);
1015 auto voidPtrTy = cir::PointerType::get(voidTy);
1016 auto voidFnTy = cir::FuncType::get({voidPtrTy}, voidTy);
1017 auto voidFnPtrTy = cir::PointerType::get(voidFnTy);
1018 auto handlePtrTy = cir::PointerType::get(handle.getSymType());
1020 cir::FuncType::get({voidFnPtrTy, voidPtrTy, handlePtrTy}, voidTy);
1021 const char *nameAtExit =
"__cxa_atexit";
1022 cir::FuncOp fnAtExit =
1023 buildRuntimeFunction(builder, nameAtExit, op.getLoc(), fnAtExitType);
1027 builder.setInsertionPointAfter(dtorCall);
1028 mlir::Value args[3];
1029 auto dtorPtrTy = cir::PointerType::get(dtorFunc.getFunctionType());
1031 args[0] = cir::GetGlobalOp::create(builder, dtorCall.getLoc(), dtorPtrTy,
1032 dtorFunc.getSymName());
1033 args[0] = cir::CastOp::create(builder, dtorCall.getLoc(), voidFnPtrTy,
1034 cir::CastKind::bitcast, args[0]);
1036 cir::CastOp::create(builder, dtorCall.getLoc(), voidPtrTy,
1037 cir::CastKind::bitcast, dtorCall.getArgOperand(0));
1038 args[2] = cir::GetGlobalOp::create(builder, handle.getLoc(), handlePtrTy,
1039 handle.getSymName());
1040 builder.
createCallOp(dtorCall.getLoc(), fnAtExit, args);
1042 mlir::Block &dtorBlock = dtorRegion.front();
1043 entryBB->getOperations().splice(entryBB->end(), dtorBlock.getOperations(),
1045 std::prev(dtorBlock.end()));
1049 builder.setInsertionPointToEnd(entryBB);
1050 mlir::Operation *yieldOp =
nullptr;
1051 if (!op.getCtorRegion().empty()) {
1052 mlir::Block &block = op.getCtorRegion().front();
1053 yieldOp = &block.getOperations().back();
1055 assert(!dtorRegion.empty());
1056 mlir::Block &block = dtorRegion.front();
1057 yieldOp = &block.getOperations().back();
1060 assert(isa<cir::YieldOp>(*yieldOp));
1061 cir::ReturnOp::create(builder, yieldOp->getLoc());
1066LoweringPreparePass::getGuardAcquireFn(cir::PointerType guardPtrTy) {
1068 CIRBaseBuilderTy builder(getContext());
1069 mlir::OpBuilder::InsertionGuard ipGuard{builder};
1070 builder.setInsertionPointToStart(mlirModule.getBody());
1071 mlir::Location loc = mlirModule.getLoc();
1072 cir::IntType intTy = cir::IntType::get(&getContext(), 32,
true);
1073 auto fnType = cir::FuncType::get({guardPtrTy}, intTy);
1074 return buildRuntimeFunction(builder,
"__cxa_guard_acquire", loc, fnType);
1078LoweringPreparePass::getGuardReleaseFn(cir::PointerType guardPtrTy) {
1080 CIRBaseBuilderTy builder(getContext());
1081 mlir::OpBuilder::InsertionGuard ipGuard{builder};
1082 builder.setInsertionPointToStart(mlirModule.getBody());
1083 mlir::Location loc = mlirModule.getLoc();
1084 cir::VoidType voidTy = cir::VoidType::get(&getContext());
1085 auto fnType = cir::FuncType::get({guardPtrTy}, voidTy);
1086 return buildRuntimeFunction(builder,
"__cxa_guard_release", loc, fnType);
1089cir::GlobalOp LoweringPreparePass::createGuardGlobalOp(
1090 CIRBaseBuilderTy &builder, mlir::Location loc, llvm::StringRef name,
1091 cir::IntType guardTy, cir::GlobalLinkageKind linkage) {
1092 mlir::OpBuilder::InsertionGuard guard(builder);
1093 builder.setInsertionPointToStart(mlirModule.getBody());
1094 cir::GlobalOp g = cir::GlobalOp::create(builder, loc, name, guardTy);
1096 cir::GlobalLinkageKindAttr::get(builder.getContext(), linkage));
1097 mlir::SymbolTable::setSymbolVisibility(
1098 g, mlir::SymbolTable::Visibility::Private);
1102void LoweringPreparePass::handleStaticLocal(cir::GlobalOp globalOp,
1103 cir::GetGlobalOp getGlobalOp) {
1104 CIRBaseBuilderTy builder(getContext());
1106 std::optional<cir::ASTVarDeclInterface> astOption = globalOp.getAst();
1107 assert(astOption.has_value());
1108 cir::ASTVarDeclInterface
varDecl = astOption.value();
1110 builder.setInsertionPointAfter(getGlobalOp);
1111 mlir::Block *getGlobalOpBlock = builder.getInsertionBlock();
1114 mlir::Operation *ret = getGlobalOpBlock->getTerminator();
1116 builder.setInsertionPointAfter(getGlobalOp);
1120 bool nonTemplateInline =
1126 if (nonTemplateInline) {
1127 globalOp->emitError(
1128 "NYI: guarded initialization for inline namespace-scope variables");
1135 bool threadsafe = astCtx->
getLangOpts().ThreadsafeStatics &&
1136 (
varDecl.isLocalVarDecl() || nonTemplateInline) &&
1141 globalOp->emitError(
"NYI: guarded initialization for thread-local statics");
1147 bool useInt8GuardVariable = !threadsafe && globalOp.hasInternalLinkage();
1148 if (useInt8GuardVariable) {
1149 globalOp->emitError(
"NYI: int8 guard variables for non-threadsafe statics");
1155 if (useARMGuardVarABI()) {
1156 globalOp->emitError(
"NYI: ARM-style guard variables for static locals");
1159 cir::IntType guardTy =
1160 cir::IntType::get(&getContext(), 64,
true);
1161 cir::CIRDataLayout dataLayout(mlirModule);
1162 clang::CharUnits guardAlignment =
1164 auto guardPtrTy = cir::PointerType::get(guardTy);
1167 cir::GlobalOp guard = getOrCreateStaticLocalDeclGuardAddress(
1168 builder, globalOp, varDecl, guardTy, guardAlignment);
1171 getGlobalOpBlock->push_back(ret);
1197 unsigned maxInlineWidthInBits =
1200 if (!threadsafe || maxInlineWidthInBits) {
1202 auto bytePtrTy = cir::PointerType::get(builder.
getSIntNTy(8));
1203 mlir::Value bytePtr = builder.
createBitcast(guardPtr, bytePtrTy);
1205 getGlobalOp.getLoc(), bytePtr, guardAlignment.
getAsAlign().value());
1214 auto loadOp = mlir::cast<cir::LoadOp>(guardLoad.getDefiningOp());
1215 loadOp.setMemOrder(cir::MemOrder::Acquire);
1216 loadOp.setSyncScope(cir::SyncScopeKind::System);
1239 if (useARMGuardVarABI()) {
1240 globalOp->emitError(
1241 "NYI: ARM-style guard variable check (bit 0 only) for static locals");
1247 getGlobalOp.getLoc(), mlir::cast<cir::IntType>(guardLoad.getType()), 0);
1248 auto needsInit = builder.
createCompare(getGlobalOp.getLoc(),
1249 cir::CmpOpKind::eq, guardLoad, zero);
1252 cir::IfOp::create(builder, globalOp.getLoc(), needsInit,
1254 [&](mlir::OpBuilder &, mlir::Location) {
1255 emitCXXGuardedInitIf(builder, globalOp, varDecl,
1256 guardPtr, guardPtrTy, threadsafe);
1261 globalOp->emitError(
"NYI: guarded init without inline atomics support");
1266 builder.getInsertionBlock()->push_back(ret);
1269void LoweringPreparePass::lowerGlobalOp(GlobalOp op) {
1271 if (op.getStaticLocalGuard())
1274 mlir::Region &ctorRegion = op.getCtorRegion();
1275 mlir::Region &dtorRegion = op.getDtorRegion();
1277 if (!ctorRegion.empty() || !dtorRegion.empty()) {
1280 cir::FuncOp f = buildCXXGlobalVarDeclInitFunc(op);
1283 ctorRegion.getBlocks().clear();
1284 dtorRegion.getBlocks().clear();
1287 dynamicInitializers.push_back(f);
1293void LoweringPreparePass::lowerThreeWayCmpOp(CmpThreeWayOp op) {
1294 CIRBaseBuilderTy builder(getContext());
1295 builder.setInsertionPointAfter(op);
1297 mlir::Location loc = op->getLoc();
1298 cir::CmpThreeWayInfoAttr cmpInfo = op.getInfo();
1307 mlir::Value transformedResult;
1308 if (cmpInfo.getOrdering() != CmpOrdering::Partial) {
1311 builder.
createCompare(loc, CmpOpKind::lt, op.getLhs(), op.getRhs());
1312 mlir::Value selectOnLt = builder.
createSelect(loc, lt, ltRes, gtRes);
1314 builder.
createCompare(loc, CmpOpKind::eq, op.getLhs(), op.getRhs());
1315 transformedResult = builder.
createSelect(loc, eq, eqRes, selectOnLt);
1319 loc, op.getType(), cmpInfo.getUnordered().value());
1322 builder.
createCompare(loc, CmpOpKind::eq, op.getLhs(), op.getRhs());
1323 mlir::Value selectOnEq = builder.
createSelect(loc, eq, eqRes, unorderedRes);
1325 builder.
createCompare(loc, CmpOpKind::gt, op.getLhs(), op.getRhs());
1326 mlir::Value selectOnGt = builder.
createSelect(loc, gt, gtRes, selectOnEq);
1328 builder.
createCompare(loc, CmpOpKind::lt, op.getLhs(), op.getRhs());
1329 transformedResult = builder.
createSelect(loc, lt, ltRes, selectOnGt);
1332 op.replaceAllUsesWith(transformedResult);
1336template <
typename AttributeTy>
1337static llvm::SmallVector<mlir::Attribute>
1341 for (
const auto &[name, priority] : list)
1342 attrs.push_back(AttributeTy::get(context, name, priority));
1346void LoweringPreparePass::buildGlobalCtorDtorList() {
1347 if (!globalCtorList.empty()) {
1348 llvm::SmallVector<mlir::Attribute> globalCtors =
1352 mlirModule->setAttr(cir::CIRDialect::getGlobalCtorsAttrName(),
1353 mlir::ArrayAttr::get(&getContext(), globalCtors));
1356 if (!globalDtorList.empty()) {
1357 llvm::SmallVector<mlir::Attribute> globalDtors =
1360 mlirModule->setAttr(cir::CIRDialect::getGlobalDtorsAttrName(),
1361 mlir::ArrayAttr::get(&getContext(), globalDtors));
1365void LoweringPreparePass::buildCXXGlobalInitFunc() {
1366 if (dynamicInitializers.empty())
1373 SmallString<256> fnName;
1381 llvm::raw_svector_ostream
out(fnName);
1382 std::unique_ptr<clang::MangleContext> mangleCtx(
1384 cast<clang::ItaniumMangleContext>(*mangleCtx)
1387 fnName +=
"_GLOBAL__sub_I_";
1391 CIRBaseBuilderTy builder(getContext());
1392 builder.setInsertionPointToEnd(&mlirModule.getBodyRegion().back());
1393 auto fnType = cir::FuncType::get({}, builder.
getVoidTy());
1395 buildRuntimeFunction(builder, fnName, mlirModule.getLoc(), fnType,
1396 cir::GlobalLinkageKind::ExternalLinkage);
1397 builder.setInsertionPointToStart(f.addEntryBlock());
1398 for (cir::FuncOp &f : dynamicInitializers)
1402 globalCtorList.emplace_back(fnName,
1403 cir::GlobalCtorAttr::getDefaultPriority());
1405 cir::ReturnOp::create(builder, f.getLoc());
1414 mlir::Operation *op, mlir::Type eltTy,
1416 mlir::Value numElements,
1417 uint64_t arrayLen,
bool isCtor) {
1418 mlir::Location loc = op->getLoc();
1419 bool isDynamic = numElements !=
nullptr;
1423 const unsigned sizeTypeSize =
1429 mlir::Value begin, end;
1432 end = cir::PtrStrideOp::create(builder, loc, eltTy, begin, numElements);
1434 mlir::Value endOffsetVal =
1436 begin = cir::CastOp::create(builder, loc, eltTy,
1437 cir::CastKind::array_to_ptrdecay, addr);
1438 end = cir::PtrStrideOp::create(builder, loc, eltTy, begin, endOffsetVal);
1441 mlir::Value start = isCtor ? begin : end;
1442 mlir::Value stop = isCtor ? end : begin;
1448 mlir::Value guardCond;
1451 guardCond = cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne,
1457 cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne, start, stop);
1459 ifOp = cir::IfOp::create(builder, loc, guardCond,
1461 [&](mlir::OpBuilder &, mlir::Location) {});
1462 builder.setInsertionPointToStart(&ifOp.getThenRegion().front());
1470 mlir::Block *bodyBlock = &op->getRegion(0).front();
1475 auto cloneRegionBodyInto = [&](mlir::Block *srcBlock,
1476 mlir::Value replacement) {
1477 mlir::IRMapping map;
1478 map.map(srcBlock->getArgument(0), replacement);
1479 for (mlir::Operation ®ionOp : *srcBlock) {
1480 if (!mlir::isa<cir::YieldOp>(®ionOp))
1481 builder.clone(regionOp, map);
1485 mlir::Block *partialDtorBlock =
nullptr;
1486 if (
auto arrayCtor = mlir::dyn_cast<cir::ArrayCtor>(op)) {
1487 mlir::Region &partialDtor = arrayCtor.getPartialDtor();
1488 if (!partialDtor.empty())
1489 partialDtorBlock = &partialDtor.front();
1492 auto emitCtorDtorLoop = [&]() {
1496 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1497 auto currentElement = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1498 auto cmp = cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne,
1499 currentElement, stop);
1503 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1504 auto currentElement = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1506 cloneRegionBodyInto(bodyBlock, currentElement);
1507 mlir::Value stride = builder.
getUnsignedInt(loc, 1, sizeTypeSize);
1508 auto nextElement = cir::PtrStrideOp::create(builder, loc, eltTy,
1509 currentElement, stride);
1512 mlir::Value stride = builder.
getSignedInt(loc, -1, sizeTypeSize);
1513 auto prevElement = cir::PtrStrideOp::create(builder, loc, eltTy,
1514 currentElement, stride);
1516 cloneRegionBodyInto(bodyBlock, prevElement);
1519 cir::YieldOp::create(
b, loc);
1523 if (partialDtorBlock) {
1524 cir::CleanupScopeOp::create(
1525 builder, loc, cir::CleanupKind::EH,
1527 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1529 cir::YieldOp::create(
b, loc);
1532 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1533 auto cur = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1535 cir::CmpOp::create(builder, loc, cir::CmpOpKind::ne, cur, begin);
1537 builder, loc, cmp,
false,
1538 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1542 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1543 auto el = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1544 auto neq = cir::CmpOp::create(
1545 builder, loc, cir::CmpOpKind::ne, el, begin);
1549 [&](mlir::OpBuilder &
b, mlir::Location loc) {
1550 auto el = cir::LoadOp::create(
b, loc, eltTy, tmpAddr);
1551 mlir::Value negOne =
1553 auto prev = cir::PtrStrideOp::create(builder, loc, eltTy,
1556 cloneRegionBodyInto(partialDtorBlock, prev);
1559 cir::YieldOp::create(builder, loc);
1561 cir::YieldOp::create(
b, loc);
1568 cir::YieldOp::create(builder, loc);
1573void LoweringPreparePass::lowerArrayDtor(cir::ArrayDtor op) {
1574 CIRBaseBuilderTy builder(getContext());
1575 builder.setInsertionPointAfter(op.getOperation());
1577 mlir::Type eltTy = op->getRegion(0).getArgument(0).getType();
1579 if (op.getNumElements()) {
1581 op.getNumElements(), 0,
1587 mlir::cast<cir::ArrayType>(op.getAddr().getType().getPointee()).getSize();
1593void LoweringPreparePass::lowerArrayCtor(cir::ArrayCtor op) {
1594 cir::CIRBaseBuilderTy builder(getContext());
1595 builder.setInsertionPointAfter(op.getOperation());
1597 mlir::Type eltTy = op->getRegion(0).getArgument(0).getType();
1599 if (op.getNumElements()) {
1601 op.getNumElements(), 0,
1607 mlir::cast<cir::ArrayType>(op.getAddr().getType().getPointee()).getSize();
1613void LoweringPreparePass::lowerTrivialCopyCall(cir::CallOp op) {
1618 std::optional<cir::CtorKind> ctorKind = funcOp.getCxxConstructorKind();
1619 if (ctorKind && *ctorKind == cir::CtorKind::Copy &&
1620 funcOp.isCxxTrivialMemberFunction()) {
1622 CIRBaseBuilderTy builder(getContext());
1623 mlir::ValueRange operands = op.getOperands();
1624 mlir::Value dest = operands[0];
1625 mlir::Value src = operands[1];
1626 builder.setInsertionPoint(op);
1632void LoweringPreparePass::lowerStoreOfConstAggregate(cir::StoreOp op) {
1634 auto constOp = op.getValue().getDefiningOp<cir::ConstantOp>();
1638 mlir::Type ty = constOp.getType();
1639 if (!mlir::isa<cir::ArrayType, cir::RecordType>(ty))
1645 auto alloca = op.getAddr().getDefiningOp<cir::AllocaOp>();
1649 mlir::TypedAttr constant = constOp.getValue();
1660 auto func = op->getParentOfType<cir::FuncOp>();
1663 llvm::StringRef funcName = func.getSymName();
1666 llvm::StringRef varName = alloca.getName();
1669 std::string
name = (
"__const." + funcName +
"." + varName).str();
1672 CIRBaseBuilderTy builder(getContext());
1675 builder.setInsertionPointToStart(mlirModule.getBody());
1679 if (!mlir::SymbolTable::lookupSymbolIn(
1680 mlirModule, mlir::StringAttr::get(&getContext(), name))) {
1681 auto gv = cir::GlobalOp::create(
1682 builder, op.getLoc(), name, ty,
1684 cir::LangAddressSpaceAttr::get(&getContext(),
1685 cir::LangAddressSpace::Default),
1686 cir::GlobalLinkageKind::PrivateLinkage);
1687 mlir::SymbolTable::setSymbolVisibility(
1688 gv, mlir::SymbolTable::Visibility::Private);
1689 gv.setInitialValueAttr(constant);
1693 builder.setInsertionPoint(op);
1695 auto ptrTy = cir::PointerType::get(ty);
1696 mlir::Value globalPtr =
1697 cir::GetGlobalOp::create(builder, op.getLoc(), ptrTy, name);
1706 if (constOp.use_empty())
1710void LoweringPreparePass::runOnOp(mlir::Operation *op) {
1711 if (
auto arrayCtor = dyn_cast<cir::ArrayCtor>(op)) {
1712 lowerArrayCtor(arrayCtor);
1713 }
else if (
auto arrayDtor = dyn_cast<cir::ArrayDtor>(op)) {
1714 lowerArrayDtor(arrayDtor);
1715 }
else if (
auto cast = mlir::dyn_cast<cir::CastOp>(op)) {
1717 }
else if (
auto complexDiv = mlir::dyn_cast<cir::ComplexDivOp>(op)) {
1718 lowerComplexDivOp(complexDiv);
1719 }
else if (
auto complexMul = mlir::dyn_cast<cir::ComplexMulOp>(op)) {
1720 lowerComplexMulOp(complexMul);
1721 }
else if (
auto glob = mlir::dyn_cast<cir::GlobalOp>(op)) {
1722 lowerGlobalOp(glob);
1723 }
else if (
auto getGlobal = mlir::dyn_cast<cir::GetGlobalOp>(op)) {
1727 if (getGlobal.getStaticLocal() &&
1728 getGlobal->getParentOfType<cir::FuncOp>()) {
1729 auto globalOp = mlir::dyn_cast_or_null<cir::GlobalOp>(
1730 mlir::SymbolTable::lookupNearestSymbolFrom(getGlobal,
1731 getGlobal.getNameAttr()));
1736 if (globalOp && globalOp.getStaticLocalGuard() &&
1737 !globalOp.getCtorRegion().empty())
1738 handleStaticLocal(globalOp, getGlobal);
1740 }
else if (
auto unaryOp = mlir::dyn_cast<cir::UnaryOpInterface>(op)) {
1741 lowerUnaryOp(unaryOp);
1742 }
else if (
auto callOp = dyn_cast<cir::CallOp>(op)) {
1743 lowerTrivialCopyCall(callOp);
1744 }
else if (
auto storeOp = dyn_cast<cir::StoreOp>(op)) {
1745 lowerStoreOfConstAggregate(storeOp);
1746 }
else if (
auto fnOp = dyn_cast<cir::FuncOp>(op)) {
1747 if (
auto globalCtor = fnOp.getGlobalCtorPriority())
1748 globalCtorList.emplace_back(fnOp.getName(), globalCtor.value());
1749 else if (
auto globalDtor = fnOp.getGlobalDtorPriority())
1750 globalDtorList.emplace_back(fnOp.getName(), globalDtor.value());
1752 if (mlir::Attribute attr =
1753 fnOp->getAttr(cir::CUDAKernelNameAttr::getMnemonic())) {
1754 auto kernelNameAttr = dyn_cast<CUDAKernelNameAttr>(attr);
1755 llvm::StringRef kernelName = kernelNameAttr.getKernelName();
1756 cudaKernelMap[kernelName] = fnOp;
1758 }
else if (
auto threeWayCmp = dyn_cast<cir::CmpThreeWayOp>(op)) {
1759 lowerThreeWayCmpOp(threeWayCmp);
1770 llvm::StringRef name) {
1771 return (
"__" + prefix + name).str();
1793void LoweringPreparePass::buildCUDAModuleCtor() {
1798 if (astCtx->
getLangOpts().GPURelocatableDeviceCode)
1799 llvm_unreachable(
"GPU RDC NYI");
1803 if (cudaKernelMap.empty())
1808 mlir::Attribute cudaBinaryHandleAttr =
1809 mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName());
1810 if (!cudaBinaryHandleAttr) {
1816 llvm::StringRef cudaGPUBinaryName =
1817 mlir::cast<CUDABinaryHandleAttr>(cudaBinaryHandleAttr)
1821 llvm::vfs::FileSystem &vfs =
1823 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> gpuBinaryOrErr =
1824 vfs.getBufferForFile(cudaGPUBinaryName);
1825 if (std::error_code ec = gpuBinaryOrErr.getError()) {
1826 mlirModule->emitError(
"cannot open GPU binary file: " + cudaGPUBinaryName +
1827 ": " + ec.message());
1830 std::unique_ptr<llvm::MemoryBuffer> gpuBinary =
1831 std::move(gpuBinaryOrErr.get());
1835 mlir::Location loc = mlirModule->getLoc();
1836 CIRBaseBuilderTy builder(getContext());
1837 builder.setInsertionPointToStart(mlirModule.getBody());
1841 PointerType voidPtrPtrTy = builder.
getPointerTo(voidPtrTy);
1843 IntType charTy = cir::IntType::get(&getContext(), astCtx->
getCharWidth(),
1849 llvm::StringRef fatbinConstName =
1850 astCtx->
getLangOpts().HIP ?
".hip_fatbin" :
".nv_fatbin";
1852 llvm::StringRef fatbinSectionName =
1853 astCtx->
getLangOpts().HIP ?
".hipFatBinSegment" :
".nvFatBinSegment";
1857 ArrayType::get(&getContext(), charTy, gpuBinary->getBuffer().size());
1859 GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType,
1861 GlobalLinkageKind::PrivateLinkage);
1862 fatbinStr.setAlignment(8);
1863 fatbinStr.setInitialValueAttr(cir::ConstArrayAttr::get(
1864 fatbinType, builder.getStringAttr(gpuBinary->getBuffer())));
1865 fatbinStr.setSection(fatbinConstName);
1866 fatbinStr.setPrivate();
1870 auto fatbinWrapperType = RecordType::get(
1871 &getContext(), {intTy, intTy, voidPtrTy, voidPtrTy},
1872 false,
false, RecordType::RecordKind::Struct);
1873 std::string fatbinWrapperName =
1875 GlobalOp fatbinWrapper = GlobalOp::create(
1876 builder, loc, fatbinWrapperName, fatbinWrapperType,
1877 true, {}, GlobalLinkageKind::PrivateLinkage);
1878 fatbinWrapper.setSection(fatbinSectionName);
1880 constexpr unsigned cudaFatMagic = 0x466243b1;
1881 constexpr unsigned hipFatMagic = 0x48495046;
1882 unsigned fatMagic =
isHIP ? hipFatMagic : cudaFatMagic;
1884 auto magicInit = IntAttr::get(intTy, fatMagic);
1885 auto versionInit = IntAttr::get(intTy, 1);
1886 auto fatbinStrSymbol =
1887 mlir::FlatSymbolRefAttr::get(fatbinStr.getSymNameAttr());
1888 auto fatbinInit = GlobalViewAttr::get(voidPtrTy, fatbinStrSymbol);
1890 fatbinWrapper.setInitialValueAttr(cir::ConstRecordAttr::get(
1892 mlir::ArrayAttr::get(&getContext(),
1893 {magicInit, versionInit, fatbinInit, unusedInit})));
1896 std::string gpubinHandleName =
1899 GlobalOp gpuBinHandle = GlobalOp::create(
1900 builder, loc, gpubinHandleName, voidPtrPtrTy,
1901 false, {}, cir::GlobalLinkageKind::InternalLinkage);
1903 gpuBinHandle.setPrivate();
1908 std::string regFuncName =
1910 FuncType regFuncType = FuncType::get({voidPtrTy}, voidPtrPtrTy);
1911 cir::FuncOp regFunc =
1912 buildRuntimeFunction(builder, regFuncName, loc, regFuncType);
1915 cir::FuncOp moduleCtor = buildRuntimeFunction(
1916 builder, moduleCtorName, loc, FuncType::get({}, voidTy),
1917 GlobalLinkageKind::InternalLinkage);
1919 globalCtorList.emplace_back(moduleCtorName,
1920 cir::GlobalCtorAttr::getDefaultPriority());
1921 builder.setInsertionPointToStart(moduleCtor.addEntryBlock());
1924 llvm_unreachable(
"HIP Module Constructor Support");
1925 }
else if (!astCtx->
getLangOpts().GPURelocatableDeviceCode) {
1933 mlir::Value fatbinVoidPtr = builder.
createBitcast(wrapper, voidPtrTy);
1934 cir::CallOp gpuBinaryHandleCall =
1936 mlir::Value gpuBinaryHandle = gpuBinaryHandleCall.getResult();
1938 mlir::Value gpuBinaryHandleGlobal = builder.
createGetGlobal(gpuBinHandle);
1939 builder.
createStore(loc, gpuBinaryHandle, gpuBinaryHandleGlobal);
1942 if (std::optional<FuncOp> regGlobal = buildCUDARegisterGlobals()) {
1943 builder.
createCallOp(loc, *regGlobal, gpuBinaryHandle);
1952 cir::CIRBaseBuilderTy globalBuilder(getContext());
1953 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
1955 buildRuntimeFunction(globalBuilder,
"__cudaRegisterFatBinaryEnd", loc,
1956 FuncType::get({voidPtrPtrTy}, voidTy));
1960 llvm_unreachable(
"GPU RDC NYI");
1965 if (std::optional<FuncOp> dtor = buildCUDAModuleDtor()) {
1968 cir::CIRBaseBuilderTy globalBuilder(getContext());
1969 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
1970 FuncOp atexit = buildRuntimeFunction(
1971 globalBuilder,
"atexit", loc,
1972 FuncType::get(PointerType::get(dtor->getFunctionType()), intTy));
1973 mlir::Value dtorFunc = GetGlobalOp::create(
1974 builder, loc, PointerType::get(dtor->getFunctionType()),
1975 mlir::FlatSymbolRefAttr::get(dtor->getSymNameAttr()));
1978 cir::ReturnOp::create(builder, loc);
1981std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() {
1982 if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName()))
1987 VoidType voidTy = VoidType::get(&getContext());
1988 PointerType voidPtrPtrTy = PointerType::get(PointerType::get(voidTy));
1990 mlir::Location loc = mlirModule.getLoc();
1992 cir::CIRBaseBuilderTy builder(getContext());
1993 builder.setInsertionPointToStart(mlirModule.getBody());
1996 std::string unregisterFuncName =
1998 FuncOp unregisterFunc = buildRuntimeFunction(
1999 builder, unregisterFuncName, loc, FuncType::get({voidPtrPtrTy}, voidTy));
2008 buildRuntimeFunction(builder, dtorName, loc, FuncType::get({}, voidTy),
2009 GlobalLinkageKind::InternalLinkage);
2011 builder.setInsertionPointToStart(dtor.addEntryBlock());
2017 GlobalOp gpubinGlobal = cast<GlobalOp>(mlirModule.lookupSymbol(gpubinName));
2019 mlir::Value gpubin = builder.
createLoad(loc, gpubinAddress);
2021 ReturnOp::create(builder, loc);
2026std::optional<FuncOp> LoweringPreparePass::buildCUDARegisterGlobals() {
2028 if (cudaKernelMap.empty())
2031 cir::CIRBaseBuilderTy builder(getContext());
2032 builder.setInsertionPointToStart(mlirModule.getBody());
2034 mlir::Location loc = mlirModule.getLoc();
2037 auto voidTy = VoidType::get(&getContext());
2038 auto voidPtrTy = PointerType::get(voidTy);
2039 auto voidPtrPtrTy = PointerType::get(voidPtrTy);
2043 std::string regGlobalFuncName =
2045 auto regGlobalFuncTy = FuncType::get({voidPtrPtrTy}, voidTy);
2046 FuncOp regGlobalFunc =
2047 buildRuntimeFunction(builder, regGlobalFuncName, loc, regGlobalFuncTy,
2048 GlobalLinkageKind::InternalLinkage);
2049 builder.setInsertionPointToStart(regGlobalFunc.addEntryBlock());
2051 buildCUDARegisterGlobalFunctions(builder, regGlobalFunc);
2055 ReturnOp::create(builder, loc);
2056 return regGlobalFunc;
2059void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
2060 cir::CIRBaseBuilderTy &builder, FuncOp regGlobalFunc) {
2061 mlir::Location loc = mlirModule.getLoc();
2063 cir::CIRDataLayout dataLayout(mlirModule);
2065 auto voidTy = VoidType::get(&getContext());
2066 auto voidPtrTy = PointerType::get(voidTy);
2067 auto voidPtrPtrTy = PointerType::get(voidPtrTy);
2069 IntType charTy = cir::IntType::get(&getContext(), astCtx->
getCharWidth(),
2073 mlir::Value fatbinHandle = *regGlobalFunc.args_begin();
2075 cir::CIRBaseBuilderTy globalBuilder(getContext());
2076 globalBuilder.setInsertionPointToStart(mlirModule.getBody());
2090 FuncOp cudaRegisterFunction = buildRuntimeFunction(
2092 FuncType::get({voidPtrPtrTy, voidPtrTy, voidPtrTy, voidPtrTy, intTy,
2093 voidPtrTy, voidPtrTy, voidPtrTy, voidPtrTy, voidPtrTy},
2096 auto makeConstantString = [&](llvm::StringRef str) -> GlobalOp {
2097 auto strType = ArrayType::get(&getContext(), charTy, 1 + str.size());
2098 auto tmpString = cir::GlobalOp::create(
2099 globalBuilder, loc, (
".str" + str).str(), strType,
2101 cir::GlobalLinkageKind::PrivateLinkage);
2104 tmpString.setInitialValueAttr(ConstArrayAttr::get(
2105 strType, StringAttr::get(&getContext(), str +
"\0")));
2106 tmpString.setPrivate();
2110 cir::ConstantOp cirNullPtr = builder.
getNullPtr(voidPtrTy, loc);
2112 for (
auto kernelName : cudaKernelMap.keys()) {
2113 FuncOp deviceStub = cudaKernelMap[kernelName];
2114 GlobalOp deviceFuncStr = makeConstantString(kernelName);
2119 llvm_unreachable(
"HIP kernel registration NYI");
2122 GetGlobalOp::create(
2123 builder, loc, PointerType::get(deviceStub.getFunctionType()),
2124 mlir::FlatSymbolRefAttr::get(deviceStub.getSymNameAttr())),
2127 loc, cudaRegisterFunction,
2128 {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
2129 ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)),
2130 cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
2135void LoweringPreparePass::runOnOperation() {
2136 mlir::Operation *op = getOperation();
2137 if (isa<::mlir::ModuleOp>(op))
2138 mlirModule = cast<::mlir::ModuleOp>(op);
2140 llvm::SmallVector<mlir::Operation *> opsToTransform;
2142 op->walk([&](mlir::Operation *op) {
2143 if (mlir::isa<cir::ArrayCtor, cir::ArrayDtor, cir::CastOp,
2144 cir::ComplexMulOp, cir::ComplexDivOp, cir::DynamicCastOp,
2145 cir::FuncOp, cir::CallOp, cir::GetGlobalOp, cir::GlobalOp,
2146 cir::StoreOp, cir::CmpThreeWayOp, cir::IncOp, cir::DecOp,
2147 cir::MinusOp, cir::NotOp>(op))
2148 opsToTransform.push_back(op);
2151 for (mlir::Operation *o : opsToTransform)
2154 buildCXXGlobalInitFunc();
2156 buildCUDAModuleCtor();
2158 buildGlobalCtorDtorList();
2162 return std::make_unique<LoweringPreparePass>();
2165std::unique_ptr<Pass>
2167 auto pass = std::make_unique<LoweringPreparePass>();
2168 pass->setASTContext(astCtx);
2169 return std::move(pass);
Defines the clang::ASTContext interface.
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)
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.
mlir::Value getSignedInt(mlir::Location loc, int64_t val, unsigned numBits)
mlir::Value createBitcast(mlir::Value src, mlir::Type newTy)
mlir::Value createGetGlobal(mlir::Location loc, cir::GlobalOp global, bool threadLocal=false)
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...
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()