https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/225224
>From dc60329cd51bf9248e3018048637391713c149c0 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Mon, 21 Sep 2026 15:37:43 -0700 Subject: [PATCH 1/2] [CIR] Consume serialized LangOptions in post-CIRGen lowering Post-CIRGen lowering (LoweringPrepare, CallConvLowering) read a handful of LangOptions facts from a live clang::LangOptions via the pass's ASTContext. That prevented a reloaded .cir from lowering the same way it was compiled, since a serialized module has no ASTContext. PR #224757 serialized those facts onto the module as #cir.lowering_lang_options; this change makes lowering consume them from there. To make sure the `lowering_lang_options` attribute is always available, I moved the constrction of langOpts from release to the constructor. The overall approach is very close to `LowerModule::getTarget()` Currently there is still a reliance to `astContext` which I plan to remove in upcoming PRs. The reliance blocks consuming .cir as an input and test `cir-opt` with some of such passes. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 39 ++++++------ .../Dialect/Transforms/LoweringPrepare.cpp | 60 +++++++++++-------- .../Transforms/TargetLowering/LowerModule.cpp | 29 +++++++-- .../Transforms/TargetLowering/LowerModule.h | 7 +++ clang/lib/CIR/Lowering/CIRPasses.cpp | 29 ++++++--- .../CIR/CodeGenCUDA/lowering-lang-options.cu | 59 ++++++++++++++++++ 6 files changed, 168 insertions(+), 55 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/lowering-lang-options.cu diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 2f032fe39e66e..df71f5a1e45ec 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -155,6 +155,27 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, theModule->setAttr(cir::CIRDialect::getIntTypeWidthAttrName(), builder.getI32IntegerAttr(target.getIntWidth())); + // Serialize the lowering-relevant LangOptions onto the ModuleOp so a reloaded + // .cir is self-describing and lowers the same way it was compiled, without a + // live clang::LangOptions. Set here (like the triple) rather than in + // release() so it is present even when codegen bails on an error, keeping it + // a hard invariant that post-CIRGen lowering can rely on. See + // #cir.lowering_lang_options. + theModule->setAttr( + cir::CIRDialect::getLoweringLangOptionsAttrName(), + cir::LoweringLangOptionsAttr::get( + &mlirContext, + /*exceptions=*/langOpts.Exceptions, + /*threadsafe_statics=*/langOpts.ThreadsafeStatics, + /*cuda=*/langOpts.CUDA, + /*cuda_is_device=*/langOpts.CUDAIsDevice, + /*hip=*/langOpts.HIP, + /*gpu_rdc=*/langOpts.GPURelocatableDeviceCode, + /*openmp=*/langOpts.OpenMP != 0, + /*openmp_is_target_device=*/langOpts.OpenMPIsTargetDevice, + /*clang_abi_compat=*/ + static_cast<int32_t>(langOpts.getClangABICompat()))); + if (cgo.OptimizationLevel > 0 || cgo.OptimizeSize > 0) theModule->setAttr(cir::CIRDialect::getOptInfoAttrName(), cir::OptInfoAttr::get(&mlirContext, @@ -3973,24 +3994,6 @@ void CIRGenModule::release() { } } - // Serialize the lowering-relevant LangOptions onto the ModuleOp, - // unconditionally, so a reloaded .cir module is self-describing. See - // #cir.lowering_lang_options. - theModule->setAttr( - cir::CIRDialect::getLoweringLangOptionsAttrName(), - cir::LoweringLangOptionsAttr::get( - &getMLIRContext(), - /*exceptions=*/langOpts.Exceptions, - /*threadsafe_statics=*/langOpts.ThreadsafeStatics, - /*cuda=*/langOpts.CUDA, - /*cuda_is_device=*/langOpts.CUDAIsDevice, - /*hip=*/langOpts.HIP, - /*gpu_rdc=*/langOpts.GPURelocatableDeviceCode, - /*openmp=*/langOpts.OpenMP != 0, - /*openmp_is_target_device=*/langOpts.OpenMPIsTargetDevice, - /*clang_abi_compat=*/ - static_cast<int32_t>(langOpts.getClangABICompat()))); - // Classic codegen calls `checkAliases` here to validate any alias // definitions emitted during codegen. assert(!cir::MissingFeatures::checkAliases()); diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index a155bd4661ce3..50b52ee9785bb 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -297,7 +297,7 @@ struct LoweringPreparePass /// AST related /// ----------- - clang::ASTContext *astCtx; + clang::ASTContext *astCtx = nullptr; /// Target/ABI facts sourced from the module's own attributes. std::unique_ptr<cir::LowerModule> lowerModule; @@ -307,6 +307,17 @@ struct LoweringPreparePass return lowerModule->getTarget(); } + /// LangOptions facts consumed by lowering, sourced from the module's + /// serialized #cir.lowering_lang_options (via LowerModule) so lowering does + /// not depend on a live clang::LangOptions and a reloaded .cir lowers the + /// same way it was compiled. CIRGen sets that attribute at module + /// construction (like the triple), so it is always present here; this + /// mirrors getTargetInfo, which likewise reads only from LowerModule. + const clang::LangOptions &getLangOpts() const { + assert(lowerModule && "LoweringPrepare requires a module with a triple"); + return lowerModule->getLangOpts(); + } + /// Tracks current module. mlir::ModuleOp mlirModule; @@ -506,7 +517,7 @@ struct LoweringPreparePass // structural, so it is only worth building when there can be one. // OG: CGF.EHStack.pushCleanup<CallGuardAbort>(EHCleanup, guard); // ... CGF.PopCleanupBlock(); - if (astCtx->getLangOpts().Exceptions) { + if (getLangOpts().Exceptions) { cir::CleanupScopeOp::create( builder, loc, cir::CleanupKind::EH, [&](mlir::OpBuilder &, mlir::Location bodyLoc) { @@ -847,7 +858,8 @@ buildRangeReductionComplexDiv(CIRBaseBuilderTy &builder, mlir::Location loc, static mlir::Type higherPrecisionElementTypeForComplexArithmetic( mlir::MLIRContext &context, clang::ASTContext &cc, - CIRBaseBuilderTy &builder, mlir::Type elementType) { + const clang::LangOptions &langOpts, CIRBaseBuilderTy &builder, + mlir::Type elementType) { auto getHigherPrecisionFPType = [&context](mlir::Type type) -> mlir::Type { if (mlir::isa<cir::FP16Type>(type)) @@ -863,7 +875,7 @@ static mlir::Type higherPrecisionElementTypeForComplexArithmetic( }; auto getFloatTypeSemantics = - [&cc](mlir::Type type) -> const llvm::fltSemantics & { + [&cc, &langOpts](mlir::Type type) -> const llvm::fltSemantics & { const clang::TargetInfo &info = cc.getTargetInfo(); if (mlir::isa<cir::FP16Type>(type)) return info.getHalfFormat(); @@ -878,13 +890,13 @@ static mlir::Type higherPrecisionElementTypeForComplexArithmetic( return info.getDoubleFormat(); if (mlir::isa<cir::LongDoubleType>(type)) { - if (cc.getLangOpts().OpenMP && cc.getLangOpts().OpenMPIsTargetDevice) + if (langOpts.OpenMP && langOpts.OpenMPIsTargetDevice) llvm_unreachable("NYI Float type semantics with OpenMP"); return info.getLongDoubleFormat(); } if (mlir::isa<cir::FP128Type>(type)) { - if (cc.getLangOpts().OpenMP && cc.getLangOpts().OpenMPIsTargetDevice) + if (langOpts.OpenMP && langOpts.OpenMPIsTargetDevice) llvm_unreachable("NYI Float type semantics with OpenMP"); return info.getFloat128Format(); } @@ -935,8 +947,8 @@ lowerComplexDiv(LoweringPreparePass &pass, CIRBaseBuilderTy &builder, if (range == cir::ComplexRangeKind::Promoted) { mlir::Type originalElementType = complexTy.getElementType(); mlir::Type higherPrecisionElementType = - higherPrecisionElementTypeForComplexArithmetic(mlirCx, cc, builder, - originalElementType); + higherPrecisionElementTypeForComplexArithmetic( + mlirCx, cc, pass.getLangOpts(), builder, originalElementType); if (!higherPrecisionElementType) return buildRangeReductionComplexDiv(builder, loc, lhsReal, lhsImag, @@ -1406,7 +1418,7 @@ void LoweringPreparePass::handleStaticLocal(cir::GlobalOp globalOp, // We only need to use thread-safe statics for local non-TLS variables and // inline variables; other global initialization is always single-threaded // or (through lazy dynamic loading in multiple threads) unsequenced. - bool threadsafe = astCtx->getLangOpts().ThreadsafeStatics && + bool threadsafe = getLangOpts().ThreadsafeStatics && (info.getLocal() || nonTemplateInline) && info.getTls() == cir::TLSKind::None; @@ -2445,8 +2457,8 @@ void LoweringPreparePass::runOnOp(mlir::Operation *op) { } } -static llvm::StringRef getCUDAPrefix(clang::ASTContext *astCtx) { - if (astCtx->getLangOpts().HIP) +static llvm::StringRef getCUDAPrefix(const clang::LangOptions &langOpts) { + if (langOpts.HIP) return "hip"; return "cuda"; } @@ -2476,9 +2488,9 @@ static std::string addUnderscoredPrefix(llvm::StringRef prefix, /// } /// \endcode void LoweringPreparePass::buildCUDAModuleCtor() { - bool isHIP = astCtx->getLangOpts().HIP; + bool isHIP = getLangOpts().HIP; - if (astCtx->getLangOpts().GPURelocatableDeviceCode) + if (getLangOpts().GPURelocatableDeviceCode) llvm_unreachable("GPU RDC NYI"); // For CUDA without -fgpu-rdc, it's safe to stop generating ctor @@ -2514,7 +2526,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { std::move(gpuBinaryOrErr.get()); // Set up common types and builder. - llvm::StringRef cudaPrefix = getCUDAPrefix(astCtx); + llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); mlir::Location loc = mlirModule->getLoc(); CIRBaseBuilderTy builder(getContext()); builder.setInsertionPointToStart(mlirModule.getBody()); @@ -2530,10 +2542,10 @@ void LoweringPreparePass::buildCUDAModuleCtor() { // The section names are different for MAC OS X. llvm::StringRef fatbinConstName = - astCtx->getLangOpts().HIP ? ".hip_fatbin" : ".nv_fatbin"; + getLangOpts().HIP ? ".hip_fatbin" : ".nv_fatbin"; llvm::StringRef fatbinSectionName = - astCtx->getLangOpts().HIP ? ".hipFatBinSegment" : ".nvFatBinSegment"; + getLangOpts().HIP ? ".hipFatBinSegment" : ".nvFatBinSegment"; // Create the fatbin string constant with GPU binary contents. auto fatbinType = @@ -2671,7 +2683,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { } return; } - if (!astCtx->getLangOpts().GPURelocatableDeviceCode) { + if (!getLangOpts().GPURelocatableDeviceCode) { // --- Create CUDA CTOR-DTOR --- // Register binary with CUDA runtime. This is substantially different in @@ -2731,7 +2743,7 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() { if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName())) return {}; - llvm::StringRef prefix = getCUDAPrefix(astCtx); + llvm::StringRef prefix = getCUDAPrefix(getLangOpts()); VoidType voidTy = VoidType::get(&getContext()); PointerType voidPtrPtrTy = PointerType::get(PointerType::get(voidTy)); @@ -2788,7 +2800,7 @@ std::optional<FuncOp> LoweringPreparePass::buildHIPModuleDtor() { if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName())) return {}; - llvm::StringRef prefix = getCUDAPrefix(astCtx); + llvm::StringRef prefix = getCUDAPrefix(getLangOpts()); VoidType voidTy = VoidType::get(&getContext()); PointerType voidPtrPtrTy = PointerType::get(PointerType::get(voidTy)); @@ -2851,7 +2863,7 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDARegisterGlobals() { builder.setInsertionPointToStart(mlirModule.getBody()); mlir::Location loc = mlirModule.getLoc(); - llvm::StringRef cudaPrefix = getCUDAPrefix(astCtx); + llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); auto voidTy = VoidType::get(&getContext()); auto voidPtrTy = PointerType::get(voidTy); @@ -2877,7 +2889,7 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDARegisterGlobals() { void LoweringPreparePass::buildCUDARegisterGlobalFunctions( cir::CIRBaseBuilderTy &builder, FuncOp regGlobalFunc) { mlir::Location loc = mlirModule.getLoc(); - llvm::StringRef cudaPrefix = getCUDAPrefix(astCtx); + llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); cir::CIRDataLayout dataLayout(mlirModule); auto voidTy = VoidType::get(&getContext()); @@ -2926,7 +2938,7 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions( }; cir::ConstantOp cirNullPtr = builder.getNullPtr(voidPtrTy, loc); - bool isHIP = astCtx->getLangOpts().HIP; + bool isHIP = getLangOpts().HIP; for (auto kernelName : cudaKernelMap.keys()) { FuncOp deviceStub = cudaKernelMap[kernelName]; GlobalOp deviceFuncStr = makeConstantString(kernelName); @@ -2964,7 +2976,7 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions( void LoweringPreparePass::buildCUDARegisterVars(cir::CIRBaseBuilderTy &builder, FuncOp regGlobalFunc) { mlir::Location loc = mlirModule.getLoc(); - llvm::StringRef cudaPrefix = getCUDAPrefix(astCtx); + llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); cir::CIRDataLayout dataLayout(mlirModule); PointerType voidPtrTy = builder.getVoidPtrTy(); @@ -3064,7 +3076,7 @@ void LoweringPreparePass::runOnOperation() { buildCXXGlobalInitFunc(); buildCXXGlobalTlsFunc(); - if (astCtx->getLangOpts().CUDA && !astCtx->getLangOpts().CUDAIsDevice) + if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice) buildCUDAModuleCtor(); buildGlobalCtorDtorList(); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 8ed0fb5b2cd26..6a15278a14b3d 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -69,7 +69,8 @@ LowerModule::LowerModule(clang::LangOptions langOpts, clang::CodeGenOptions codeGenOpts, mlir::ModuleOp &module, std::unique_ptr<clang::TargetInfo> target) - : module(module), target(std::move(target)), abi(createCXXABI(*this)) {} + : module(module), langOpts(std::move(langOpts)), target(std::move(target)), + abi(createCXXABI(*this)) {} const TargetLoweringInfo &LowerModule::getTargetLoweringInfo() { if (!targetLoweringInfo) @@ -93,11 +94,29 @@ std::unique_ptr<LowerModule> createLowerModule(mlir::ModuleOp module) { targetOptions.Triple = triple.str(); auto targetInfo = clang::targets::AllocateTarget(triple, targetOptions); - // FIXME(cir): This just uses the default language options. We need to account - // for custom options. - // Create context. - assert(!cir::MissingFeatures::lowerModuleLangOpts()); + // Populate the lowering-relevant LangOptions from the module's + // #cir.lowering_lang_options attribute so a reloaded .cir lowers the same + // way it was compiled, without a live clang::LangOptions. When the attribute + // is absent (e.g. hand-written CIR) the defaults are kept; the follow-up that + // enables .cir as a cc1 input adds the create-vs-load consistency diagnostic. + // Other LangOptions members remain unpopulated (see getCXXABIKind, which + // still carries the lowerModuleLangOpts marker for that residual gap). clang::LangOptions langOpts; + if (auto loweringLangOpts = + mlir::dyn_cast_if_present<cir::LoweringLangOptionsAttr>( + module->getAttr( + cir::CIRDialect::getLoweringLangOptionsAttrName()))) { + langOpts.Exceptions = loweringLangOpts.getExceptions(); + langOpts.ThreadsafeStatics = loweringLangOpts.getThreadsafeStatics(); + langOpts.CUDA = loweringLangOpts.getCuda(); + langOpts.CUDAIsDevice = loweringLangOpts.getCudaIsDevice(); + langOpts.HIP = loweringLangOpts.getHip(); + langOpts.GPURelocatableDeviceCode = loweringLangOpts.getGpuRdc(); + langOpts.OpenMP = loweringLangOpts.getOpenmp(); + langOpts.OpenMPIsTargetDevice = loweringLangOpts.getOpenmpIsTargetDevice(); + langOpts.setClangABICompat(static_cast<clang::LangOptions::ClangABI>( + loweringLangOpts.getClangAbiCompat())); + } // FIXME(cir): This just uses the default code generation options. We need to // account for custom options. diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.h index ab3a648683279..56fcd0a9b58cd 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.h @@ -28,6 +28,7 @@ namespace cir { class LowerModule { mlir::ModuleOp module; + const clang::LangOptions langOpts; const std::unique_ptr<clang::TargetInfo> target; std::unique_ptr<TargetLoweringInfo> targetLoweringInfo; std::unique_ptr<CIRCXXABI> abi; @@ -38,6 +39,12 @@ class LowerModule { std::unique_ptr<clang::TargetInfo> target); ~LowerModule() = default; + // The lowering-relevant LangOptions, populated by createLowerModule() from + // the module's #cir.lowering_lang_options attribute when present (so a + // reloaded .cir lowers the same way without a live clang::LangOptions), + // otherwise left at their defaults. + const clang::LangOptions &getLangOpts() const { return langOpts; } + clang::TargetCXXABI::Kind getCXXABIKind() const { assert(!cir::MissingFeatures::lowerModuleLangOpts()); return target->getCXXABI().getKind(); diff --git a/clang/lib/CIR/Lowering/CIRPasses.cpp b/clang/lib/CIR/Lowering/CIRPasses.cpp index 9b5a17dec67af..b612b64b71476 100644 --- a/clang/lib/CIR/Lowering/CIRPasses.cpp +++ b/clang/lib/CIR/Lowering/CIRPasses.cpp @@ -15,6 +15,7 @@ #include "clang/AST/ASTContext.h" #include "clang/Basic/LangOptions.h" #include "clang/Basic/TargetInfo.h" +#include "clang/CIR/Dialect/IR/CIRDialect.h" #include "clang/CIR/Dialect/Passes.h" #include "llvm/Support/TimeProfiler.h" #include "llvm/TargetParser/Triple.h" @@ -43,10 +44,10 @@ static llvm::abi::X86AVXABILevel getX86AVXABILevel(llvm::StringRef abi) { /// Whether `__attribute__((target(...)))` on a function may raise its AVX ABI /// level above the command line's. A target that opts out, and any ABI older /// than the rule, stay at the module level. -static bool allowsX86TargetAttrAvx(const clang::ASTContext &astContext) { +static bool allowsX86TargetAttrAvx(const clang::ASTContext &astContext, + clang::LangOptions::ClangABI compat) { return !astContext.getTargetInfo().getTriple().isPS() && - astContext.getLangOpts().getClangABICompat() > - clang::LangOptions::ClangABI::Ver23; + compat > clang::LangOptions::ClangABI::Ver23; } /// The x86_64 ABI-compatibility flags, derived from the target and the @@ -55,10 +56,9 @@ static bool allowsX86TargetAttrAvx(const clang::ASTContext &astContext) { /// modern Linux target, so leaving it at the default classifies a union larger /// than an eightbyte as though every member spanned its size. static llvm::abi::X86ABICompatInfo -getX86ABICompatInfo(const clang::ASTContext &astContext) { +getX86ABICompatInfo(const clang::ASTContext &astContext, + clang::LangOptions::ClangABI compat) { const llvm::Triple &triple = astContext.getTargetInfo().getTriple(); - const clang::LangOptions &langOpts = astContext.getLangOpts(); - clang::LangOptions::ClangABI compat = langOpts.getClangABICompat(); llvm::abi::X86ABICompatInfo abiCompat; abiCompat.HonorsRevision98 = !triple.isOSDarwin(); abiCompat.ClassifyIntegerMMXAsSSE = @@ -120,10 +120,23 @@ runCIRToCIRPasses(mlir::ModuleOp theModule, mlir::MLIRContext &mlirContext, // is implemented; other targets are left unchanged. const clang::TargetInfo &targetInfo = astContext.getTargetInfo(); CallConvTarget target = getCallConvTarget(targetInfo.getTriple()); - if (target != CallConvTarget::None) + if (target != CallConvTarget::None) { + // Source the ABI-compatibility version from the module's serialized + // #cir.lowering_lang_options so a reloaded .cir classifies the same way + // it was compiled, without a live clang::LangOptions. CIRGen sets this + // attribute at module construction; if it is absent fall back to the + // LangOptions default, matching how LowerModule reads the same attribute. + auto compat = clang::LangOptions::ClangABI::Latest; + if (auto loweringLangOpts = + theModule->getAttrOfType<cir::LoweringLangOptionsAttr>( + cir::CIRDialect::getLoweringLangOptionsAttrName())) + compat = static_cast<clang::LangOptions::ClangABI>( + loweringLangOpts.getClangAbiCompat()); pm.addPass(mlir::createCallConvLoweringPass( target, getX86AVXABILevel(targetInfo.getABI()), - allowsX86TargetAttrAvx(astContext), getX86ABICompatInfo(astContext))); + allowsX86TargetAttrAvx(astContext, compat), + getX86ABICompatInfo(astContext, compat))); + } } pm.enableVerifier(enableVerifier); diff --git a/clang/test/CIR/CodeGenCUDA/lowering-lang-options.cu b/clang/test/CIR/CodeGenCUDA/lowering-lang-options.cu new file mode 100644 index 0000000000000..7b3f23853244c --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/lowering-lang-options.cu @@ -0,0 +1,59 @@ +// RUN: echo -n "GPU binary would be here." > %t.fatbin + +// The lowering-relevant LangOptions are serialized onto the module as +// #cir.lowering_lang_options (see clang/test/CIR/CodeGen/lowering-lang-options.cpp), +// and post-CIRGen lowering consumes them from there via LowerModule rather than +// from a live clang::LangOptions. This test exercises both halves end-to-end: +// CIRGen records the CUDA host/device configuration in the attribute, and +// LoweringPrepare gates CUDA module-ctor synthesis on it +// (cuda && !cuda_is_device), so the registration ctor appears for the host +// compile and is suppressed for the device compile. + +//===----------------------------------------------------------------------===// +// Host compilation: cuda = true, cuda_is_device = false. +//===----------------------------------------------------------------------===// + +// The serialized attribute records the host configuration: +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-cir %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=HOST-ATTR +// HOST-ATTR: cir.lowering_lang_options = #cir.lowering_lang_options< +// HOST-ATTR-SAME: cuda = true +// HOST-ATTR-SAME: cuda_is_device = false + +// -emit-cir stops before the passes, so the ctor only appears once lowering has +// consumed the attribute. cuda && !cuda_is_device holds, so it is synthesized: +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=HOST-LOWER +// HOST-LOWER: __cuda_module_ctor +// HOST-LOWER: __cudaRegisterFatBinary + +//===----------------------------------------------------------------------===// +// Device compilation: cuda = true, cuda_is_device = true. +//===----------------------------------------------------------------------===// + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -emit-cir %s -x cuda \ +// RUN: -fcuda-is-device -target-sdk-version=12.3 -o - \ +// RUN: | FileCheck %s --check-prefix=DEV-ATTR +// DEV-ATTR: cir.lowering_lang_options = #cir.lowering_lang_options< +// DEV-ATTR-SAME: cuda = true +// DEV-ATTR-SAME: cuda_is_device = true + +// cuda_is_device flips the gate, so lowering must NOT synthesize the host-only +// CUDA module ctor: +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -emit-llvm %s -x cuda \ +// RUN: -fcuda-is-device -target-sdk-version=12.3 -o - \ +// RUN: | FileCheck %s --check-prefix=DEV-LOWER +// DEV-LOWER: @_Z6kernelv +// DEV-LOWER-NOT: __cuda_module_ctor + +// Minimal CUDA runtime declarations so the host-side launch stub can be built +// without the real CUDA headers. +typedef unsigned long size_t; +struct dim3 { unsigned x, y, z; }; +extern "C" int cudaLaunchKernel(const void *, dim3, dim3, void **, size_t, + void *); +extern "C" int __cudaPopCallConfiguration(dim3 *, dim3 *, size_t *, void *); + +__attribute__((global)) void kernel() {} >From 10d88ef90438eef2f7d6c8532d3b4083adde9e1f Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Mon, 21 Sep 2026 15:48:02 -0700 Subject: [PATCH 2/2] Trim comment line --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 5 +---- clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp | 8 +------- .../CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp | 5 +---- 3 files changed, 3 insertions(+), 15 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index df71f5a1e45ec..10a2019941292 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -157,10 +157,7 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, // Serialize the lowering-relevant LangOptions onto the ModuleOp so a reloaded // .cir is self-describing and lowers the same way it was compiled, without a - // live clang::LangOptions. Set here (like the triple) rather than in - // release() so it is present even when codegen bails on an error, keeping it - // a hard invariant that post-CIRGen lowering can rely on. See - // #cir.lowering_lang_options. + // live clang::LangOptions. theModule->setAttr( cir::CIRDialect::getLoweringLangOptionsAttrName(), cir::LoweringLangOptionsAttr::get( diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index 50b52ee9785bb..d1a7c406c62d7 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -307,14 +307,8 @@ struct LoweringPreparePass return lowerModule->getTarget(); } - /// LangOptions facts consumed by lowering, sourced from the module's - /// serialized #cir.lowering_lang_options (via LowerModule) so lowering does - /// not depend on a live clang::LangOptions and a reloaded .cir lowers the - /// same way it was compiled. CIRGen sets that attribute at module - /// construction (like the triple), so it is always present here; this - /// mirrors getTargetInfo, which likewise reads only from LowerModule. const clang::LangOptions &getLangOpts() const { - assert(lowerModule && "LoweringPrepare requires a module with a triple"); + assert(lowerModule && "LoweringPrepare requires a module with LangOptions"); return lowerModule->getLangOpts(); } diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 6a15278a14b3d..04afa86844034 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -97,10 +97,7 @@ std::unique_ptr<LowerModule> createLowerModule(mlir::ModuleOp module) { // Populate the lowering-relevant LangOptions from the module's // #cir.lowering_lang_options attribute so a reloaded .cir lowers the same // way it was compiled, without a live clang::LangOptions. When the attribute - // is absent (e.g. hand-written CIR) the defaults are kept; the follow-up that - // enables .cir as a cc1 input adds the create-vs-load consistency diagnostic. - // Other LangOptions members remain unpopulated (see getCXXABIKind, which - // still carries the lowerModuleLangOpts marker for that residual gap). + // is absent (e.g. hand-written CIR) the defaults are kept; clang::LangOptions langOpts; if (auto loweringLangOpts = mlir::dyn_cast_if_present<cir::LoweringLangOptionsAttr>( _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
