https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/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 >From 9debfd9eb656f81103b2386241396dcb4163a902 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 15 Oct 2026 06:09:27 +0200 Subject: [PATCH] [CIR][SPIR-V] Use the program address space for function global views 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 --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 6 +++++- clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip | 4 ++++ 2 files changed, 9 insertions(+), 1 deletion(-) 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)> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
