https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/230813
CIR treated HIP `__constant__` like OpenCL constant, which is addrspace(2) on SPIR-V Align behavior with classic codegen >From 17178290a937d608d64f3dd1a237bc5b1263b126 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Sat, 10 Oct 2026 19:10:33 +0200 Subject: [PATCH] [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; } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
