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 <[email protected]> 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() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
