[clang] [CIR] Consume serialized LangOptions in post-CIRGen lowering (PR #225224)
Konstantinos Parasyris via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 21 15:48:49 PDT 2026
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 <konstantinos.parasyris at intel.com>
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) <noreply at anthropic.com>
---
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 <konstantinos.parasyris at intel.com>
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>(
More information about the cfe-commits
mailing list