llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Alias initializer pointer was hardcoded to address space 0, which trips an assertion on spirv64-amd-amdhsa where functions live in space 4 --- Full diff: https://github.com/llvm/llvm-project/pull/233720.diff 2 Files Affected: - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+3-3) - (modified) clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip (+5) ``````````diff diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 40e6a5624366c..ac005a4e838d1 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -3058,9 +3058,9 @@ mlir::LogicalResult CIRToLLVMFuncOpLowering::matchAndRewriteAlias( mlir::OpBuilder builder(op.getContext()); mlir::Block *block = builder.createBlock(&aliasOp.getInitializerRegion()); builder.setInsertionPointToStart(block); - // The type of AddressOfOp is always a pointer. - assert(!cir::MissingFeatures::addressSpace()); - mlir::Type ptrTy = mlir::LLVM::LLVMPointerType::get(ty.getContext()); + cir::CIRDataLayout cirDataLayout(aliasOp->getParentOfType<mlir::ModuleOp>()); + mlir::Type ptrTy = mlir::LLVM::LLVMPointerType::get( + ty.getContext(), cirDataLayout.getProgramAddressSpace()); auto addrOp = mlir::LLVM::AddressOfOp::create(builder, loc, ptrTy, aliasee); mlir::LLVM::ReturnOp::create(builder, loc, addrOp); diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index 1ddc16332f65c..ac9dc614cbfb8 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -26,6 +26,7 @@ // 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] +// LLVM-DAG: @_Z5aliasv = alias i32 (), ptr addrspace(4) @alias_target __constant__ int tab[2] = {1, 2}; __device__ int dev = 3; @@ -60,3 +61,7 @@ __device__ int get_shared() { // 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; } + +// Function aliases point into the program address space 4. +extern "C" __device__ int alias_target() { return 1; } +[[gnu::alias("alias_target")]] __device__ int alias(); `````````` </details> https://github.com/llvm/llvm-project/pull/233720 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
