[clang] [CIR][SPIR-V] Move global address space selection to a target hook (PR #231098)

[Arseniy Obolenskiy via cfe-commits]

https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/231098

Discussed in https://github.com/llvm/llvm-project/pull/230813

>From 5fa9dd54af4fa815127a7c923c51a047046d10f5 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 12 Oct 2026 10:44:41 +0200
Subject: [PATCH] [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; }