https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231098
>From 5fa9dd54af4fa815127a7c923c51a047046d10f5 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 12 Oct 2026 10:44:41 +0200 Subject: [PATCH 1/2] [CIR][SPIR-V][NFC] Move global address space selection to a target hook --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 9 +-------- clang/lib/CIR/CodeGen/TargetInfo.h | 7 +++++++ clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 12 ++++++++++++ .../test/CIR/CodeGenHIP/spirv-constant-addrspace.hip | 11 +++++++++++ 4 files changed, 31 insertions(+), 8 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index db505dc9824d0..0c1a2dfccd0dc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2330,14 +2330,7 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s, 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); + return getTargetCIRGenInfo().getGlobalCIRAddressSpace(as); } // TODO(cir): this could be a common AST helper for both CIR and LLVM codegen. diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 91d585fe00a19..09f41fcd3cf9b 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -70,6 +70,13 @@ class TargetCIRGenInfo { cir::LangAddressSpace::Default); } + /// Get the CIR address space attribute for a global in AST address space + /// \p as. + virtual mlir::ptr::MemorySpaceAttrInterface + getGlobalCIRAddressSpace(clang::LangAS as) const { + return cir::toCIRAddressSpaceAttr(info->cgt.getMLIRContext(), as); + } + /// Get the CIR value of a null pointer of type \p ptrTy, where \p qt is the /// source pointer type. virtual mlir::Value getNullPointer(CIRGenModule &cgm, cir::PointerType ptrTy, diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 8e36b25301dc9..2231d96550cdc 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -40,6 +40,18 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { cir::LangAddressSpace::OffloadPrivate); } + mlir::ptr::MemorySpaceAttrInterface + getGlobalCIRAddressSpace(LangAS as) const override { + // The cir::LangAddressSpace fold loses the SPIRDefIsGenMap split between + // cuda_constant and opencl_constant, so use the target map like classic. + if (as == LangAS::cuda_device || as == LangAS::cuda_constant || + as == LangAS::cuda_shared) + return cir::TargetAddressSpaceAttr::get( + &getABIInfo().cgt.getMLIRContext(), + getABIInfo().cgt.getASTContext().getTargetAddressSpace(as)); + return TargetCIRGenInfo::getGlobalCIRAddressSpace(as); + } + void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, CIRGenModule &cgm) const override { auto func = mlir::dyn_cast<cir::FuncOp>(global); diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index 92d10dcbf1e86..8c3cbe27254af 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -7,12 +7,18 @@ #define __device__ __attribute__((device)) #define __constant__ __attribute__((constant)) +#define __shared__ __attribute__((shared)) // 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] +// CIR-DAG: cir.global {{.*}} target_address_space(1) @dev +// CIR-DAG: cir.global {{.*}} target_address_space(3) @_ZZ10get_sharedvE2sh // LLVM-DAG: @_ZZ10static_fmtiE3fmt = internal addrspace(1) constant [3 x i8] c"%d\00" +// LLVM-DAG: @dev = addrspace(1) externally_initialized global i32 3 +// LLVM-DAG: @_ZZ10get_sharedvE2sh = internal addrspace(3) global [4 x i32] __constant__ int tab[2] = {1, 2}; +__device__ int dev = 3; // CIR-LABEL: cir.func{{.*}} @_Z10static_fmti // CIR: %[[G:.*]] = cir.get_global @_ZZ10static_fmtiE3fmt : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> @@ -33,6 +39,11 @@ __device__ const char *static_fmt(int i) { // LLVM: addrspacecast (ptr addrspace(1) @tab to ptr addrspace(4)) __device__ const int *get_tab() { return tab; } +__device__ int get_shared() { + __shared__ int sh[4]; + return dev + sh[0]; +} + // 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 e4df90303bf2b922ff0ab65fcd4f76a51db7ee07 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 12 Oct 2026 16:58:07 +0200 Subject: [PATCH 2/2] comment --- clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index 8c3cbe27254af..9794a4c0dff06 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -4,6 +4,15 @@ // 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 +// RUN: %clang_cc1 -triple spir64 -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=SPIR %s +// RUN: %clang_cc1 -triple spir64 -x hip -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=SPIR %s + +// SPIR-DAG: @tab = externally_initialized constant [2 x i32] +// SPIR-DAG: @dev = externally_initialized global i32 3 +// SPIR-DAG: @_ZZ10static_fmtiE3fmt = internal constant [3 x i8] +// SPIR-DAG: @_ZZ10get_sharedvE2sh = internal global [4 x i32] #define __device__ __attribute__((device)) #define __constant__ __attribute__((constant)) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
