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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
