https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/233720
Alias initializer pointer was hardcoded to address space 0, which trips an assertion on spirv64-amd-amdhsa where functions live in space 4 >From 1a903aa9e324bcc824285cd7e8d1b7092631ca22 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Tue, 20 Oct 2026 08:59:01 +0200 Subject: [PATCH] [CIR][SPIR-V] Use the program address space for function aliases Alias initializer pointer was hardcoded to address space 0, which trips an assertion on spirv64-amd-amdhsa where functions live in space 4 --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 6 +++--- clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip | 5 +++++ 2 files changed, 8 insertions(+), 3 deletions(-) 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(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
