[clang] fa0638e - [CIR] Cast alloca for function parameter to the right address space (#220836)

[via cfe-commits]

Author: Mariya Podchishchaeva
Date: 2026-09-08T12:47:58+02:00
New Revision: fa0638edc6628ce2c386519ffdcbdcf35513e67e

URL: https://github.com/llvm/llvm-project/commit/fa0638edc6628ce2c386519ffdcbdcf35513e67e
DIFF: https://github.com/llvm/llvm-project/commit/fa0638edc6628ce2c386519ffdcbdcf35513e67e.diff

LOG: [CIR] Cast alloca for function parameter to the right address space (#220836)

Function parameters are stored into temporary allocas and for
address-space aware targets allocas may yield pointers to address spaces
that are different from default address space or the address space of
the parameter type. Do a cast to avoid mismatch. The address space
mismatch was reproduced using cir.ternary op returning a function
parameter and a local variable. For local variables and return
temporaries we already cast alloca address space to temporary address
space but not for function parameter allocas.

Added: 
    clang/test/CIR/CodeGenHIP/ternary-struct-addrspace.hip

Modified: 
    clang/lib/CIR/CodeGen/CIRGenFunction.cpp
    clang/lib/CIR/CodeGen/CIRGenFunction.h
    clang/test/CIR/CodeGenOpenMP/target-map.c

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
index 25ae08364713b..6da9e8ac88e16 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -488,10 +488,17 @@ void CIRGenFunction::emitFunctionProlog(const FunctionArgList &args,
                    convertType(paramVar->getType()), paramLoc, alignment,
                    /*insertIntoFnEntryBlock=*/true);
 
-    declare(addrVal, paramVar, paramVar->getType(), paramLoc, alignment,
+    mlir::ptr::MemorySpaceAttrInterface destAddrSpace =
+        cir::toCIRAddressSpaceAttr(getMLIRContext(),
+                                   paramVar->getType().getAddressSpace());
+    Address addr = Address(addrVal, alignment);
+    addr = maybeCastStackAddressSpace(addr, destAddrSpace);
+
+    declare(addr.getPointer(), paramVar, paramVar->getType(), paramLoc,
+            alignment,
             /*isParam=*/true);
 
-    setAddrOfLocalVar(paramVar, Address(addrVal, alignment));
+    setAddrOfLocalVar(paramVar, addr);
 
     bool isPromoted = isa<ParmVarDecl>(paramVar) &&
                       cast<ParmVarDecl>(paramVar)->isKNRPromoted();

diff  --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index fe5ff90c25bef..8547a77e86faf 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -2548,7 +2548,7 @@ class CIRGenFunction : public CIRGenTypeCache {
   Address
   maybeCastStackAddressSpace(Address alloca,
                              mlir::ptr::MemorySpaceAttrInterface destAddrSpace,
-                             mlir::Value arraySize);
+                             mlir::Value arraySize = nullptr);
   Address createDefaultAlignTempAlloca(mlir::Type ty, mlir::Location loc,
                                        const Twine &name);
 

diff  --git a/clang/test/CIR/CodeGenHIP/ternary-struct-addrspace.hip b/clang/test/CIR/CodeGenHIP/ternary-struct-addrspace.hip
new file mode 100644
index 0000000000000..8dfbe163fa090
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/ternary-struct-addrspace.hip
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-cir %s -o - | FileCheck %s
+
+// Verify that a ternary expression returning a struct parameter from a
+// __device__ function produces a well-typed cir.ternary: both branch yields
+// and the result must all carry the same pointer type (no address-space
+// mismatch between the private-address-space alloca and the cast result).
+
+struct S { int a; };
+
+__attribute__((device)) bool cmp(S, S);
+
+__attribute__((device)) S choose(const S x) {
+  S y = x;
+  return cmp(x, y) ? x : y;
+}
+
+
+// CHECK: %[[X:.*]] = cir.alloca "x" {{.*}} : !cir.ptr<!rec_S, target_address_space(5)>
+// CHECK: %[[Y:.*]] = cir.alloca "y" {{.*}} : !cir.ptr<!rec_S, target_address_space(5)> loc(#loc18)
+
+// CHECK: %[[Y_CAST:.*]] = cir.cast address_space %[[Y]] : !cir.ptr<!rec_S, target_address_space(5)> -> !cir.ptr<!rec_S>
+// CHECK: %[[X_CAST:.*]] = cir.cast address_space %[[X]] : !cir.ptr<!rec_S, target_address_space(5)> -> !cir.ptr<!rec_S>
+
+// CHECK: %[[TERNARY:.*]] = cir.ternary({{.*}}, true {
+// CHECK:   cir.yield %[[X_CAST]] : !cir.ptr<!rec_S>
+// CHECK: }, false {
+// CHECK:   cir.yield %[[Y_CAST]] : !cir.ptr<!rec_S>
+// CHECK: }) : (!cir.bool) -> !cir.ptr<!rec_S>

diff  --git a/clang/test/CIR/CodeGenOpenMP/target-map.c b/clang/test/CIR/CodeGenOpenMP/target-map.c
index a4218a8c4a97d..11a00e3fd094d 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-map.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-map.c
@@ -89,9 +89,9 @@ void target_map_multiple(int a, int b) {
   // CIR-DEVICE: cir.func{{.*}}@target_map_multiple
   // CIR-DEVICE-DAG: %[[A_ALLOCA:.*]] = cir.alloca "a" align(4) init : !cir.ptr<!s32i, target_address_space(5)>
   // CIR-DEVICE-DAG: %[[B_ALLOCA:.*]] = cir.alloca "b" align(4) init : !cir.ptr<!s32i, target_address_space(5)>
+  // CIR-DEVICE: %[[CAST_B:.*]] = cir.cast address_space %[[B_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
   // CIR-DEVICE: %[[CAST_A:.*]] = cir.cast address_space %[[A_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
   // CIR-DEVICE: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[CAST_A]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) name("a") -> !cir.ptr<!s32i>
-  // CIR-DEVICE: %[[CAST_B:.*]] = cir.cast address_space %[[B_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i>
   // CIR-DEVICE: %[[MAP_B:.*]] = omp.map.info var_ptr(%[[CAST_B]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) name("b") -> !cir.ptr<!s32i>
   // CIR-DEVICE: omp.target kernel_type(generic) map_entries(%[[MAP_A]] -> %[[ARG_A:.*]], %[[MAP_B]] -> %[[ARG_B:.*]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
   // CIR-DEVICE: omp.terminator