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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
