[clang] [CIR][HIP][SPIR-V] Emit constant globals in addrspace(1) (PR #230813)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 12 00:50:26 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/230813
>From 17178290a937d608d64f3dd1a237bc5b1263b126 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Sat, 10 Oct 2026 19:10:33 +0200
Subject: [PATCH 1/2] [CIR][HIP][SPIR-V] Emit constant globals in addrspace(1)
CIR treated HIP `__constant__` like OpenCL constant, which is addrspace(2) on SPIR-V
Align behavior with classic codegen
---
clang/lib/CIR/CodeGen/CIRGenDecl.cpp | 4 +-
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 22 ++++++++---
clang/lib/CIR/CodeGen/CIRGenModule.h | 2 +
.../CodeGenHIP/spirv-constant-addrspace.hip | 38 +++++++++++++++++++
4 files changed, 59 insertions(+), 7 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 01b9f810fbed4..f841cecbe0a85 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -538,8 +538,8 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &d,
else
init = emitNullConstantAttr(ty);
- mlir::ptr::MemorySpaceAttrInterface addrSpace = cir::toCIRAddressSpaceAttr(
- getMLIRContext(), getGlobalVarAddressSpace(&d));
+ mlir::ptr::MemorySpaceAttrInterface addrSpace =
+ getGlobalCIRAddressSpace(getGlobalVarAddressSpace(&d));
cir::GlobalOp gv =
builder.createVersionedGlobal(getModule(), getLoc(d.getLocation()), name,
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 1f80f4e0ebf08..db505dc9824d0 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1294,7 +1294,8 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
assert(!cir::MissingFeatures::openMP());
if (entry.getSymType() == ty &&
- cir::isMatchingAddressSpace(entryCIRAS, langAS))
+ cir::normalizeDefaultAddressSpace(entryCIRAS) ==
+ cir::normalizeDefaultAddressSpace(getGlobalCIRAddressSpace(langAS)))
return entry;
// If there are two attempts to define the same mangled name, issue an
@@ -1340,7 +1341,7 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
}
mlir::ptr::MemorySpaceAttrInterface declCIRAS =
- cir::toCIRAddressSpaceAttr(getMLIRContext(), getGlobalVarAddressSpace(d));
+ getGlobalCIRAddressSpace(getGlobalVarAddressSpace(d));
// mlir::SymbolTable::Visibility::Public is the default, no need to explicitly
// mark it as such.
@@ -2221,7 +2222,7 @@ generateStringLiteral(mlir::Location loc, mlir::TypedAttr c,
LangAS as = CodeGenUtils::getGlobalConstantAddressSpace(cgm.getLangOpts(),
cgm.getTarget());
mlir::ptr::MemorySpaceAttrInterface addrSpace =
- cir::toCIRAddressSpaceAttr(cgm.getMLIRContext(), as);
+ cgm.getGlobalCIRAddressSpace(as);
// Create a global variable for this string
// FIXME(cir): check for insertion point in module level.
@@ -2327,6 +2328,18 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s,
return builder.getGlobalViewAttr(ptrTy, gv);
}
+mlir::ptr::MemorySpaceAttrInterface
+CIRGenModule::getGlobalCIRAddressSpace(LangAS as) {
+ // The cir::LangAddressSpace fold loses the SPIRDefIsGenMap split between
+ // cuda_constant and opencl_constant, so use the target map like classic.
+ if (langOpts.CUDAIsDevice && getTriple().isSPIRV() &&
+ (as == LangAS::cuda_device || as == LangAS::cuda_constant ||
+ as == LangAS::cuda_shared))
+ return cir::TargetAddressSpaceAttr::get(
+ &getMLIRContext(), astContext.getTargetAddressSpace(as));
+ return cir::toCIRAddressSpaceAttr(getMLIRContext(), as);
+}
+
// TODO(cir): this could be a common AST helper for both CIR and LLVM codegen.
LangAS CIRGenModule::getLangTempAllocaAddressSpace() const {
if (getLangOpts().OpenCL)
@@ -3940,8 +3953,7 @@ void CIRGenModule::release() {
auto int8Ty = cir::IntType::get(&getMLIRContext(), 8, /*isSigned=*/false);
auto loc = builder.getUnknownLoc();
mlir::ptr::MemorySpaceAttrInterface addrSpace =
- cir::LangAddressSpaceAttr::get(&getMLIRContext(),
- getGlobalVarAddressSpace(nullptr));
+ getGlobalCIRAddressSpace(getGlobalVarAddressSpace(nullptr));
auto gv = createGlobalOp(loc, cuidName, int8Ty,
/*isConstant=*/false, addrSpace);
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 05ba8b5ee212e..3e0bc06bf191e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -470,6 +470,8 @@ class CIRGenModule : public CIRGenTypeCache {
/// space, which may lead to type mismatches in other parts of the IR.
LangAS getLangTempAllocaAddressSpace() const;
+ mlir::ptr::MemorySpaceAttrInterface getGlobalCIRAddressSpace(LangAS as);
+
/// Set attributes which are common to any form of a global definition (alias,
/// Objective-C method, function, global variable).
///
diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
new file mode 100644
index 0000000000000..92d10dcbf1e86
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+
+#define __device__ __attribute__((device))
+#define __constant__ __attribute__((constant))
+
+// CIR-DAG: cir.global {{.*}} target_address_space(1) @tab
+// CIR-DAG: cir.global {{.*}} target_address_space(1) @_ZZ10static_fmtiE3fmt
+// LLVM-DAG: @tab = addrspace(1) externally_initialized constant [2 x i32] [i32 1, i32 2]
+// LLVM-DAG: @_ZZ10static_fmtiE3fmt = internal addrspace(1) constant [3 x i8] c"%d\00"
+__constant__ int tab[2] = {1, 2};
+
+// CIR-LABEL: cir.func{{.*}} @_Z10static_fmti
+// CIR: %[[G:.*]] = cir.get_global @_ZZ10static_fmtiE3fmt : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)>
+// CIR: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)>
+
+// LLVM-LABEL: define{{.*}} ptr addrspace(4) @_Z10static_fmti
+// LLVM: addrspacecast (ptr addrspace(1) @_ZZ10static_fmtiE3fmt to ptr addrspace(4))
+__device__ const char *static_fmt(int i) {
+ static const char fmt[] = "%d";
+ return fmt + i;
+}
+
+// CIR-LABEL: cir.func{{.*}} @_Z7get_tabv
+// CIR: %[[G:.*]] = cir.get_global @tab : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(1)>
+// CIR: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(1)> -> !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)>
+
+// LLVM-LABEL: define{{.*}} ptr addrspace(4) @_Z7get_tabv
+// LLVM: addrspacecast (ptr addrspace(1) @tab to ptr addrspace(4))
+__device__ const int *get_tab() { return tab; }
+
+// An explicit opencl_constant qualifier keeps UniformConstant.
+// LLVM-LABEL: define{{.*}} i32 @_Z6get_ocPU3AS2i(ptr addrspace(2)
+__device__ int get_oc(__attribute__((opencl_constant)) int *p) { return *p; }
>From e219897143983f2d5c07970b8b101bd3a365ca66 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 12 Oct 2026 09:50:14 +0200
Subject: [PATCH 2/2] comment
---
clang/lib/CIR/CodeGen/CIRGenModule.h | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 3e0bc06bf191e..22c3c391ee7d8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -470,6 +470,8 @@ class CIRGenModule : public CIRGenTypeCache {
/// space, which may lead to type mismatches in other parts of the IR.
LangAS getLangTempAllocaAddressSpace() const;
+ /// Return the CIR address space attribute for a global in \p as, honoring
+ /// the target address space map where the LangAS fold would lose it.
mlir::ptr::MemorySpaceAttrInterface getGlobalCIRAddressSpace(LangAS as);
/// Set attributes which are common to any form of a global definition (alias,
More information about the cfe-commits
mailing list