llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/232173.diff 2 Files Affected: - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+5-1) - (modified) clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip (+4) ``````````diff 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)> `````````` </details> https://github.com/llvm/llvm-project/pull/232173 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
