[clang] 1639504 - [CIR] Attach address space to global variables (#226455)

[via cfe-commits]

Author: Steffen Larsen
Date: 2026-09-25T23:46:41-04:00
New Revision: 163950475fc724cd5d62e868bdb46ba0c39e5519

URL: https://github.com/llvm/llvm-project/commit/163950475fc724cd5d62e868bdb46ba0c39e5519
DIFF: https://github.com/llvm/llvm-project/commit/163950475fc724cd5d62e868bdb46ba0c39e5519.diff

LOG: [CIR] Attach address space to global variables (#226455)

Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>

Added: 
    

Modified: 
    clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
    clang/lib/CIR/CodeGen/CIRGenDecl.cpp
    clang/test/CIR/CodeGenCUDA/address-spaces.cu

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
index 29f1a64ad1d17..36d583cfe9fbe 100644
--- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
+++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
@@ -491,10 +491,10 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
 
   cir::GetGlobalOp createGetGlobal(mlir::Location loc, cir::GlobalOp global,
                                    bool threadLocal = false) {
-    assert(!cir::MissingFeatures::addressSpace());
-    return cir::GetGlobalOp::create(*this, loc,
-                                    getPointerTo(global.getSymType()),
-                                    global.getSymNameAttr(), threadLocal);
+    return cir::GetGlobalOp::create(
+        *this, loc,
+        getPointerTo(global.getSymType(), global.getAddrSpaceAttr()),
+        global.getSymNameAttr(), threadLocal);
   }
 
   cir::GetGlobalOp createGetGlobal(cir::GlobalOp global,

diff  --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 689ead0c68a67..451f6f8af7fd1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -528,7 +528,6 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &d,
   std::string name = getStaticDeclName(*this, d);
 
   mlir::Type lty = getTypes().convertTypeForMem(ty);
-  assert(!cir::MissingFeatures::addressSpace());
 
   // OpenCL variables in local address space and CUDA shared
   // variables cannot have an initializer.
@@ -539,8 +538,12 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &d,
   else
     init = builder.getZeroInitAttr(convertType(ty));
 
-  cir::GlobalOp gv = builder.createVersionedGlobal(
-      getModule(), getLoc(d.getLocation()), name, lty, false, linkage);
+  mlir::ptr::MemorySpaceAttrInterface addrSpace = cir::toCIRAddressSpaceAttr(
+      getMLIRContext(), getGlobalVarAddressSpace(&d));
+
+  cir::GlobalOp gv =
+      builder.createVersionedGlobal(getModule(), getLoc(d.getLocation()), name,
+                                    lty, false, linkage, addrSpace);
   insertGlobalSymbol(gv);
   // TODO(cir): infer visibility from linkage in global op builder.
   gv.setVisibility(getMLIRVisibilityFromCIRLinkage(linkage));

diff  --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 9e923547c21c4..6637100fd76c9 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -45,8 +45,8 @@
 
 // Verifies CIR emits correct address spaces for CUDA globals.
 
-// CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j = #cir.undef
-// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef
+// CIR-DEVICE: cir.global "private" internal dso_local target_address_space(3) @_ZZ2fnvE1j = #cir.undef
+// LLVM-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef
 
 // 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>
@@ -161,16 +161,16 @@ __global__ void fn() {
 // CIR-DEVICE:   %[[ALLOCA:.*]] = cir.alloca "i" {{.*}} init : !cir.ptr<!s32i>
 // CIR-DEVICE:   %[[ZERO:.*]] = cir.const #cir.int<0> : !s32i
 // CIR-DEVICE:   cir.store {{.*}}%[[ZERO]], %[[ALLOCA]] : !s32i, !cir.ptr<!s32i>
-// CIR-DEVICE:   %[[J:.*]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i>
+// CIR-DEVICE:   %[[J:.*]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i, target_address_space(3)>
 // CIR-DEVICE:   %[[VAL:.*]] = cir.load {{.*}}%[[ALLOCA]] : !cir.ptr<!s32i>, !s32i
-// CIR-DEVICE:   cir.store {{.*}}%[[VAL]], %[[J]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE:   cir.store {{.*}}%[[VAL]], %[[J]] : !s32i, !cir.ptr<!s32i, target_address_space(3)>
 // CIR-DEVICE:   cir.return
 
 // LLVM-DEVICE: define dso_local ptx_kernel void @_Z2fnv()
 // LLVM-DEVICE:   %[[ALLOCA:.*]] = alloca i32, 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:   store i32 %[[VAL]], ptr addrspace(3) @_ZZ2fnvE1j, align 4
 // LLVM-DEVICE:   ret void
 
 // OGCG-DEVICE: define dso_local ptx_kernel void @_Z2fnv()