[clang] [CIR][SPIR-V] Use the program address space for function aliases (PR #233720)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

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