[clang] [CIR] Cast alloca for function parameter to the right address space (PR #220836)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 3 00:59:17 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: Mariya Podchishchaeva (Fznamznon)
<details>
<summary>Changes</summary>
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.
Assited-by: claude in OCL test cases updating
---
Patch is 22.31 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/220836.diff
7 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.cpp (+13)
- (added) clang/test/CIR/CodeGenHIP/ternary-struct-addrspace.hip (+28)
- (modified) clang/test/CIR/CodeGenOpenCL/address-space-conversions.cl (+30-9)
- (modified) clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp (+3-2)
- (modified) clang/test/CIR/CodeGenOpenCL/as_type.cl (+12-8)
- (modified) clang/test/CIR/CodeGenOpenCL/vector.cl (+28-20)
- (modified) clang/test/CIR/CodeGenOpenMP/target-map.c (+1-1)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
index 8301627ad8123..a55ea643b7040 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -487,6 +487,19 @@ void CIRGenFunction::emitFunctionProlog(const FunctionArgList &args,
convertType(paramVar->getType()), paramLoc, alignment,
/*insertIntoFnEntryBlock=*/true);
+
+
+ mlir::ptr::MemorySpaceAttrInterface srcAddrSpace = getCIRAllocaAddressSpace();
+ mlir::ptr::MemorySpaceAttrInterface destAddrSpace =
+ cir::toCIRAddressSpaceAttr(getMLIRContext(),
+ paramVar->getType().getAddressSpace());
+ if (srcAddrSpace != destAddrSpace) {
+ mlir::Type destPtrTy = builder.getPointerTo(
+ (cast<cir::PointerType>(addrVal.getType())).getPointee(),
+ destAddrSpace);
+ addrVal = performAddrSpaceCast(addrVal, destPtrTy);
+ }
+
declare(addrVal, paramVar, paramVar->getType(), paramLoc, alignment,
/*isParam=*/true);
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/CodeGenOpenCL/address-space-conversions.cl b/clang/test/CIR/CodeGenOpenCL/address-space-conversions.cl
index 8b4f51196fdf1..ef37b194f1742 100644
--- a/clang/test/CIR/CodeGenOpenCL/address-space-conversions.cl
+++ b/clang/test/CIR/CodeGenOpenCL/address-space-conversions.cl
@@ -12,12 +12,33 @@ void address_space_conversions(global int *global_ptr,
}
// CIR-LABEL: cir.func dso_local @address_space_conversions
-// CIR: cir.cast address_space
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_global)>
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_generic)>
-// CIR: cir.cast address_space
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_private)>
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_generic)>
-// CIR: cir.cast address_space
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_generic)>
-// CIR-SAME: !cir.ptr<!s32i, lang_address_space(offload_global)>
+
+// Alloca slots for each parameter (outer type is pointer-to-pointer, no outer
+// address space yet — the inner pointer carries the OpenCL address space).
+// CIR: %[[GLOBAL_PTR_ALLOCA:.*]] = cir.alloca "global_ptr" {{.*}} : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>>
+// CIR: %[[GENERIC_PTR_ALLOCA:.*]] = cir.alloca "generic_ptr" {{.*}} : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>>
+// CIR: %[[PRIVATE_PTR_ALLOCA:.*]] = cir.alloca "private_ptr" {{.*}} : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_private)>>
+
+// Each alloca is immediately cast to add offload_private on the outer pointer.
+// All subsequent loads/stores go through these cast results.
+// CIR: %[[GLOBAL_SLOT:.*]] = cir.cast address_space %[[GLOBAL_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>> -> !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>, lang_address_space(offload_private)>
+// CIR: cir.store %arg0, %[[GLOBAL_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_global)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>, lang_address_space(offload_private)>
+// CIR: %[[GENERIC_SLOT:.*]] = cir.cast address_space %[[GENERIC_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>> -> !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>, lang_address_space(offload_private)>
+// CIR: cir.store %arg1, %[[GENERIC_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_generic)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>, lang_address_space(offload_private)>
+// CIR: %[[PRIVATE_SLOT:.*]] = cir.cast address_space %[[PRIVATE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_private)>> -> !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_private)>, lang_address_space(offload_private)>
+// CIR: cir.store %arg2, %[[PRIVATE_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_private)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_private)>, lang_address_space(offload_private)>
+
+// generic_ptr = global_ptr --> load global_ptr slot, cast global -> generic, store into generic_ptr slot
+// CIR: %[[GLOBAL_VAL:.*]] = cir.load {{.*}} %[[GLOBAL_SLOT]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>, lang_address_space(offload_private)>, !cir.ptr<!s32i, lang_address_space(offload_global)>
+// CIR: %[[GLOBAL_TO_GENERIC:.*]] = cir.cast address_space %[[GLOBAL_VAL]] : !cir.ptr<!s32i, lang_address_space(offload_global)> -> !cir.ptr<!s32i, lang_address_space(offload_generic)>
+// CIR: cir.store {{.*}} %[[GLOBAL_TO_GENERIC]], %[[GENERIC_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_generic)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>, lang_address_space(offload_private)>
+
+// generic_ptr = private_ptr --> load private_ptr slot, cast private -> generic, store into generic_ptr slot
+// CIR: %[[PRIVATE_VAL:.*]] = cir.load {{.*}} %[[PRIVATE_SLOT]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_private)>, lang_address_space(offload_private)>, !cir.ptr<!s32i, lang_address_space(offload_private)>
+// CIR: %[[PRIVATE_TO_GENERIC:.*]] = cir.cast address_space %[[PRIVATE_VAL]] : !cir.ptr<!s32i, lang_address_space(offload_private)> -> !cir.ptr<!s32i, lang_address_space(offload_generic)>
+// CIR: cir.store {{.*}} %[[PRIVATE_TO_GENERIC]], %[[GENERIC_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_generic)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>, lang_address_space(offload_private)>
+
+// global_ptr = (global int *)generic_ptr --> load generic_ptr slot, cast generic -> global, store into global_ptr slot
+// CIR: %[[GENERIC_VAL:.*]] = cir.load {{.*}} %[[GENERIC_SLOT]] : !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_generic)>, lang_address_space(offload_private)>, !cir.ptr<!s32i, lang_address_space(offload_generic)>
+// CIR: %[[GENERIC_TO_GLOBAL:.*]] = cir.cast address_space %[[GENERIC_VAL]] : !cir.ptr<!s32i, lang_address_space(offload_generic)> -> !cir.ptr<!s32i, lang_address_space(offload_global)>
+// CIR: cir.store {{.*}} %[[GENERIC_TO_GLOBAL]], %[[GLOBAL_SLOT]] : !cir.ptr<!s32i, lang_address_space(offload_global)>, !cir.ptr<!cir.ptr<!s32i, lang_address_space(offload_global)>, lang_address_space(offload_private)>
diff --git a/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp b/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
index bfce737207e87..8232f4b6304b7 100644
--- a/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
+++ b/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
@@ -14,8 +14,9 @@
// CIR: %[[R_ALLOCA:.*]] = cir.alloca "r" {{.*}} init const : !cir.ptr<!cir.ptr<!s32i>>
// CIR: %[[R:.*]] = cir.cast address_space %[[R_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>> -> !cir.ptr<!cir.ptr<!s32i>>
// CIR: %[[GR:.*]] = cir.cast address_space %[[GR_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>> -> !cir.ptr<!cir.ptr<!s32i>>
-// CIR: cir.store %arg0, %[[GP]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
-// CIR: %[[DEREF:.*]] = cir.load deref {{.*}} %[[GP]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
+// CIR: %[[GP_CAST:.*]] = cir.cast address_space %[[GP]] : !cir.ptr<!cir.ptr<!s32i>> -> !cir.ptr<!cir.ptr<!s32i>>
+// CIR: cir.store %arg0, %[[GP_CAST]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+// CIR: %[[DEREF:.*]] = cir.load deref {{.*}} %[[GP_CAST]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
// CIR: cir.store {{.*}} %[[DEREF]], %[[GR]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
// CIR: %[[GR_VAL:.*]] = cir.load %[[GR]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
// CIR: %[[CAST:.*]] = cir.cast address_space %[[GR_VAL]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i>
diff --git a/clang/test/CIR/CodeGenOpenCL/as_type.cl b/clang/test/CIR/CodeGenOpenCL/as_type.cl
index 303ebebc1f258..7cc1a03813546 100644
--- a/clang/test/CIR/CodeGenOpenCL/as_type.cl
+++ b/clang/test/CIR/CodeGenOpenCL/as_type.cl
@@ -17,8 +17,9 @@ char4 f4(int x) {
// CIR: cir.func {{.*}} @f4
// CIR: %[[X_ADDR:.*]] = cir.alloca "x" {{.*}} init : !cir.ptr<!s32i>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<4 x !s8i>>
-// CIR: cir.store %{{.*}}, %[[X_ADDR]] : !s32i, !cir.ptr<!s32i>
-// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_ADDR]] : !cir.ptr<!s32i>, !s32i
+// CIR: %[[X_CAST:.*]] = cir.cast address_space %[[X_ADDR]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i>
+// CIR: cir.store %{{.*}}, %[[X_CAST]] : !s32i, !cir.ptr<!s32i>
+// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_CAST]] : !cir.ptr<!s32i>, !s32i
// CIR: %[[X_V4_I8:.*]] = cir.cast bitcast %[[TMP_X]] : !s32i -> !cir.vector<4 x !s8i>
// CIR: cir.store %[[X_V4_I8]], %[[RET_ADDR]] : !cir.vector<4 x !s8i>, !cir.ptr<!cir.vector<4 x !s8i>>
// CIR: %[[TMP_RET:.*]] = cir.load %[[RET_ADDR]] : !cir.ptr<!cir.vector<4 x !s8i>>, !cir.vector<4 x !s8i>
@@ -35,8 +36,9 @@ int f6(char4 x) {
// CIR: cir.func {{.*}} @f6
// CIR: %[[X_ADDR:.*]] = cir.alloca "x" {{.*}} init : !cir.ptr<!cir.vector<4 x !s8i>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!s32i>
-// CIR: cir.store %{{.*}}, %[[X_ADDR]] : !cir.vector<4 x !s8i>, !cir.ptr<!cir.vector<4 x !s8i>>
-// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_ADDR]] : !cir.ptr<!cir.vector<4 x !s8i>>, !cir.vector<4 x !s8i>
+// CIR: %[[X_CAST:.*]] = cir.cast address_space %[[X_ADDR]] : !cir.ptr<!cir.vector<4 x !s8i>> -> !cir.ptr<!cir.vector<4 x !s8i>>
+// CIR: cir.store %{{.*}}, %[[X_CAST]] : !cir.vector<4 x !s8i>, !cir.ptr<!cir.vector<4 x !s8i>>
+// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_CAST]] : !cir.ptr<!cir.vector<4 x !s8i>>, !cir.vector<4 x !s8i>
// CIR: %[[X_S32I:.*]] = cir.cast bitcast %[[TMP_X]] : !cir.vector<4 x !s8i> -> !s32i
// CIR: cir.store %[[X_S32I]], %[[RET_ADDR]] : !s32i, !cir.ptr<!s32i>
// CIR: %[[TMP_RET:.*]] = cir.load %[[RET_ADDR]] : !cir.ptr<!s32i>, !s32i
@@ -53,8 +55,9 @@ int* int_to_ptr(int x) {
// CIR: cir.func {{.*}} @int_to_ptr
// CIR: %[[X_ADDR:.*]] = cir.alloca "x" {{.*}} init : !cir.ptr<!s32i>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.ptr<!s32i>>
-// CIR: cir.store %{{.*}}, %[[X_ADDR]] : !s32i, !cir.ptr<!s32i>
-// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_ADDR]] : !cir.ptr<!s32i>, !s32i
+// CIR: %[[X_CAST:.*]] = cir.cast address_space %[[X_ADDR]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i>
+// CIR: cir.store %{{.*}}, %[[X_CAST]] : !s32i, !cir.ptr<!s32i>
+// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_CAST]] : !cir.ptr<!s32i>, !s32i
// CIR: %[[X_PTR:.*]] = cir.cast int_to_ptr %[[TMP_X]] : !s32i -> !cir.ptr<!s32i>
// CIR: cir.store %[[X_PTR]], %[[RET_ADDR]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
// CIR: %[[TMP_RET]] = cir.load %[[RET_ADDR]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
@@ -71,8 +74,9 @@ char3 vec4_to_vec_3(char4 x) {
// CIR: cir.func {{.*}} @vec4_to_vec_3
// CIR: %[[X_ADDR:.*]] = cir.alloca "x" {{.*}} init : !cir.ptr<!cir.vector<4 x !s8i>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<3 x !s8i>>
-// CIR: cir.store %{{.*}}, %[[X_ADDR]] : !cir.vector<4 x !s8i>, !cir.ptr<!cir.vector<4 x !s8i>>
-// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_ADDR]] : !cir.ptr<!cir.vector<4 x !s8i>>, !cir.vector<4 x !s8i>
+// CIR: %[[X_CAST:.*]] = cir.cast address_space %[[X_ADDR]] : !cir.ptr<!cir.vector<4 x !s8i>> -> !cir.ptr<!cir.vector<4 x !s8i>>
+// CIR: cir.store %{{.*}}, %[[X_CAST]] : !cir.vector<4 x !s8i>, !cir.ptr<!cir.vector<4 x !s8i>>
+// CIR: %[[TMP_X:.*]] = cir.load {{.*}} %[[X_CAST]] : !cir.ptr<!cir.vector<4 x !s8i>>, !cir.vector<4 x !s8i>
// CIR: %[[POISON:.*]] = cir.const #cir.poison : !cir.vector<4 x !s8i>
// CIR: %[[RESULT:.*]] = cir.vec.shuffle(%[[TMP_X]], %[[POISON]] : !cir.vector<4 x !s8i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<2> : !s32i] : !cir.vector<3 x !s8i>
// CIR: cir.store %[[RESULT]], %[[RET_ADDR]] : !cir.vector<3 x !s8i>, !cir.ptr<!cir.vector<3 x !s8i>>
diff --git a/clang/test/CIR/CodeGenOpenCL/vector.cl b/clang/test/CIR/CodeGenOpenCL/vector.cl
index 4013449cb49b2..baa58419508c1 100644
--- a/clang/test/CIR/CodeGenOpenCL/vector.cl
+++ b/clang/test/CIR/CodeGenOpenCL/vector.cl
@@ -18,12 +18,15 @@ int4 vec_ternary(int4 c, int4 a, int4 b) {
// CIR: %[[A_ADDR:.*]] = cir.alloca "a" {{.*}} init : !cir.ptr<!cir.vector<4 x !s32i>>
// CIR: %[[B_ADDR:.*]] = cir.alloca "b" {{.*}} init : !cir.ptr<!cir.vector<4 x !s32i>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[C_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[A_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[B_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: %[[TMP_C:.*]] = cir.load {{.*}} %[[C_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
-// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
-// CIR: %[[TMP_C_2:.*]] = cir.load {{.*}} %[[C_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[C_CAST:.*]] = cir.cast address_space %[[C_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: cir.store %{{.*}}, %[[C_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[A_CAST:.*]] = cir.cast address_space %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: cir.store %{{.*}}, %[[A_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[B_CAST:.*]] = cir.cast address_space %[[B_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: cir.store %{{.*}}, %[[B_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[TMP_C:.*]] = cir.load {{.*}} %[[C_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[TMP_C_2:.*]] = cir.load {{.*}} %[[C_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
// CIR: %[[CONST_0_VEC:.*]] = cir.const #cir.zero : !cir.vector<4 x !s32i>
// CIR: %[[C_CMP:.*]] = cir.vec.cmp(lt, %[[TMP_C]], %[[CONST_0_VEC]]) : !cir.vector<4 x !s32i>, !cir.vector<4 x !s32i>
// CIR: %[[C_CMP_NOT:.*]] = cir.not %[[C_CMP]] : !cir.vector<4 x !s32i>
@@ -48,12 +51,15 @@ float4 vec_ternary_f4(int4 c, float4 a, float4 b) {
// CIR: %[[A_ADDR:.*]] = cir.alloca "a" {{.*}} init : !cir.ptr<!cir.vector<4 x !cir.float>>
// CIR: %[[B_ADDR:.*]] = cir.alloca "b" {{.*}} init : !cir.ptr<!cir.vector<4 x !cir.float>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<4 x !cir.float>>
-// CIR: cir.store %{{.*}}, %[[C_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[A_ADDR]] : !cir.vector<4 x !cir.float>, !cir.ptr<!cir.vector<4 x !cir.float>>
-// CIR: cir.store %{{.*}}, %[[B_ADDR]] : !cir.vector<4 x !cir.float>, !cir.ptr<!cir.vector<4 x !cir.float>>
-// CIR: %[[TMP_C:.*]] = cir.load {{.*}} %[[C_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
-// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !cir.float>>, !cir.vector<4 x !cir.float>
-// CIR: %[[TMP_B:.*]] = cir.load {{.*}} %[[B_ADDR]] : !cir.ptr<!cir.vector<4 x !cir.float>>, !cir.vector<4 x !cir.float>
+// CIR: %[[C_CAST:.*]] = cir.cast address_space %[[C_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: cir.store %{{.*}}, %[[C_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[A_CAST:.*]] = cir.cast address_space %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !cir.float>> -> !cir.ptr<!cir.vector<4 x !cir.float>>
+// CIR: cir.store %{{.*}}, %[[A_CAST]] : !cir.vector<4 x !cir.float>, !cir.ptr<!cir.vector<4 x !cir.float>>
+// CIR: %[[B_CAST:.*]] = cir.cast address_space %[[B_ADDR]] : !cir.ptr<!cir.vector<4 x !cir.float>> -> !cir.ptr<!cir.vector<4 x !cir.float>>
+// CIR: cir.store %{{.*}}, %[[B_CAST]] : !cir.vector<4 x !cir.float>, !cir.ptr<!cir.vector<4 x !cir.float>>
+// CIR: %[[TMP_C:.*]] = cir.load {{.*}} %[[C_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_CAST]] : !cir.ptr<!cir.vector<4 x !cir.float>>, !cir.vector<4 x !cir.float>
+// CIR: %[[TMP_B:.*]] = cir.load {{.*}} %[[B_CAST]] : !cir.ptr<!cir.vector<4 x !cir.float>>, !cir.vector<4 x !cir.float>
// CIR: %[[CONST_0_VEC:.*]] = cir.const #cir.zero : !cir.vector<4 x !s32i>
// CIR: %[[C_CMP:.*]] = cir.vec.cmp(lt, %[[TMP_C]], %[[CONST_0_VEC]]) : !cir.vector<4 x !s32i>, !cir.vector<4 x !s32i>
// CIR: %[[C_CMP_NOT:.*]] = cir.not %[[C_CMP]] : !cir.vector<4 x !s32i>
@@ -78,11 +84,12 @@ int4 vec_unary_inc(int4 a) {
// CIR: %[[A_ADDR:.*]] = cir.alloca "a" {{.*}} init : !cir.ptr<!cir.vector<4 x !s32i>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[A_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[A_CAST:.*]] = cir.cast address_space %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: cir.store %{{.*}}, %[[A_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
// CIR: %[[VEC_INC:.*]] = cir.inc %[[TMP_A]] : !cir.vector<4 x !s32i>
-// CIR: cir.store {{.*}} %[[VEC_INC]], %[[A_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: %[[RESULT:.*]] = cir.load {{.*}} %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: cir.store {{.*}} %[[VEC_INC]], %[[A_CAST]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
+// CIR: %[[RESULT:.*]] = cir.load {{.*}} %[[A_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
// CIR: cir.store %[[RESULT]], %[[RET_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
// LLVM: %[[VEC_INC:.*]] = add <4 x i32> %[[A:.*]], splat (i32 1)
@@ -95,11 +102,12 @@ int4 vec_unary_dec(int4 a) {
// CIR: %[[A_ADDR:.*]] = cir.alloca "a" {{.*}} init : !cir.ptr<!cir.vector<4 x !s32i>>
// CIR: %[[RET_ADDR:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: cir.store %{{.*}}, %[[A_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>>
-// CIR: %[[TMP_A:.*]] = cir.load {{.*}} %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i>
+// CIR: %[[A_CAST:.*]] = cir.cast address_space %[[A_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/220836
More information about the cfe-commits
mailing list