[clang] 09fa829 - [CIR][SPIR-V] Use the program address space for function global views (#232173)

[via cfe-commits]

Author: Arseniy Obolenskiy
Date: 2026-10-15T13:00:26+02:00
New Revision: 09fa829c48bc5bede870c653c2f2c1f791d0751d

URL: https://github.com/llvm/llvm-project/commit/09fa829c48bc5bede870c653c2f2c1f791d0751d
DIFF: https://github.com/llvm/llvm-project/commit/09fa829c48bc5bede870c653c2f2c1f791d0751d.diff

LOG: [CIR][SPIR-V] Use the program address space for function global views (#232173)

CIR assumed functions are in address space 0 when taking their address
in a constant

On spirv64-amd-amdhsa functions are in address space 4, so
`__attribute__((used))` device function crashed while building llvm.used

Added: 
    

Modified: 
    clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
    clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
    clang/test/CIR/Lowering/address-space.cir

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index fb2be40281883..f0c1021fcd6c4 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -40,6 +40,7 @@
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "clang/CIR/Dialect/IR/CIRAttrs.h"
+#include "clang/CIR/Dialect/IR/CIRDataLayout.h"
 #include "clang/CIR/Dialect/IR/CIRDialect.h"
 #include "clang/CIR/Dialect/IR/CIRTypes.h"
 #include "clang/CIR/Dialect/Passes.h"
@@ -997,7 +998,8 @@ mlir::Value CIRAttrToValue::visitCirAttr(cir::ConstVectorAttr attr) {
 CIRAttrToValue::LoweredGlobalSymbol
 CIRAttrToValue::getAddressOfSymbol(mlir::FlatSymbolRefAttr symbol) {
   auto moduleOp = parentOp->getParentOfType<mlir::ModuleOp>();
-  mlir::DataLayout dataLayout(moduleOp);
+  cir::CIRDataLayout cirDataLayout(moduleOp);
+  const mlir::DataLayout &dataLayout = cirDataLayout.layout;
   mlir::Type sourceType;
   unsigned addrSpace = 0;
   llvm::StringRef symName;
@@ -1016,9 +1018,11 @@ CIRAttrToValue::getAddressOfSymbol(mlir::FlatSymbolRefAttr symbol) {
   } else if (auto llvmFun = dyn_cast<mlir::LLVM::LLVMFuncOp>(sourceSymbol)) {
     sourceType = llvmFun.getFunctionType();
     symName = llvmFun.getSymName();
+    addrSpace = cirDataLayout.getProgramAddressSpace();
   } else if (auto fun = dyn_cast<cir::FuncOp>(sourceSymbol)) {
     sourceType = converter->convertType(fun.getFunctionType());
     symName = fun.getSymName();
+    addrSpace = cirDataLayout.getProgramAddressSpace();
   } else if (auto alias = dyn_cast<mlir::LLVM::AliasOp>(sourceSymbol)) {
     sourceType = alias.getType();
     symName = alias.getSymName();

diff  --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
index 9794a4c0dff06..1ddc16332f65c 100644
--- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
+++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
@@ -29,6 +29,10 @@
 __constant__ int tab[2] = {1, 2};
 __device__ int dev = 3;
 
+// Functions live in the program address space 4.
+// LLVM-DAG: @llvm.used = {{.*}}ptr addrspace(4) @_Z7used_fnv
+__attribute__((used)) __device__ void used_fn() {}
+
 // 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)>

diff  --git a/clang/test/CIR/Lowering/address-space.cir b/clang/test/CIR/Lowering/address-space.cir
index 5aadac7f41b54..9a978d8eeb89b 100644
--- a/clang/test/CIR/Lowering/address-space.cir
+++ b/clang/test/CIR/Lowering/address-space.cir
@@ -1,5 +1,7 @@
 // RUN: cir-translate %s -cir-to-llvmir --target spirv64-unknown-unknown -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
+// RUN: cir-translate %s -cir-to-llvmir --target spirv64-amd-amdhsa -o %t.amd.ll
+// RUN: FileCheck --input-file=%t.amd.ll %s -check-prefix=AMD
 
 !s32i = !cir.int<s, 32>
 
@@ -7,6 +9,19 @@ module {
   cir.global external target_address_space(7) @addrspace3 = #cir.int<3> : !s32i
   // LLVM: @addrspace3 = addrspace(7) global i32
 
+  // Views of an llvm.func (fn_first) and a cir.func (fn_last). Keep before
+  // other global views, which would cache fn_first as a cir.func.
+  cir.func @fn_first() {
+    cir.return
+  }
+  cir.global external @fn_view_after = #cir.global_view<@fn_first> : !cir.ptr<!cir.void>
+  // AMD: @fn_view_after = global ptr addrspacecast (ptr addrspace(4) @fn_first to ptr)
+  cir.global external @fn_view_before = #cir.global_view<@fn_last> : !cir.ptr<!cir.void>
+  // AMD: @fn_view_before = global ptr addrspacecast (ptr addrspace(4) @fn_last to ptr)
+  cir.func @fn_last() {
+    cir.return
+  }
+
   // Test GlobalViewAttr with address space cast.
   cir.global external target_address_space(1) @global_in_as1 = #cir.int<42> : !s32i
   // LLVM: @global_in_as1 = addrspace(1) global i32 42