[clang] [llvm] [CIR][CUDA] Shadow variables, internalize device side variables, lower poison attribute (PR #190087)
Zaky Hermawan via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 30 22:31:35 PDT 2026
https://github.com/ZakyHermawan updated https://github.com/llvm/llvm-project/pull/190087
>From a30d3830c28be0ee50f76c836c06393709dfb8c4 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Thu, 2 Apr 2026 06:49:09 +0700
Subject: [PATCH 01/12] [CIR][CUDA] Shadow variables, lower poison attribute,
improve test readability
Poison PoisonAtt already been introduced, but no lowering exist, this commit address that,
Improve readability of checks in address-spaces.cu by removing information that we did not care,
Note:
- CIR->LLVM initialize global variables with poison,
while OGCG initialize global variables with undef.
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
.../clang/CIR/Dialect/IR/CIRCUDAAttrs.td | 18 ++++++
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 60 ++++++++++++++++++
clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp | 19 ++++++
clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h | 13 +++-
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 49 ++++++++++++++-
clang/lib/CIR/CodeGen/CIRGenModule.h | 7 +++
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 24 ++++---
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 63 ++++++++++++++-----
8 files changed, 225 insertions(+), 28 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
index 5932db8323196..d68fb61fd115c 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
@@ -37,6 +37,24 @@ def CIR_CUDAKernelNameAttr : CIR_Attr<"CUDAKernelName", "cu.kernel_name"> {
let canHaveIllegalCXXABIType = 0;
}
+def CUDAShadowNameAttr : CIR_Attr<"CUDAShadowName",
+ "cu.shadow_name"> {
+ let summary = "Device-side global variable name for this shadow.";
+ let description =
+ [{
+ This attribute is attached to global variable definitions and records the
+ mangled name of the global variable used on the device.
+
+ In CUDA, __device__, __constant__ and __shared__ variables, as well as
+ surface and texture variables, will generate a shadow symbol on host.
+ We must preserve the correspodence in order to generate registration
+ functions.
+ }];
+
+ let parameters = (ins "std::string":$device_side_name);
+ let assemblyFormat = "`<` $device_side_name `>`";
+}
+
def CUDAExternallyInitializedAttr : CIR_Attr<"CUDAExternallyInitialized",
"cu.externally_initialized"> {
let summary = "The marked variable is externally initialized.";
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index 8b8e99023eceb..8fb7191e7a89e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -12,6 +12,7 @@
//===----------------------------------------------------------------------===//
#include "CIRGenCUDARuntime.h"
+#include "CIRGenCXXABI.h"
#include "CIRGenFunction.h"
#include "CIRGenModule.h"
#include "mlir/IR/Operation.h"
@@ -64,6 +65,11 @@ class CIRGenNVCUDARuntime : public CIRGenCUDARuntime {
void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
FunctionArgList &args) override;
+
+ void internalizeDeviceSideVar(const VarDecl *d,
+ cir::GlobalLinkageKind &linkage) override;
+
+ std::string getDeviceSideName(const NamedDecl *nd) override;
};
} // namespace
@@ -342,3 +348,57 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
return globalOp;
}
+
+void CIRGenNVCUDARuntime::internalizeDeviceSideVar(
+ const VarDecl *d, cir::GlobalLinkageKind &linkage) {
+ if (cgm.getLangOpts().GPURelocatableDeviceCode)
+ cgm.errorNYI(
+ "internalizeDeviceSideVar: GPU Relocatable Deviced Code (RDC)");
+
+ // __shared__ variables are odd. Shadows do get created, but
+ // they are not registered with the CUDA runtime, so they
+ // can't really be used to access their device-side
+ // counterparts. It's not clear yet whether it's nvcc's bug or
+ // a feature, but we've got to do the same for compatibility.
+ if (d->hasAttr<CUDADeviceAttr>() || d->hasAttr<CUDAConstantAttr>() ||
+ d->hasAttr<CUDASharedAttr>()) {
+ linkage = cir::GlobalLinkageKind::InternalLinkage;
+ }
+
+ if (d->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ d->getType()->isCUDADeviceBuiltinTextureType())
+ cgm.errorNYI("internalizeDeviceSideVar: CUDA Surface/Texture support");
+}
+
+std::string CIRGenNVCUDARuntime::getDeviceSideName(const NamedDecl *nd) {
+ GlobalDecl gd;
+ // nd could be either a kernel or a variable.
+ if (auto *fd = dyn_cast<FunctionDecl>(nd))
+ gd = GlobalDecl(fd, KernelReferenceKind::Kernel);
+ else
+ gd = GlobalDecl(nd);
+ std::string deviceSideName;
+ MangleContext *mc;
+ if (cgm.getLangOpts().CUDAIsDevice)
+ mc = &cgm.getCXXABI().getMangleContext();
+ else
+ mc = deviceMC.get();
+ if (mc->shouldMangleDeclName(nd)) {
+ SmallString<256> buffer;
+ llvm::raw_svector_ostream out(buffer);
+ mc->mangleName(gd, out);
+ deviceSideName = std::string(out.str());
+ } else
+ deviceSideName = std::string(nd->getIdentifier()->getName());
+
+ // Make unique name for device side static file-scope variable for HIP.
+ if (cgm.getASTContext().shouldExternalize(nd) &&
+ cgm.getLangOpts().GPURelocatableDeviceCode) {
+ SmallString<256> buffer;
+ llvm::raw_svector_ostream out(buffer);
+ out << deviceSideName;
+ cgm.printPostfixForExternalizedDecl(out, nd);
+ deviceSideName = std::string(out.str());
+ }
+ return deviceSideName;
+}
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp
index 25d981ef2f64b..8898071a35c12 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp
@@ -20,6 +20,25 @@
using namespace clang;
using namespace CIRGen;
+static std::unique_ptr<MangleContext> initDeviceMC(CIRGenModule &cgm) {
+ // If the host and device have different C++ ABIs, mark it as the device
+ // mangle context so that the mangling needs to retrieve the additional
+ // device lambda mangling number instead of the regular host one.
+ if (cgm.getASTContext().getAuxTargetInfo() &&
+ cgm.getASTContext().getTargetInfo().getCXXABI().isMicrosoft() &&
+ cgm.getASTContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily()) {
+ return std::unique_ptr<MangleContext>(
+ cgm.getASTContext().createDeviceMangleContext(
+ *cgm.getASTContext().getAuxTargetInfo()));
+ }
+
+ return std::unique_ptr<MangleContext>(cgm.getASTContext().createMangleContext(
+ cgm.getASTContext().getAuxTargetInfo()));
+}
+
+CIRGenCUDARuntime::CIRGenCUDARuntime(CIRGenModule &cgm)
+ : cgm(cgm), deviceMC(initDeviceMC(cgm)) {}
+
CIRGenCUDARuntime::~CIRGenCUDARuntime() {}
RValue CIRGenCUDARuntime::emitCUDAKernelCallExpr(CIRGenFunction &cgf,
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
index ba33602511e3b..39b6571849f29 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
@@ -33,8 +33,11 @@ class CIRGenCUDARuntime {
protected:
CIRGenModule &cgm;
+ /// Mangle context for device.
+ std::unique_ptr<MangleContext> deviceMC;
+
public:
- CIRGenCUDARuntime(CIRGenModule &cgm) : cgm(cgm) {}
+ CIRGenCUDARuntime(CIRGenModule &cgm);
virtual ~CIRGenCUDARuntime();
virtual void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
@@ -47,6 +50,14 @@ class CIRGenCUDARuntime {
virtual mlir::Operation *getKernelHandle(cir::FuncOp fn, GlobalDecl gd) = 0;
virtual mlir::Operation *getKernelStub(mlir::Operation *handle) = 0;
+
+ /// Adjust linkage of shadow variables in host compilation
+ virtual void internalizeDeviceSideVar(const VarDecl *d,
+ cir::GlobalLinkageKind &linkage) = 0;
+
+ /// Returns function or variable name on device side even if the current
+ /// compilation is for host.
+ virtual std::string getDeviceSideName(const NamedDecl *nd) = 0;
};
CIRGenCUDARuntime *createNVCUDARuntime(CIRGenModule &cgm);
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 2bc33c191bb32..88c4ac5d27115 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -434,6 +434,26 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
global->getType()->isCUDADeviceBuiltinTextureType();
}
+void CIRGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &os,
+ const Decl *d) {
+ // ptxas does not allow '.' in symbol names. On the other hand, HIP prefers
+ // postfix beginning with '.' since the symbol name can be demangled.
+ if (langOpts.HIP)
+ os << (isa<VarDecl>(d) ? ".static." : ".intern.");
+ else
+ os << (isa<VarDecl>(d) ? "__static__" : "__intern__");
+
+ // If the CUID is not specified we try to generate a unique postfix.
+ if (getLangOpts().CUID.empty()) {
+ // TODO: Once we add 'PreprocessorOpts' into CIRGenModule this part can be
+ // brought in from OG.
+ errorNYI(d->getSourceRange(),
+ "printPostfixForExternalizedDecl: CUID is not specified");
+ } else {
+ os << getASTContext().getCUIDHash();
+ }
+}
+
void CIRGenModule::emitGlobal(clang::GlobalDecl gd) {
if (const auto *cd = dyn_cast<clang::OpenACCConstructDecl>(gd.getDecl())) {
emitGlobalOpenACCDecl(cd);
@@ -1233,15 +1253,40 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
cir::CUDAExternallyInitializedAttr::get(&getMLIRContext()));
}
} else {
- // TODO(cir):
// Adjust linkage of shadow variables in host compilation
- // getCUDARuntime().internalizeDeviceSideVar(vd, linkage);
+ getCUDARuntime().internalizeDeviceSideVar(vd, linkage);
}
// TODO(cir):
// Handle variable registration
// getCUDARuntime().handleVarRegistration(vd, gv);
}
+ // Decorate CUDA shadow variables with the cu.shadow_name attribute so we know
+ // how to register them when lowering.
+ if (langOpts.CUDA && !langOpts.CUDAIsDevice &&
+ (vd->hasAttr<CUDAConstantAttr>() || vd->hasAttr<CUDADeviceAttr>())) {
+ // Shadow variables and their properties must be registered with CUDA
+ // runtime. Skip Extern global variables, which will be registered in
+ // the TU where they are defined.
+ //
+ // Don't register a C++17 inline variable. The local symbol can be
+ // discarded and referencing a discarded local symbol from outside the
+ // comdat (__cuda_register_globals) is disallowed by the ELF spec.
+ //
+ // HIP managed variables need to be always recorded in device and host
+ // compilations for transformation.
+ //
+ // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
+ // added to llvm.compiler-used, therefore they are safe to be registered.
+ if ((!vd->hasExternalStorage() && !vd->isInline()) ||
+ getASTContext().CUDADeviceVarODRUsedByHost.contains(vd) ||
+ vd->hasAttr<HIPManagedAttr>()) {
+ auto shadowName = cudaRuntime->getDeviceSideName(cast<NamedDecl>(vd));
+ auto attr = cir::CUDAShadowNameAttr::get(&getMLIRContext(), shadowName);
+ gv->setAttr(cir::CUDAShadowNameAttr::getMnemonic(), attr);
+ }
+ }
+
// Set initializer and finalize emission
CIRGenModule::setInitializer(gv, init);
if (emitter)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 266510de84fd0..388b78f2a75e6 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -632,6 +632,13 @@ class CIRGenModule : public CIRGenTypeCache {
// related attributes.
bool shouldEmitCUDAGlobalVar(const VarDecl *global) const;
+ /// Print the postfix for externalized static variable or kernels for single
+ /// source offloading languages CUDA and HIP. The unique postfix is created
+ /// using either the CUID argument, or the file's UniqueID and active macros.
+ /// The fallback method without a CUID requires that the offloading toolchain
+ /// does not define separate macros via the -cc1 options.
+ void printPostfixForExternalizedDecl(llvm::raw_ostream &os, const Decl *d);
+
/// Replace all uses of the old global with the new global, updating types
/// and references as needed. Erases the old global when done.
void replaceGlobal(cir::GlobalOp oldGV, cir::GlobalOp newGV);
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index ba89fbe3091bc..6881fef5d9f49 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -25,6 +25,7 @@
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinDialect.h"
#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/Location.h"
#include "mlir/IR/Types.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
@@ -390,7 +391,7 @@ class CIRAttrToValue {
.Case<cir::BoolAttr, cir::IntAttr, cir::FPAttr, cir::ConstComplexAttr,
cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr,
cir::ConstPtrAttr, cir::GlobalViewAttr, cir::TypeInfoAttr,
- cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>(
+ cir::UndefAttr, cir::PoisonAttr, cir::VTableAttr, cir::ZeroAttr>(
[&](auto attrT) { return visitCirAttr(attrT); })
.Default([&](auto attrT) { return mlir::Value(); });
}
@@ -406,6 +407,7 @@ class CIRAttrToValue {
mlir::Value visitCirAttr(cir::GlobalViewAttr attr);
mlir::Value visitCirAttr(cir::TypeInfoAttr attr);
mlir::Value visitCirAttr(cir::UndefAttr attr);
+ mlir::Value visitCirAttr(cir::PoisonAttr attr);
mlir::Value visitCirAttr(cir::VTableAttr attr);
mlir::Value visitCirAttr(cir::ZeroAttr attr);
@@ -767,6 +769,13 @@ mlir::Value CIRAttrToValue::visitCirAttr(cir::UndefAttr undefAttr) {
rewriter, loc, converter->convertType(undefAttr.getType()));
}
+/// PoisonAttr visitor.
+mlir::Value CIRAttrToValue::visitCirAttr(cir::PoisonAttr poisonAttr) {
+ mlir::Location loc = parentOp->getLoc();
+ return mlir::LLVM::PoisonOp::create(
+ rewriter, loc, converter->convertType(poisonAttr.getType()));
+}
+
// VTableAttr visitor.
mlir::Value CIRAttrToValue::visitCirAttr(cir::VTableAttr vtableArr) {
mlir::Type llvmTy = converter->convertType(vtableArr.getType());
@@ -2603,11 +2612,10 @@ CIRToLLVMGlobalOpLowering::matchAndRewriteRegionInitializedGlobal(
cir::GlobalOp op, mlir::Attribute init,
mlir::ConversionPatternRewriter &rewriter) const {
// TODO: Generalize this handling when more types are needed here.
- assert(
- (isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr,
- cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr,
- cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>(
- init)));
+ assert((isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr,
+ cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr,
+ cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr,
+ cir::VTableAttr, cir::ZeroAttr>(init)));
// TODO(cir): once LLVM's dialect has proper equivalent attributes this
// should be updated. For now, we use a custom op to initialize globals
@@ -2674,8 +2682,8 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
} else if (mlir::isa<cir::ConstArrayAttr, cir::ConstVectorAttr,
cir::ConstRecordAttr, cir::ConstPtrAttr,
cir::ConstComplexAttr, cir::GlobalViewAttr,
- cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr,
- cir::ZeroAttr>(init.value())) {
+ cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr,
+ cir::VTableAttr, cir::ZeroAttr>(init.value())) {
// TODO(cir): once LLVM's dialect has proper equivalent attributes this
// should be updated. For now, we use a custom op to initialize globals
// to the appropriate value.
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 1ed52378b99ac..65fa86ac4790c 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -15,20 +15,49 @@
// RUN: -I%S/Inputs/ %s -o %t.ll
// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s
-// CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j = #cir.undef : !s32i {alignment = 4 : i64}
-// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef, align 4
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
+// RUN: %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: %s -o %t.cir
+// RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: %s -o %t.cir
+// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.cir %s
+
+// CIR-DEVICE: cir.global {{.*}} @_ZZ2fnvE1j = #cir.undef
+// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef
__device__ int a;
-// CIR-DEVICE: cir.global external lang_address_space(offload_global) @[[DEV:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
-// LLVM-DEVICE: @[[DEV_LD:.*]] = externally_initialized global i32 0, align 4
-// OGCG-DEVICE: @[[DEV_OD:.*]] = addrspace(1) externally_initialized global i32 0, align 4
+// CIR-DEVICE: cir.global external lang_address_space(offload_global) @a = #cir.int<0> : !s32i {{{.*}}, cu.externally_initialized = #cir.cu.externally_initialized}
+// LLVM-DEVICE: @a = externally_initialized global i32 0
+// OGCG-DEVICE: @a = addrspace(1) externally_initialized global i32 0
+// CIR-HOST: cir.global {{.*}} @a = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<a>}
+// LLVM-HOST: @a = internal global i32 poison
+// OGCG-HOST: @a = internal global i32 undef
+
+__shared__ int b;
+// CIR-DEVICE: cir.global external lang_address_space(offload_local) @b = #cir.poison {{.*}}
+// LLVM-DEVICE: @b = global i32 poison
+// OGCG-DEVICE: @b = addrspace(3) global i32 undef
+// CIR-HOST: cir.global {{.*}} @b = #cir.poison
+// LLVM-HOST: @b = internal global i32 poison
+// OGCG-HOST: @b = internal global i32 undef
__constant__ int c;
-// CIR-DEVICE: cir.global constant external lang_address_space(offload_constant) @[[CONST:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
-// LLVM-DEVICE: @[[CONST_LL:.*]] = externally_initialized constant i32 0, align 4
-// OGCG-DEVICE: @[[CONST_OD:.*]] = addrspace(4) externally_initialized constant i32 0, align 4
+// CIR-DEVICE: cir.global constant external lang_address_space(offload_constant) @c = #cir.int<0> : !s32i {{{.*}}, cu.externally_initialized = #cir.cu.externally_initialized}
+// LLVM-DEVICE: @c = externally_initialized constant i32 0
+// OGCG-DEVICE: @c = addrspace(4) externally_initialized constant i32 0
+// CIR-HOST: cir.global {{.*}} @c = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<c>}
+// LLVM-HOST: @c = internal global i32 poison
+// OGCG-HOST: @c = internal global i32 undef
-// OGCG-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef, align 4
+// OGCG-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef
__global__ void fn() {
int i = 0;
@@ -46,16 +75,16 @@ __global__ void fn() {
// CIR-DEVICE: cir.return
// LLVM-DEVICE: define dso_local void @_Z2fnv()
-// LLVM-DEVICE: %[[ALLOCA:.*]] = alloca i32, i64 1, align 4
-// LLVM-DEVICE: store i32 0, ptr %[[ALLOCA]], align 4
-// LLVM-DEVICE: %[[VAL:.*]] = load i32, ptr %[[ALLOCA]], align 4
-// LLVM-DEVICE: store i32 %[[VAL]], ptr @_ZZ2fnvE1j, align 4
+// LLVM-DEVICE: %[[ALLOCA:.*]] = alloca i32, i64 1
+// LLVM-DEVICE: store i32 0, ptr %[[ALLOCA]]
+// LLVM-DEVICE: %[[VAL:.*]] = load i32, ptr %[[ALLOCA]]
+// LLVM-DEVICE: store i32 %[[VAL]], ptr @_ZZ2fnvE1j
// LLVM-DEVICE: ret void
// OGCG-DEVICE: define dso_local ptx_kernel void @_Z2fnv()
// OGCG-DEVICE: entry:
-// OGCG-DEVICE: %[[ALLOCA:.*]] = alloca i32, align 4
-// OGCG-DEVICE: store i32 0, ptr %[[ALLOCA]], align 4
-// OGCG-DEVICE: %[[VAL:.*]] = load i32, ptr %[[ALLOCA]], align 4
-// OGCG-DEVICE: store i32 %[[VAL]], ptr addrspacecast (ptr addrspace(3) @_ZZ2fnvE1j to ptr), align 4
+// OGCG-DEVICE: %[[ALLOCA:.*]] = alloca i32
+// OGCG-DEVICE: store i32 0, ptr %[[ALLOCA]]
+// OGCG-DEVICE: %[[VAL:.*]] = load i32, ptr %[[ALLOCA]]
+// OGCG-DEVICE: store i32 %[[VAL]], ptr addrspacecast (ptr addrspace(3) @_ZZ2fnvE1j to ptr)
// OGCG-DEVICE: ret void
>From af76d01740c17941d67c67b0f8b8001cf7efb682 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Thu, 2 Apr 2026 07:17:17 +0700
Subject: [PATCH 02/12] [CIR][CUDA][NFC] Add source range to NYI diagnostics
for internalizeDeviceSideVar
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index 8fb7191e7a89e..773924e6ca301 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -352,8 +352,8 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
void CIRGenNVCUDARuntime::internalizeDeviceSideVar(
const VarDecl *d, cir::GlobalLinkageKind &linkage) {
if (cgm.getLangOpts().GPURelocatableDeviceCode)
- cgm.errorNYI(
- "internalizeDeviceSideVar: GPU Relocatable Deviced Code (RDC)");
+ cgm.errorNYI(d->getSourceRange(),
+ "internalizeDeviceSideVar: GPU Relocatable Device Code (RDC)");
// __shared__ variables are odd. Shadows do get created, but
// they are not registered with the CUDA runtime, so they
@@ -367,7 +367,8 @@ void CIRGenNVCUDARuntime::internalizeDeviceSideVar(
if (d->getType()->isCUDADeviceBuiltinSurfaceType() ||
d->getType()->isCUDADeviceBuiltinTextureType())
- cgm.errorNYI("internalizeDeviceSideVar: CUDA Surface/Texture support");
+ cgm.errorNYI(d->getSourceRange(),
+ "internalizeDeviceSideVar: CUDA Surface/Texture support");
}
std::string CIRGenNVCUDARuntime::getDeviceSideName(const NamedDecl *nd) {
>From c590ae596917604ef17d50f50762ed9e36e6029b Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 3 Apr 2026 03:07:33 +0700
Subject: [PATCH 03/12] [CIR][NFC] Try resolve conflict
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 79 ++++++++++++--------
1 file changed, 46 insertions(+), 33 deletions(-)
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index b2b20c3a7771b..38ebef92f4fc4 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -10,8 +10,9 @@
// RUN: -mmlir -mlir-print-ir-before=cir-target-lowering %s -o %t.cir 2> %t-pre.cir
// RUN: FileCheck --check-prefix=CIR-PRE --input-file=%t-pre.cir %s
-// TODO: Add CIR (post target lowering) and LLVM checks once NVPTX TargetLoweringInfo
-// is implemented.
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-POST --input-file=%t.cir %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
@@ -44,44 +45,56 @@
// Verifies CIR emits correct address spaces for CUDA globals.
-// CIR-DEVICE: cir.global {{.*}} @_ZZ2fnvE1j = #cir.undef
+// CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j = #cir.undef
// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef
-__device__ int a;
-// CIR-PRE: cir.global external lang_address_space(offload_global) @a = #cir.int<0> : !s32i {{{.*}}, cu.externally_initialized = #cir.cu.externally_initialized}
-// LLVM-DEVICE: @a = externally_initialized global i32 0
-// OGCG-DAG: @a = addrspace(1) externally_initialized global i32 0
-// OGCG-DEVICE: @a = addrspace(1) externally_initialized global i32 0
-// CIR-HOST: cir.global {{.*}} @a = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<a>}
-// LLVM-HOST: @a = internal global i32 poison
-// OGCG-HOST: @a = internal global i32 undef
-
-__shared__ int b;
-// CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.poison {{.*}}
-// LLVM-DEVICE: @b = global i32 poison
-// OGCG-DEVICE: @b = addrspace(3) global i32 undef
-// CIR-HOST: cir.global {{.*}} @b = #cir.poison
-// LLVM-HOST: @b = internal global i32 poison
-// OGCG-HOST: @b = internal global i32 undef
-
-__constant__ int c;
-// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @c = #cir.int<0> : !s32i {{{.*}}, cu.externally_initialized = #cir.cu.externally_initialized}
-// LLVM-DEVICE: @c = externally_initialized constant i32 0
-// OGCG-DAG: @c = addrspace(4) externally_initialized constant i32 0
-// OGCG-DEVICE: @c = addrspace(4) externally_initialized constant i32 0
-// CIR-HOST: cir.global {{.*}} @c = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<c>}
-// LLVM-HOST: @c = internal global i32 poison
-// OGCG-HOST: @c = internal global i32 undef
+// CIR-PRE: cir.global external lang_address_space(offload_global) @i = #cir.int<0> : !s32i
+// CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0> : !s32i
+// LLVM-DEVICE-DAG: @i = addrspace(1) {{.*}}global i32 0
+// OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0
+// CIR-HOST: cir.global {{.*}} @i = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<a>}
+// LLVM-HOST: @i = internal global i32 poison
+// OGCG-HOST: @i = internal global i32 undef
+__device__ int i;
+
+// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0> : !s32i
+// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0> : !s32i
+// LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0
+// OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0
+// CIR-HOST: cir.global {{.*}} @j = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<c>}
+// LLVM-HOST: @j = internal global i32 poison
+// OGCG-HOST: @j = internal global i32 undef
+__constant__ int j;
+
+// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison : !s32i
+// CIR-POST: cir.global external target_address_space(3) @k = #cir.poison : !s32i
+// LLVM-DEVICE-DAG: @k = addrspace(3) global i32 {{undef|poison}}
+// OGCG-DAG: @k = addrspace(3) global i32 undef
+// CIR-HOST: cir.global {{.*}} @k = #cir.poison
+// LLVM-HOST: @k = internal global i32 poison
+// OGCG-HOST: @k = internal global i32 undef
+__shared__ int k;
+
+// CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.poison : !cir.float
+// CIR-POST: cir.global external target_address_space(3) @b = #cir.poison : !cir.float
+// LLVM-DEVICE-DAG: @b = addrspace(3) global float {{undef|poison}}
+// OGCG-DAG: @b = addrspace(3) global float undef
+__shared__ float b;
__device__ void foo() {
- // CIR-PRE: cir.get_global @a : !cir.ptr<!s32i, lang_address_space(offload_global)>
- a++;
+ // CIR-PRE: cir.get_global @i : !cir.ptr<!s32i, lang_address_space(offload_global)>
+ // CIR-POST: cir.get_global @i : !cir.ptr<!s32i, target_address_space(1)>
+ i++;
+
+ // CIR-PRE: cir.get_global @j : !cir.ptr<!s32i, lang_address_space(offload_constant)>
+ // CIR-POST: cir.get_global @j : !cir.ptr<!s32i, target_address_space(4)>
+ j++;
- // CIR-PRE: cir.get_global @c : !cir.ptr<!s32i, lang_address_space(offload_constant)>
- c++;
+ // CIR-PRE: cir.get_global @k : !cir.ptr<!s32i, lang_address_space(offload_local)>
+ // CIR-POST: cir.get_global @k : !cir.ptr<!s32i, target_address_space(3)>
+ k++;
}
-// OGCG-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef
__global__ void fn() {
int i = 0;
__shared__ int j;
>From 803287de08ee0ca8c2a8dc8617674f69f6cab740 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 3 Apr 2026 14:15:59 +0700
Subject: [PATCH 04/12] trigger GitHub actions
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
>From 6f9744b8c14713c1694962130db5a902b2436077 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Sat, 4 Apr 2026 18:46:40 +0700
Subject: [PATCH 05/12] trigger GitHub actions
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
>From 9f7465cc274c42d35bb5dd71942ac177159e7b7e Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Sat, 4 Apr 2026 23:19:51 +0700
Subject: [PATCH 06/12] Fix test
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 18 +++++++++---------
1 file changed, 9 insertions(+), 9 deletions(-)
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 38ebef92f4fc4..ca3cbd358dc89 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -30,7 +30,7 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
-// RUN: %s -o %t.cir
+// RUN: -I%S/Inputs/ %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
@@ -48,26 +48,26 @@
// CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j = #cir.undef
// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef
-// CIR-PRE: cir.global external lang_address_space(offload_global) @i = #cir.int<0> : !s32i
-// CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0> : !s32i
+// CIR-PRE: cir.global external lang_address_space(offload_global) @i = #cir.int<0>
+// CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0>
// LLVM-DEVICE-DAG: @i = addrspace(1) {{.*}}global i32 0
// OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0
-// CIR-HOST: cir.global {{.*}} @i = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<a>}
+// CIR-HOST: cir.global {{.*}} @i = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<i>}
// LLVM-HOST: @i = internal global i32 poison
// OGCG-HOST: @i = internal global i32 undef
__device__ int i;
-// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0> : !s32i
-// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0> : !s32i
+// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0>
+// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0>
// LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0
// OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0
-// CIR-HOST: cir.global {{.*}} @j = #cir.poison : !s32i {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<c>}
+// CIR-HOST-NEXT: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<j>}
// LLVM-HOST: @j = internal global i32 poison
// OGCG-HOST: @j = internal global i32 undef
__constant__ int j;
-// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison : !s32i
-// CIR-POST: cir.global external target_address_space(3) @k = #cir.poison : !s32i
+// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison
+// CIR-POST: cir.global external target_address_space(3) @k = #cir.poison
// LLVM-DEVICE-DAG: @k = addrspace(3) global i32 {{undef|poison}}
// OGCG-DAG: @k = addrspace(3) global i32 undef
// CIR-HOST: cir.global {{.*}} @k = #cir.poison
>From 48b09725a99774d36419dfb03fc1b79ede95b6b9 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Sun, 5 Apr 2026 06:15:10 +0700
Subject: [PATCH 07/12] remove next in check to follow the pattern
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index ca3cbd358dc89..c0f48e4c22aae 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -61,7 +61,7 @@ __device__ int i;
// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0>
// LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0
// OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0
-// CIR-HOST-NEXT: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<j>}
+// CIR-HOST: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<j>}
// LLVM-HOST: @j = internal global i32 poison
// OGCG-HOST: @j = internal global i32 undef
__constant__ int j;
>From eb7e490ffe5051ff6b5970aff832f69f3363ba3d Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 02:59:25 +0700
Subject: [PATCH 08/12] [CIR] Annotate device shadow variable using
CUDAVarRegistrationInfo
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
.../clang/CIR/Dialect/IR/CIRCUDAAttrs.td | 48 +++++++-----
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 77 +++++++++++++++++++
clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h | 7 ++
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 20 +++--
clang/lib/CIR/CodeGen/CIRGenModule.h | 6 ++
clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 66 ++++++++++++++++
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 4 +-
7 files changed, 202 insertions(+), 26 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
index d68fb61fd115c..5bd37e6471b7a 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td
@@ -37,25 +37,7 @@ def CIR_CUDAKernelNameAttr : CIR_Attr<"CUDAKernelName", "cu.kernel_name"> {
let canHaveIllegalCXXABIType = 0;
}
-def CUDAShadowNameAttr : CIR_Attr<"CUDAShadowName",
- "cu.shadow_name"> {
- let summary = "Device-side global variable name for this shadow.";
- let description =
- [{
- This attribute is attached to global variable definitions and records the
- mangled name of the global variable used on the device.
-
- In CUDA, __device__, __constant__ and __shared__ variables, as well as
- surface and texture variables, will generate a shadow symbol on host.
- We must preserve the correspodence in order to generate registration
- functions.
- }];
-
- let parameters = (ins "std::string":$device_side_name);
- let assemblyFormat = "`<` $device_side_name `>`";
-}
-
-def CUDAExternallyInitializedAttr : CIR_Attr<"CUDAExternallyInitialized",
+def CIR_CUDAExternallyInitializedAttr : CIR_Attr<"CUDAExternallyInitialized",
"cu.externally_initialized"> {
let summary = "The marked variable is externally initialized.";
let description =
@@ -69,4 +51,32 @@ def CUDAExternallyInitializedAttr : CIR_Attr<"CUDAExternallyInitialized",
let canHaveIllegalCXXABIType = 0;
}
+// Enum for device variable kinds
+def CIR_CUDADeviceVarKind : I32EnumAttr<"CUDADeviceVarKind",
+ "CUDA device variable kind", [
+ I32EnumAttrCase<"Variable", 0>,
+ I32EnumAttrCase<"Surface", 1>, // Future
+ I32EnumAttrCase<"Texture", 2>, // Future
+]> {
+ let cppNamespace = "::cir";
+}
+
+// Attribute carrying device variable registration flags
+def CIR_CUDAVarRegistrationInfoAttr : CIR_Attr<"CUDAVarRegistrationInfo", "cu.var_registration"> {
+ let summary = "Device variable registration flags.";
+ let parameters = (ins
+ "std::string":$device_side_name,
+ "CUDADeviceVarKind":$kind,
+ // Note: isExtern is only exercisable once HIP managed variable support
+ // lands. The only path to isExtern=true requires an extern __managed__
+ // variable, which bypasses the external storage filter in
+ // handleVarRegistration via hasAttr<HIPManagedAttr>().
+ "bool":$isExtern,
+ "bool":$isConstant,
+ "bool":$isManaged
+ );
+
+ let hasCustomAssemblyFormat = 1;
+}
+
#endif // CLANG_CIR_DIALECT_IR_CIRCUDAATTRS_TD
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index 773924e6ca301..3b598be1bc222 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -66,10 +66,25 @@ class CIRGenNVCUDARuntime : public CIRGenCUDARuntime {
void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
FunctionArgList &args) override;
+ void handleVarRegistration(const VarDecl *vd, cir::GlobalOp var) override;
+ void finalizeModule() override;
+
void internalizeDeviceSideVar(const VarDecl *d,
cir::GlobalLinkageKind &linkage) override;
std::string getDeviceSideName(const NamedDecl *nd) override;
+
+ void registerDeviceVar(const VarDecl *vd, cir::GlobalOp &var, bool isExtern,
+ bool isConstant) {
+ // Attach the device var attribute to the GlobalOp
+ auto &builder = cgm.getBuilder();
+ var->setAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic(),
+ cir::CUDAVarRegistrationInfoAttr::get(
+ builder.getContext(),
+ getDeviceSideName(cast<NamedDecl>(vd)),
+ cir::CUDADeviceVarKind::Variable, isExtern, isConstant,
+ vd->hasAttr<HIPManagedAttr>()));
+ }
};
} // namespace
@@ -403,3 +418,65 @@ std::string CIRGenNVCUDARuntime::getDeviceSideName(const NamedDecl *nd) {
}
return deviceSideName;
}
+
+void CIRGenNVCUDARuntime::handleVarRegistration(const VarDecl *vd,
+ cir::GlobalOp var) {
+ if (vd->hasAttr<CUDADeviceAttr>() || vd->hasAttr<CUDAConstantAttr>()) {
+ // Shadow variables and their properties must be registered with CUDA
+ // runtime. Skip Extern global variables, which will be registered in
+ // the TU where they are defined.
+ //
+ // Don't register a C++17 inline variable. The local symbol can be
+ // discarded and referencing a discarded local symbol from outside the
+ // comdat (__cuda_register_globals) is disallowed by the ELF spec.
+ //
+ // HIP managed variables need to be always recorded in device and host
+ // compilations for transformation.
+ //
+ // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
+ // added to llvm.compiler-used, therefore they are safe to be registered.
+ if ((!vd->hasExternalStorage() && !vd->isInline()) ||
+ cgm.getASTContext().CUDADeviceVarODRUsedByHost.contains(vd) ||
+ vd->hasAttr<HIPManagedAttr>()) {
+ registerDeviceVar(vd, var, !vd->hasDefinition(),
+ vd->hasAttr<CUDAConstantAttr>());
+ }
+ } else if (vd->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ vd->getType()->isCUDADeviceBuiltinTextureType()) {
+ // Builtin surfaces and textures and their template arguments are
+ // also registered with CUDA runtime.
+ cgm.errorNYI(vd->getSourceRange(),
+ "handleVarRegistration: Surface and Texture registration");
+ }
+}
+
+void CIRGenNVCUDARuntime::finalizeModule() {
+ if (!cgm.getLangOpts().CUDAIsDevice)
+ return;
+
+ // Mark ODR-used device variables as compiler used to prevent them from being
+ // eliminated by optimization. This is necessary for device variables
+ // ODR-used by host functions. Sema correctly marks them as ODR-used no
+ // matter whether they are ODR-used by device or host functions.
+ //
+ // We do not need to do this if the variable has used attribute since it
+ // has already been added.
+ //
+ // Static device variables have been externalized at this point, therefore
+ // variables with private or internal linkage need not be added.
+ for (auto globalOp : cgm.getModule().getOps<cir::GlobalOp>()) {
+ auto regAttr = globalOp->getAttrOfType<cir::CUDAVarRegistrationInfoAttr>(
+ cir::CUDAVarRegistrationInfoAttr::getMnemonic());
+ if (!regAttr)
+ continue;
+
+ auto kind = regAttr.getKind();
+ if (!globalOp.isDeclaration() &&
+ !cir::isLocalLinkage(globalOp.getLinkage()) &&
+ (kind == cir::CUDADeviceVarKind::Variable ||
+ kind == cir::CUDADeviceVarKind::Surface ||
+ kind == cir::CUDADeviceVarKind::Texture)) {
+ cgm.addCompilerUsedGlobal(globalOp);
+ }
+ }
+}
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
index 39b6571849f29..c5e91ed36e1b6 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h
@@ -55,6 +55,13 @@ class CIRGenCUDARuntime {
virtual void internalizeDeviceSideVar(const VarDecl *d,
cir::GlobalLinkageKind &linkage) = 0;
+ /// Check whether a variable is a device variable and register it if true.
+ virtual void handleVarRegistration(const VarDecl *vd, cir::GlobalOp var) = 0;
+
+ /// Perform module finalization: on device side, mark ODR-used device
+ /// variables as compiler-used. Mirrors OG's CGCUDARuntime::finalizeModule.
+ virtual void finalizeModule() {}
+
/// Returns function or variable name on device side even if the current
/// compilation is for host.
virtual std::string getDeviceSideName(const NamedDecl *nd) = 0;
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index c023a69870970..6fe8c8198ea1f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1269,9 +1269,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
// Adjust linkage of shadow variables in host compilation
getCUDARuntime().internalizeDeviceSideVar(vd, linkage);
}
- // TODO(cir):
- // Handle variable registration
- // getCUDARuntime().handleVarRegistration(vd, gv);
+ getCUDARuntime().handleVarRegistration(vd, gv);
}
// Decorate CUDA shadow variables with the cu.shadow_name attribute so we know
@@ -1295,8 +1293,11 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
getASTContext().CUDADeviceVarODRUsedByHost.contains(vd) ||
vd->hasAttr<HIPManagedAttr>()) {
auto shadowName = cudaRuntime->getDeviceSideName(cast<NamedDecl>(vd));
- auto attr = cir::CUDAShadowNameAttr::get(&getMLIRContext(), shadowName);
- gv->setAttr(cir::CUDAShadowNameAttr::getMnemonic(), attr);
+ auto attr = cir::CUDAVarRegistrationInfoAttr::get(
+ &getMLIRContext(), shadowName, cir::CUDADeviceVarKind::Variable,
+ vd->hasDefinition(), vd->hasAttr<CUDAConstantAttr>(),
+ vd->hasAttr<HIPManagedAttr>());
+ gv->setAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic(), attr);
}
}
@@ -3152,6 +3153,9 @@ void CIRGenModule::release() {
(getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD))
emitAMDGPUMetadata();
+ if (astContext.getLangOpts().CUDA && cudaRuntime)
+ getCUDARuntime().finalizeModule();
+
// There's a lot of code that is not implemented yet.
assert(!cir::MissingFeatures::cgmRelease());
}
@@ -3319,6 +3323,12 @@ void CIRGenModule::updateResolvedBlockAddress(cir::BlockAddressOp op,
it->second = newLabel;
}
+void CIRGenModule::addCompilerUsedGlobal(cir::GlobalOp gv) {
+ assert(!gv.isDeclaration() &&
+ "Only globals with definition can force usage.");
+ llvmCompilerUsed.emplace_back(gv);
+}
+
cir::LabelOp
CIRGenModule::lookupBlockAddressInfo(cir::BlockAddrInfoAttr blockInfo) {
return blockAddressInfoToLabel.lookup(blockInfo);
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 388b78f2a75e6..b4714b097d915 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -176,6 +176,10 @@ class CIRGenModule : public CIRGenTypeCache {
void mapResolvedBlockAddress(cir::BlockAddressOp op, cir::LabelOp);
void updateResolvedBlockAddress(cir::BlockAddressOp op,
cir::LabelOp newLabel);
+
+ /// Add a global value to the LLVMCompilerUsed list.
+ void addCompilerUsedGlobal(cir::GlobalOp gv);
+
/// Tell the consumer that this variable has been instantiated.
void handleCXXStaticMemberVarInstantiation(VarDecl *vd);
@@ -330,6 +334,8 @@ class CIRGenModule : public CIRGenTypeCache {
void emitVTable(const CXXRecordDecl *rd);
+ std::vector<cir::GlobalOp> llvmCompilerUsed;
+
/// Return the appropriate linkage for the vtable, VTT, and type information
/// of the given class.
cir::GlobalLinkageKind getVTableLinkage(const CXXRecordDecl *rd);
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index 27cba6a20b445..3e7b342bc401a 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -429,6 +429,72 @@ ConstComplexAttr::verify(function_ref<InFlightDiagnostic()> emitError,
return success();
}
+//===----------------------------------------------------------------------===//
+// CIR_CUDAVarRegistrationInfoAttr definitions
+//===----------------------------------------------------------------------===//
+
+void CUDAVarRegistrationInfoAttr::print(AsmPrinter &p) const {
+ p << "<" << getDeviceSideName();
+ p << ", " << stringifyEnum(getKind());
+ if (getIsExtern())
+ p << ", extern";
+ if (getIsConstant())
+ p << ", constant";
+ if (getIsManaged())
+ p << ", managed";
+ p << ">";
+}
+
+Attribute CUDAVarRegistrationInfoAttr::parse(AsmParser &parser, Type odsType) {
+ if (parser.parseLess())
+ return {};
+
+ std::string deviceSideName;
+ if (parser.parseString(&deviceSideName))
+ return {};
+
+ // Parse the device variable kind (Variable, Surface, Texture)
+ StringRef kindStr;
+ if (parser.parseKeyword(&kindStr))
+ return {};
+
+ auto kind = symbolizeCUDADeviceVarKind(kindStr);
+ if (!kind) {
+ parser.emitError(parser.getCurrentLocation(),
+ "unknown device variable kind: ")
+ << kindStr;
+ return {};
+ }
+
+ // Parse optional flags: extern, constant, managed
+ bool isExtern = false;
+ bool isConstant = false;
+ bool isManaged = false;
+
+ while (parser.parseOptionalGreater().failed()) {
+ if (parser.parseComma())
+ return {};
+
+ StringRef flag;
+ if (parser.parseKeyword(&flag))
+ return {};
+
+ if (flag == "extern")
+ isExtern = true;
+ else if (flag == "constant")
+ isConstant = true;
+ else if (flag == "managed")
+ isManaged = true;
+ else {
+ parser.emitError(parser.getCurrentLocation(), "unknown flag: ") << flag;
+ return {};
+ }
+ }
+
+ return get(parser.getContext(), deviceSideName, *kind, isExtern, isConstant,
+ isManaged);
+}
+
//===----------------------------------------------------------------------===//
// DataMemberAttr definitions
//===----------------------------------------------------------------------===//
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index c0f48e4c22aae..1208bce0f5c31 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -52,7 +52,7 @@
// CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0>
// LLVM-DEVICE-DAG: @i = addrspace(1) {{.*}}global i32 0
// OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0
-// CIR-HOST: cir.global {{.*}} @i = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<i>}
+// CIR-HOST: cir.global {{.*}} @i = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
// LLVM-HOST: @i = internal global i32 poison
// OGCG-HOST: @i = internal global i32 undef
__device__ int i;
@@ -61,7 +61,7 @@ __device__ int i;
// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0>
// LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0
// OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0
-// CIR-HOST: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.shadow_name = #cir.cu.shadow_name<j>}
+// CIR-HOST: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<j, Variable, extern, constant>}
// LLVM-HOST: @j = internal global i32 poison
// OGCG-HOST: @j = internal global i32 undef
__constant__ int j;
>From b8a3ea73a4c4323ea5ef0ec1a4a23cff7e97b669 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 03:22:59 +0700
Subject: [PATCH 09/12] [CIR] Remove git conflict symbols
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 5 +--
clang/lib/CIR/CodeGen/CIRGenModule.h | 5 ---
dump.cir | 48 ++++++++++++++++++++++++++
3 files changed, 49 insertions(+), 9 deletions(-)
create mode 100644 dump.cir
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 990461b0c025a..d3c1ecfcaa530 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1319,10 +1319,7 @@ static void emitUsed(CIRGenModule &cgm, StringRef name,
gv.setSectionAttr(builder.getStringAttr("llvm.metadata"));
}
-void CIRGenModule::emitLLVMUsed() {
- emitUsed(*this, "llvm.used", llvmUsed);
- emitUsed(*this, "llvm.compiler.used", llvmCompilerUsed);
-}
+void CIRGenModule::emitLLVMUsed() { emitUsed(*this, "llvm.used", llvmUsed); }
void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
bool isTentative) {
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 9c78768268c1a..0a14ae79bab6c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -209,10 +209,6 @@ class CIRGenModule : public CIRGenTypeCache {
void updateResolvedBlockAddress(cir::BlockAddressOp op,
cir::LabelOp newLabel);
-<<<<<<< HEAD
- /// Add a global value to the LLVMCompilerUsed list.
- void addCompilerUsedGlobal(cir::GlobalOp gv);
-=======
/// Add a global value to the llvmUsed list.
void addUsedGlobal(cir::CIRGlobalValueInterface gv);
@@ -224,7 +220,6 @@ class CIRGenModule : public CIRGenTypeCache {
/// Emit llvm.used and llvm.compiler.used globals.
void emitLLVMUsed();
->>>>>>> 2fdb09cf65e680094bdb3573435467ae544c7490
/// Tell the consumer that this variable has been instantiated.
void handleCXXStaticMemberVarInstantiation(VarDecl *vd);
diff --git a/dump.cir b/dump.cir
new file mode 100644
index 0000000000000..30835eb5071d5
--- /dev/null
+++ b/dump.cir
@@ -0,0 +1,48 @@
+!rec_cudaStream = !cir.record<struct "cudaStream" incomplete>
+!s32i = !cir.int<s, 32>
+!u32i = !cir.int<u, 32>
+!u64i = !cir.int<u, 64>
+!void = !cir.void
+!rec_dim3 = !cir.record<struct "dim3" {!u32i, !u32i, !u32i}>
+module @"/home/zaky/workspace/llvm-project/clang/test/CIR/CodeGenCUDA/address-spaces.cu" attributes {cir.lang = #cir.lang<cxx>, cir.module_asm = [], cir.triple = "x86_64-unknown-linux-gnu", dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>} {
+ cir.global "private" internal dso_local @i = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>} loc(#loc12)
+ cir.global "private" internal dso_local @j = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<j, Variable, extern, constant>} loc(#loc13)
+ cir.global "private" internal dso_local @k = #cir.poison : !s32i {alignment = 4 : i64} loc(#loc14)
+ cir.global "private" internal dso_local @b = #cir.poison : !cir.float {alignment = 4 : i64} loc(#loc15)
+ cir.func private dso_local @__cudaPopCallConfiguration(!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_cudaStream>>) -> !s32i loc(#loc)
+ cir.func private dso_local @cudaLaunchKernel(!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_cudaStream>) -> !u32i loc(#loc)
+ cir.func no_inline dso_local @_Z17__device_stub__fnv() attributes {cu.kernel_name = #cir.cu.kernel_name<_Z2fnv>, nothrow} {
+ %0 = cir.alloca !cir.array<!cir.ptr<!void> x 0>, !cir.ptr<!cir.array<!cir.ptr<!void> x 0>>, ["kernel_args"] {alignment = 16 : i64} loc(#loc16)
+ %1 = cir.cast array_to_ptrdecay %0 : !cir.ptr<!cir.array<!cir.ptr<!void> x 0>> -> !cir.ptr<!cir.ptr<!void>> loc(#loc16)
+ %2 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["grid_dim"] {alignment = 8 : i64} loc(#loc16)
+ %3 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["block_dim"] {alignment = 8 : i64} loc(#loc16)
+ %4 = cir.alloca !u64i, !cir.ptr<!u64i>, ["shared_mem"] {alignment = 8 : i64} loc(#loc16)
+ %5 = cir.alloca !cir.ptr<!rec_cudaStream>, !cir.ptr<!cir.ptr<!rec_cudaStream>>, ["stream"] {alignment = 8 : i64} loc(#loc16)
+ %6 = cir.call @__cudaPopCallConfiguration(%2, %3, %4, %5) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_cudaStream>>) -> !s32i loc(#loc16)
+ %7 = cir.get_global @_Z17__device_stub__fnv : !cir.ptr<!cir.func<()>> loc(#loc16)
+ %8 = cir.cast bitcast %7 : !cir.ptr<!cir.func<()>> -> !cir.ptr<!void> loc(#loc16)
+ %9 = cir.load %4 : !cir.ptr<!u64i>, !u64i loc(#loc16)
+ %10 = cir.load %5 : !cir.ptr<!cir.ptr<!rec_cudaStream>>, !cir.ptr<!rec_cudaStream> loc(#loc16)
+ %11 = cir.load align(8) %2 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc11)
+ %12 = cir.load align(8) %3 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc11)
+ %13 = cir.call @cudaLaunchKernel(%8, %11, %12, %1, %9, %10) : (!cir.ptr<!void> {llvm.noundef}, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>> {llvm.noundef}, !u64i {llvm.noundef}, !cir.ptr<!rec_cudaStream> {llvm.noundef}) -> (!u32i {llvm.noundef}) loc(#loc11)
+ cir.return loc(#loc10)
+ } loc(#loc16)
+} loc(#loc)
+#loc = loc("/home/zaky/workspace/llvm-project/clang/test/CIR/CodeGenCUDA/address-spaces.cu":0:0)
+#loc1 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":58:1)
+#loc2 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":58:16)
+#loc3 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":67:1)
+#loc4 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":67:18)
+#loc5 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":76:1)
+#loc6 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":76:16)
+#loc7 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":82:1)
+#loc8 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":82:18)
+#loc9 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":98:1)
+#loc10 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":102:1)
+#loc11 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":98:17)
+#loc12 = loc(fused[#loc1, #loc2])
+#loc13 = loc(fused[#loc3, #loc4])
+#loc14 = loc(fused[#loc5, #loc6])
+#loc15 = loc(fused[#loc7, #loc8])
+#loc16 = loc(fused[#loc9, #loc10])
>From 88d330b41d4ef9bfcfa03c7ad2281fd668413def Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 03:30:34 +0700
Subject: [PATCH 10/12] Merge branch 'main' of github.com:llvm/llvm-project
into shadow-variables
---
clang/lib/CIR/CodeGen/CIRGenModule.h | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 0a14ae79bab6c..2bd6010a5afed 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -405,8 +405,6 @@ class CIRGenModule : public CIRGenTypeCache {
void emitVTable(const CXXRecordDecl *rd);
- std::vector<cir::GlobalOp> llvmCompilerUsed;
-
/// Return the appropriate linkage for the vtable, VTT, and type information
/// of the given class.
cir::GlobalLinkageKind getVTableLinkage(const CXXRecordDecl *rd);
>From e420bd14914519a50f53cf33d80e9a1622b0832a Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 04:11:32 +0700
Subject: [PATCH 11/12] Fix wrong addCompilerUsedGlobal implementation
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 ----
clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 5 +--
clang/test/CIR/IR/var-registration.cir | 8 +++++
dump.cir | 48 --------------------------
4 files changed, 11 insertions(+), 56 deletions(-)
create mode 100644 clang/test/CIR/IR/var-registration.cir
delete mode 100644 dump.cir
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index d3c1ecfcaa530..9dbe3d664b6aa 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -3585,12 +3585,6 @@ void CIRGenModule::updateResolvedBlockAddress(cir::BlockAddressOp op,
it->second = newLabel;
}
-void CIRGenModule::addCompilerUsedGlobal(cir::GlobalOp gv) {
- assert(!gv.isDeclaration() &&
- "Only globals with definition can force usage.");
- llvmCompilerUsed.emplace_back(gv);
-}
-
cir::LabelOp
CIRGenModule::lookupBlockAddressInfo(cir::BlockAddrInfoAttr blockInfo) {
return blockAddressInfoToLabel.lookup(blockInfo);
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index fa86dbf2eb0a9..f8d4f61452cd3 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -449,9 +449,10 @@ Attribute CUDAVarRegistrationInfoAttr::parse(AsmParser &parser, Type odsType) {
if (parser.parseLess())
return {};
- std::string deviceSideName;
- if (parser.parseString(&deviceSideName))
+ StringRef deviceSideNameRef;
+ if (parser.parseKeyword(&deviceSideNameRef))
return {};
+ std::string deviceSideName = deviceSideNameRef.str();
// Parse the device variable kind (Variable, Surface, Texture)
StringRef kindStr;
diff --git a/clang/test/CIR/IR/var-registration.cir b/clang/test/CIR/IR/var-registration.cir
new file mode 100644
index 0000000000000..630f86f522d1c
--- /dev/null
+++ b/clang/test/CIR/IR/var-registration.cir
@@ -0,0 +1,8 @@
+// RUN: cir-opt %s --verify-roundtrip | FileCheck %s
+
+!s32i = !cir.int<s, 32>
+
+module {
+ // CHECK: cir.global {{.*}} @i = #cir.poison : !s32i {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
+ cir.global "private" internal dso_local @i = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
+}
diff --git a/dump.cir b/dump.cir
deleted file mode 100644
index 30835eb5071d5..0000000000000
--- a/dump.cir
+++ /dev/null
@@ -1,48 +0,0 @@
-!rec_cudaStream = !cir.record<struct "cudaStream" incomplete>
-!s32i = !cir.int<s, 32>
-!u32i = !cir.int<u, 32>
-!u64i = !cir.int<u, 64>
-!void = !cir.void
-!rec_dim3 = !cir.record<struct "dim3" {!u32i, !u32i, !u32i}>
-module @"/home/zaky/workspace/llvm-project/clang/test/CIR/CodeGenCUDA/address-spaces.cu" attributes {cir.lang = #cir.lang<cxx>, cir.module_asm = [], cir.triple = "x86_64-unknown-linux-gnu", dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>} {
- cir.global "private" internal dso_local @i = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>} loc(#loc12)
- cir.global "private" internal dso_local @j = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<j, Variable, extern, constant>} loc(#loc13)
- cir.global "private" internal dso_local @k = #cir.poison : !s32i {alignment = 4 : i64} loc(#loc14)
- cir.global "private" internal dso_local @b = #cir.poison : !cir.float {alignment = 4 : i64} loc(#loc15)
- cir.func private dso_local @__cudaPopCallConfiguration(!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_cudaStream>>) -> !s32i loc(#loc)
- cir.func private dso_local @cudaLaunchKernel(!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_cudaStream>) -> !u32i loc(#loc)
- cir.func no_inline dso_local @_Z17__device_stub__fnv() attributes {cu.kernel_name = #cir.cu.kernel_name<_Z2fnv>, nothrow} {
- %0 = cir.alloca !cir.array<!cir.ptr<!void> x 0>, !cir.ptr<!cir.array<!cir.ptr<!void> x 0>>, ["kernel_args"] {alignment = 16 : i64} loc(#loc16)
- %1 = cir.cast array_to_ptrdecay %0 : !cir.ptr<!cir.array<!cir.ptr<!void> x 0>> -> !cir.ptr<!cir.ptr<!void>> loc(#loc16)
- %2 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["grid_dim"] {alignment = 8 : i64} loc(#loc16)
- %3 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["block_dim"] {alignment = 8 : i64} loc(#loc16)
- %4 = cir.alloca !u64i, !cir.ptr<!u64i>, ["shared_mem"] {alignment = 8 : i64} loc(#loc16)
- %5 = cir.alloca !cir.ptr<!rec_cudaStream>, !cir.ptr<!cir.ptr<!rec_cudaStream>>, ["stream"] {alignment = 8 : i64} loc(#loc16)
- %6 = cir.call @__cudaPopCallConfiguration(%2, %3, %4, %5) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_cudaStream>>) -> !s32i loc(#loc16)
- %7 = cir.get_global @_Z17__device_stub__fnv : !cir.ptr<!cir.func<()>> loc(#loc16)
- %8 = cir.cast bitcast %7 : !cir.ptr<!cir.func<()>> -> !cir.ptr<!void> loc(#loc16)
- %9 = cir.load %4 : !cir.ptr<!u64i>, !u64i loc(#loc16)
- %10 = cir.load %5 : !cir.ptr<!cir.ptr<!rec_cudaStream>>, !cir.ptr<!rec_cudaStream> loc(#loc16)
- %11 = cir.load align(8) %2 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc11)
- %12 = cir.load align(8) %3 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc11)
- %13 = cir.call @cudaLaunchKernel(%8, %11, %12, %1, %9, %10) : (!cir.ptr<!void> {llvm.noundef}, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>> {llvm.noundef}, !u64i {llvm.noundef}, !cir.ptr<!rec_cudaStream> {llvm.noundef}) -> (!u32i {llvm.noundef}) loc(#loc11)
- cir.return loc(#loc10)
- } loc(#loc16)
-} loc(#loc)
-#loc = loc("/home/zaky/workspace/llvm-project/clang/test/CIR/CodeGenCUDA/address-spaces.cu":0:0)
-#loc1 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":58:1)
-#loc2 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":58:16)
-#loc3 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":67:1)
-#loc4 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":67:18)
-#loc5 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":76:1)
-#loc6 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":76:16)
-#loc7 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":82:1)
-#loc8 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":82:18)
-#loc9 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":98:1)
-#loc10 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":102:1)
-#loc11 = loc("clang/test/CIR/CodeGenCUDA/address-spaces.cu":98:17)
-#loc12 = loc(fused[#loc1, #loc2])
-#loc13 = loc(fused[#loc3, #loc4])
-#loc14 = loc(fused[#loc5, #loc6])
-#loc15 = loc(fused[#loc7, #loc8])
-#loc16 = loc(fused[#loc9, #loc10])
>From cf8ec86b540f263540805486954ecdda66baefb2 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 12:31:01 +0700
Subject: [PATCH 12/12] [CIR] Fix parser
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 12 +++++++++---
clang/test/CIR/IR/var-registration.cir | 2 +-
2 files changed, 10 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index f8d4f61452cd3..f842aeaaeacd3 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h"
#include "clang/CIR/Dialect/IR/CIRDialect.h"
+#include "mlir/IR/Attributes.h"
#include "mlir/IR/DialectImplementation.h"
#include "llvm/ADT/TypeSwitch.h"
@@ -449,10 +450,15 @@ Attribute CUDAVarRegistrationInfoAttr::parse(AsmParser &parser, Type odsType) {
if (parser.parseLess())
return {};
- StringRef deviceSideNameRef;
- if (parser.parseKeyword(&deviceSideNameRef))
+ std::string deviceSideName;
+ if (parser.parseKeywordOrString(&deviceSideName)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "expected device variable name");
+ return {};
+ }
+
+ if (parser.parseComma())
return {};
- std::string deviceSideName = deviceSideNameRef.str();
// Parse the device variable kind (Variable, Surface, Texture)
StringRef kindStr;
diff --git a/clang/test/CIR/IR/var-registration.cir b/clang/test/CIR/IR/var-registration.cir
index 630f86f522d1c..a6edf942701fa 100644
--- a/clang/test/CIR/IR/var-registration.cir
+++ b/clang/test/CIR/IR/var-registration.cir
@@ -3,6 +3,6 @@
!s32i = !cir.int<s, 32>
module {
- // CHECK: cir.global {{.*}} @i = #cir.poison : !s32i {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
+ // CHECK: cir.global {{.*}} @i = #cir.poison : !s32i {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
cir.global "private" internal dso_local @i = #cir.poison : !s32i {alignment = 4 : i64, cu.var_registration = #cir.cu.var_registration<i, Variable, extern>}
}
More information about the cfe-commits
mailing list