https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/177623
Reverts /llvm/llvm-project#177059 and llvm/llvm-project/#177491. The new test may have an issue. >From d41c3ce0b22ce5673a2f8c47311e1325b550291f Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Fri, 23 Jan 2026 09:27:28 -0800 Subject: [PATCH 1/2] Revert "[NFC][OpenMP] Mark new mapper test as XFAIL on intelgpu. (#177491)" This reverts commit 7d5622f7917815d224b780309432ffe4729e4852. --- offload/test/mapping/declare_mapper_target_checks.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp index 5133d2f6abfef..1c5edae99a7c8 100644 --- a/offload/test/mapping/declare_mapper_target_checks.cpp +++ b/offload/test/mapping/declare_mapper_target_checks.cpp @@ -1,6 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu - #include <omp.h> #include <stdio.h> >From 66aa65f0d1ea8f12104b65d58a0e64d7cf71a9cd Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Fri, 23 Jan 2026 09:28:41 -0800 Subject: [PATCH 2/2] Revert "[OpenMP][Mappers] Fix ref-count tracking for maps inserted by mappers. (#177059)" This reverts commit a7b57cf7bba13cd7851bca1f2f3a820c74110c0d. --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 +- clang/test/OpenMP/declare_mapper_codegen.cpp | 25 +- ..._of_structs_with_nested_mapper_codegen.cpp | 220 +++++++------- ..._of_structs_with_nested_mapper_codegen.cpp | 286 +++++++++--------- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 8 + mlir/test/Target/LLVMIR/omptarget-llvm.mlir | 5 +- offload/libomptarget/omptarget.cpp | 8 +- .../mapping/declare_mapper_target_checks.cpp | 164 ---------- 8 files changed, 300 insertions(+), 422 deletions(-) delete mode 100644 offload/test/mapping/declare_mapper_target_checks.cpp diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index e21d67010521a..8981a0de6d0e4 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -10465,8 +10465,10 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { /// void *base, void *begin, /// int64_t size, int64_t type, /// void *name = nullptr) { -/// // Allocate space for an array section first. -/// if ((size > 1 || (base != begin)) && !maptype.IsDelete) +/// // Allocate space for an array section first or add a base/begin for +/// // pointer dereference. +/// if ((size > 1 || (base != begin && maptype.IsPtrAndObj)) && +/// !maptype.IsDelete) /// __tgt_push_mapper_component(rt_mapper_handle, base, begin, /// size*sizeof(Ty), clearToFromMember(type)); /// // Map members. diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index 69f9dae7d4988..7dc32d0ae12ff 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -92,7 +92,10 @@ class C { // CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK0-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK0-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 +// CK0-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 +// CK0-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] +// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] // CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK0-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -589,7 +592,10 @@ class C { // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK1-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK1-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 +// CK1-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 +// CK1-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] +// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] // CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK1-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -696,7 +702,10 @@ class C { // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK2-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK2-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 +// CK2-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 +// CK2-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] +// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] // CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK2-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -888,7 +897,10 @@ class C { // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK4-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK4-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 +// CK4-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 +// CK4-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] +// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] // CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -1075,7 +1087,10 @@ void foo(){ // CK5-DAG: [[PTREND:%.+]] = getelementptr %struct.myvec, ptr [[BEGIN]], i64 [[SIZE]] // CK5-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK5-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK5-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 +// CK5-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 +// CK5-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] +// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] // CK5-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK5-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK5-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp index 054e8f22633cb..5df1e958ad55a 100644 --- a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp +++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp @@ -37,15 +37,15 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@_Z3foov // CHECK-SAME: () #[[ATTR0:[0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[SA:%.*]] = alloca [10 x [[STRUCT_D:%.*]]], align 4 +// CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 1 -// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 +// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 // CHECK-NEXT: store i32 111, ptr [[E]], align 4 -// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 1 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 // CHECK-NEXT: store i32 222, ptr [[A]], align 4 @@ -99,10 +99,10 @@ void foo() { // CHECK-NEXT: [[SA_ADDR:%.*]] = alloca ptr, align 8 // CHECK-NEXT: store ptr [[SA]], ptr [[SA_ADDR]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8, !nonnull [[META5:![0-9]+]], !align [[META6:![0-9]+]] -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [[STRUCT_D:%.*]]], ptr [[TMP0]], i64 0, i64 1 -// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 +// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 // CHECK-NEXT: store i32 333, ptr [[E]], align 4 -// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[TMP0]], i64 0, i64 1 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 // CHECK-NEXT: store i32 444, ptr [[A]], align 4 @@ -117,15 +117,18 @@ void foo() { // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 // CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8 // CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK-NEXT: [[TMP10:%.*]] = and i64 [[TMP4]], 16 +// CHECK-NEXT: [[TMP11:%.*]] = icmp ne i64 [[TMP10]], 0 +// CHECK-NEXT: [[TMP12:%.*]] = and i1 [[TMP9]], [[TMP11]] +// CHECK-NEXT: [[TMP13:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP12]] // CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 -// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-NEXT: [[TMP14:%.*]] = and i1 [[TMP13]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP14]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP15:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP17:%.*]] = or i64 [[TMP16]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP15]], i64 [[TMP17]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] @@ -135,115 +138,115 @@ void foo() { // CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 // CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2 -// CHECK-NEXT: [[TMP15:%.*]] = getelementptr i32, ptr [[H]], i32 1 -// CHECK-NEXT: [[TMP16:%.*]] = ptrtoint ptr [[TMP15]] to i64 -// CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[E]] to i64 -// CHECK-NEXT: [[TMP18:%.*]] = sub i64 [[TMP16]], [[TMP17]] -// CHECK-NEXT: [[TMP19:%.*]] = sdiv exact i64 [[TMP18]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK-NEXT: [[TMP20:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP21:%.*]] = shl i64 [[TMP20]], 48 -// CHECK-NEXT: [[TMP22:%.*]] = add nuw i64 0, [[TMP21]] -// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP23]], 0 -// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[H]], i32 1 +// CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP18]] to i64 +// CHECK-NEXT: [[TMP20:%.*]] = ptrtoint ptr [[E]] to i64 +// CHECK-NEXT: [[TMP21:%.*]] = sub i64 [[TMP19]], [[TMP20]] +// CHECK-NEXT: [[TMP22:%.*]] = sdiv exact i64 [[TMP21]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK-NEXT: [[TMP23:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP24:%.*]] = shl i64 [[TMP23]], 48 +// CHECK-NEXT: [[TMP25:%.*]] = add nuw i64 0, [[TMP24]] +// CHECK-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP27:%.*]] = icmp eq i64 [[TMP26]], 0 +// CHECK-NEXT: br i1 [[TMP27]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP22]], -4 +// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP25]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END:%.*]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP23]], 1 -// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP29:%.*]] = icmp eq i64 [[TMP26]], 1 +// CHECK-NEXT: br i1 [[TMP29]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP22]], -3 +// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP25]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP23]], 2 -// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP26]], 2 +// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP22]], -2 +// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP25]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP25]], [[OMP_TYPE_ALLOC]] ], [ [[TMP27]], [[OMP_TYPE_TO]] ], [ [[TMP29]], [[OMP_TYPE_FROM]] ], [ [[TMP22]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP19]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0 -// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP28]], [[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]], [[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP22]], i64 [[OMP_MAPTYPE]], ptr null) +// CHECK-NEXT: [[TMP33:%.*]] = add nuw i64 281474976711171, [[TMP24]] +// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0 +// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] // CHECK: omp.type.alloc1: -// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4 +// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP33]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]] // CHECK: omp.type.alloc.else2: -// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1 -// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]] +// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 1 +// CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]] // CHECK: omp.type.to3: -// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3 +// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP33]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.to.else4: -// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2 -// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]] +// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP34]], 2 +// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]] // CHECK: omp.type.from5: -// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2 +// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP33]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.end6: -// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ] +// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP36]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP38]], [[OMP_TYPE_TO3]] ], [ [[TMP40]], [[OMP_TYPE_FROM5]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE4]] ] // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null) -// CHECK-NEXT: [[TMP38:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP39:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP40:%.*]] = icmp eq i64 [[TMP39]], 0 -// CHECK-NEXT: br i1 [[TMP40]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] +// CHECK-NEXT: [[TMP41:%.*]] = add nuw i64 281474976711171, [[TMP24]] +// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0 +// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] // CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP41:%.*]] = and i64 [[TMP38]], -4 +// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP41]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]] // CHECK: omp.type.alloc.else9: -// CHECK-NEXT: [[TMP42:%.*]] = icmp eq i64 [[TMP39]], 1 -// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] +// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1 +// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] // CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP43:%.*]] = and i64 [[TMP38]], -3 +// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP41]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END13]] // CHECK: omp.type.to.else11: -// CHECK-NEXT: [[TMP44:%.*]] = icmp eq i64 [[TMP39]], 2 -// CHECK-NEXT: br i1 [[TMP44]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] +// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2 +// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] // CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP45:%.*]] = and i64 [[TMP38]], -2 +// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP41]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END13]] // CHECK: omp.type.end13: -// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP41]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP43]], [[OMP_TYPE_TO10]] ], [ [[TMP45]], [[OMP_TYPE_FROM12]] ], [ [[TMP38]], [[OMP_TYPE_TO_ELSE11]] ] +// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP46]], [[OMP_TYPE_TO10]] ], [ [[TMP48]], [[OMP_TYPE_FROM12]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE11]] ] // CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null) #[[ATTR3]] -// CHECK-NEXT: [[TMP46:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP48:%.*]] = icmp eq i64 [[TMP47]], 0 -// CHECK-NEXT: br i1 [[TMP48]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] +// CHECK-NEXT: [[TMP49:%.*]] = add nuw i64 281474976711171, [[TMP24]] +// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP50]], 0 +// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] // CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP46]], -4 +// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP49]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.alloc.else16: -// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP47]], 1 -// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] +// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP50]], 1 +// CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] // CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP46]], -3 +// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP49]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.to.else18: -// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP47]], 2 -// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] +// CHECK-NEXT: [[TMP55:%.*]] = icmp eq i64 [[TMP50]], 2 +// CHECK-NEXT: br i1 [[TMP55]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] // CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP46]], -2 +// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP49]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.end20: -// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP49]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP51]], [[OMP_TYPE_TO17]] ], [ [[TMP53]], [[OMP_TYPE_FROM19]] ], [ [[TMP46]], [[OMP_TYPE_TO_ELSE18]] ] +// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP52]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP54]], [[OMP_TYPE_TO17]] ], [ [[TMP56]], [[OMP_TYPE_FROM19]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE18]] ] // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], ptr null) // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 // CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK: omp.arraymap.exit: // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP54]], 0 -// CHECK-NEXT: [[TMP55:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP55]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP57]], 0 +// CHECK-NEXT: [[TMP58:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP56:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP58:%.*]] = or i64 [[TMP57]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP56]], i64 [[TMP58]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP59:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP61:%.*]] = or i64 [[TMP60]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP59]], i64 [[TMP61]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void @@ -257,15 +260,18 @@ void foo() { // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 // CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8 // CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK-NEXT: [[TMP10:%.*]] = and i64 [[TMP4]], 16 +// CHECK-NEXT: [[TMP11:%.*]] = icmp ne i64 [[TMP10]], 0 +// CHECK-NEXT: [[TMP12:%.*]] = and i1 [[TMP9]], [[TMP11]] +// CHECK-NEXT: [[TMP13:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP12]] // CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 -// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-NEXT: [[TMP14:%.*]] = and i1 [[TMP13]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP14]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 4 -// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP15:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP17:%.*]] = or i64 [[TMP16]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP15]], i64 [[TMP17]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] @@ -273,44 +279,44 @@ void foo() { // CHECK: omp.arraymap.body: // CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ] // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 -// CHECK-NEXT: [[TMP15:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP16:%.*]] = shl i64 [[TMP15]], 48 -// CHECK-NEXT: [[TMP17:%.*]] = add nuw i64 1, [[TMP16]] -// CHECK-NEXT: [[TMP18:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP19:%.*]] = icmp eq i64 [[TMP18]], 0 -// CHECK-NEXT: br i1 [[TMP19]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP18:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP19:%.*]] = shl i64 [[TMP18]], 48 +// CHECK-NEXT: [[TMP20:%.*]] = add nuw i64 1, [[TMP19]] +// CHECK-NEXT: [[TMP21:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP22:%.*]] = icmp eq i64 [[TMP21]], 0 +// CHECK-NEXT: br i1 [[TMP22]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP20:%.*]] = and i64 [[TMP17]], -4 +// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP20]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP18]], 1 -// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP21]], 1 +// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP17]], -3 +// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP20]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP23:%.*]] = icmp eq i64 [[TMP18]], 2 -// CHECK-NEXT: br i1 [[TMP23]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP21]], 2 +// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP24:%.*]] = and i64 [[TMP17]], -2 +// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP20]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP23]], [[OMP_TYPE_ALLOC]] ], [ [[TMP25]], [[OMP_TYPE_TO]] ], [ [[TMP27]], [[OMP_TYPE_FROM]] ], [ [[TMP20]], [[OMP_TYPE_TO_ELSE]] ] // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null) // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 // CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK: omp.arraymap.exit: // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0 -// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP26]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP28]], 0 +// CHECK-NEXT: [[TMP29:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP29]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4 -// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP30:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP32:%.*]] = or i64 [[TMP31]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP30]], i64 [[TMP32]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void diff --git a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp index a017fb098ced9..0fc6de0e4279a 100644 --- a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp +++ b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp @@ -33,19 +33,19 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@_Z3foov // CHECK-SAME: () #[[ATTR0:[0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[SA:%.*]] = alloca [10 x [[STRUCT_D:%.*]]], align 4 +// CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 1 -// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 +// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 // CHECK-NEXT: store i32 111, ptr [[E]], align 4 -// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 1 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 // CHECK-NEXT: store i32 222, ptr [[A]], align 4 -// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 0 +// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [10 x %struct.D], ptr [[SA]], i64 0, i64 0 // CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 // CHECK-NEXT: store ptr [[SA]], ptr [[TMP0]], align 8 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 @@ -95,11 +95,11 @@ void foo() { // CHECK-NEXT: entry: // CHECK-NEXT: [[SA_ADDR:%.*]] = alloca ptr, align 8 // CHECK-NEXT: store ptr [[SA]], ptr [[SA_ADDR]], align 8 -// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8, !nonnull [[META5:![0-9]+]], !align [[META6:![0-9]+]] -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [[STRUCT_D:%.*]]], ptr [[TMP0]], i64 0, i64 1 -// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX]], i32 0, i32 0 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 +// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 // CHECK-NEXT: store i32 333, ptr [[E]], align 4 -// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [[STRUCT_D]]], ptr [[TMP0]], i64 0, i64 1 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 // CHECK-NEXT: store i32 444, ptr [[A]], align 4 @@ -109,138 +109,141 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1D.default // CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2:[0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 12 -// CHECK-NEXT: [[TMP7:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP2]], i64 [[TMP6]] -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] -// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 -// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 12 +// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP2]], i64 [[TMP10]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 +// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] +// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 +// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 +// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] +// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] +// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 +// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 12 +// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: -// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] +// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] // CHECK: omp.arraymap.body: -// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END20:%.*]] ] +// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ] // CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 // CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2 -// CHECK-NEXT: [[TMP15:%.*]] = getelementptr i32, ptr [[H]], i32 1 -// CHECK-NEXT: [[TMP16:%.*]] = ptrtoint ptr [[TMP15]] to i64 -// CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[E]] to i64 -// CHECK-NEXT: [[TMP18:%.*]] = sub i64 [[TMP16]], [[TMP17]] -// CHECK-NEXT: [[TMP19:%.*]] = sdiv exact i64 [[TMP18]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK-NEXT: [[TMP20:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP21:%.*]] = shl i64 [[TMP20]], 48 -// CHECK-NEXT: [[TMP22:%.*]] = add nuw i64 0, [[TMP21]] -// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP23]], 0 -// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP24:%.*]] = getelementptr i32, ptr [[H]], i32 1 +// CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64 +// CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[E]] to i64 +// CHECK-NEXT: [[TMP27:%.*]] = sub i64 [[TMP25]], [[TMP26]] +// CHECK-NEXT: [[TMP28:%.*]] = sdiv exact i64 [[TMP27]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK-NEXT: [[TMP29:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP30:%.*]] = shl i64 [[TMP29]], 48 +// CHECK-NEXT: [[TMP31:%.*]] = add nuw i64 0, [[TMP30]] +// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 +// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP22]], -4 +// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP31]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END:%.*]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP23]], 1 -// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 +// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP22]], -3 +// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP31]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP23]], 2 -// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 +// CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP22]], -2 +// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP31]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP25]], [[OMP_TYPE_ALLOC]] ], [ [[TMP27]], [[OMP_TYPE_TO]] ], [ [[TMP29]], [[OMP_TYPE_FROM]] ], [ [[TMP22]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP19]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0 -// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC]] ], [ [[TMP36]], [[OMP_TYPE_TO]] ], [ [[TMP38]], [[OMP_TYPE_FROM]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP28]], i64 [[OMP_MAPTYPE]], ptr null) +// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP30]] +// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 +// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC6:%.*]], label [[OMP_TYPE_ALLOC_ELSE7:%.*]] // CHECK: omp.type.alloc1: -// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]] +// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END11:%.*]] // CHECK: omp.type.alloc.else2: -// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1 -// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]] +// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 +// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO8:%.*]], label [[OMP_TYPE_TO_ELSE9:%.*]] // CHECK: omp.type.to3: -// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END6]] +// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END11]] // CHECK: omp.type.to.else4: -// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2 -// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]] +// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 +// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM10:%.*]], label [[OMP_TYPE_END11]] // CHECK: omp.type.from5: -// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END6]] +// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END11]] // CHECK: omp.type.end6: -// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null) -// CHECK-NEXT: [[TMP38:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP39:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP40:%.*]] = icmp eq i64 [[TMP39]], 0 -// CHECK-NEXT: br i1 [[TMP40]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC6]] ], [ [[TMP44]], [[OMP_TYPE_TO8]] ], [ [[TMP46]], [[OMP_TYPE_FROM10]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE9]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE12]], ptr null) +// CHECK-NEXT: [[TMP47:%.*]] = add nuw i64 281474976711171, [[TMP30]] +// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 +// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC13:%.*]], label [[OMP_TYPE_ALLOC_ELSE14:%.*]] // CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP41:%.*]] = and i64 [[TMP38]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]] +// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP47]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END18:%.*]] // CHECK: omp.type.alloc.else9: -// CHECK-NEXT: [[TMP42:%.*]] = icmp eq i64 [[TMP39]], 1 -// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] +// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 +// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_TO15:%.*]], label [[OMP_TYPE_TO_ELSE16:%.*]] // CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP43:%.*]] = and i64 [[TMP38]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] +// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP47]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END18]] // CHECK: omp.type.to.else11: -// CHECK-NEXT: [[TMP44:%.*]] = icmp eq i64 [[TMP39]], 2 -// CHECK-NEXT: br i1 [[TMP44]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] +// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 +// CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_FROM17:%.*]], label [[OMP_TYPE_END18]] // CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP45:%.*]] = and i64 [[TMP38]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] +// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP47]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END18]] // CHECK: omp.type.end13: -// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP41]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP43]], [[OMP_TYPE_TO10]] ], [ [[TMP45]], [[OMP_TYPE_FROM12]] ], [ [[TMP38]], [[OMP_TYPE_TO_ELSE11]] ] -// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null) #[[ATTR3]] -// CHECK-NEXT: [[TMP46:%.*]] = add nuw i64 281474976711171, [[TMP21]] -// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP48:%.*]] = icmp eq i64 [[TMP47]], 0 -// CHECK-NEXT: br i1 [[TMP48]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC13]] ], [ [[TMP52]], [[OMP_TYPE_TO15]] ], [ [[TMP54]], [[OMP_TYPE_FROM17]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE16]] ] +// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE19]], ptr null) #[[ATTR3]] +// CHECK-NEXT: [[TMP55:%.*]] = add nuw i64 281474976711171, [[TMP30]] +// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP57:%.*]] = icmp eq i64 [[TMP56]], 0 +// CHECK-NEXT: br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]] // CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP46]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] +// CHECK-NEXT: [[TMP58:%.*]] = and i64 [[TMP55]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END25]] // CHECK: omp.type.alloc.else16: -// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP47]], 1 -// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] +// CHECK-NEXT: [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1 +// CHECK-NEXT: br i1 [[TMP59]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]] // CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP46]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] +// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP55]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END25]] // CHECK: omp.type.to.else18: -// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP47]], 2 -// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] +// CHECK-NEXT: [[TMP61:%.*]] = icmp eq i64 [[TMP56]], 2 +// CHECK-NEXT: br i1 [[TMP61]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]] // CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP46]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] +// CHECK-NEXT: [[TMP62:%.*]] = and i64 [[TMP55]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END25]] // CHECK: omp.type.end20: -// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP49]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP51]], [[OMP_TYPE_TO17]] ], [ [[TMP53]], [[OMP_TYPE_FROM19]] ], [ [[TMP46]], [[OMP_TYPE_TO_ELSE18]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], ptr null) +// CHECK-NEXT: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP60]], [[OMP_TYPE_TO22]] ], [ [[TMP62]], [[OMP_TYPE_FROM24]] ], [ [[TMP55]], [[OMP_TYPE_TO_ELSE23]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE26]], ptr null) // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 -// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] +// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK: omp.arraymap.exit: -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP54]], 0 -// CHECK-NEXT: [[TMP55:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP55]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY27:%.*]] = icmp sgt i64 [[TMP10]], 1 +// CHECK-NEXT: [[TMP63:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP63]], 0 +// CHECK-NEXT: [[TMP64:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY27]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP64]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP56:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP58:%.*]] = or i64 [[TMP57]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP56]], i64 [[TMP58]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP65:%.*]] = mul nuw i64 [[TMP10]], 12 +// CHECK-NEXT: [[TMP66:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP67:%.*]] = or i64 [[TMP66]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP65]], i64 [[TMP67]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void @@ -249,65 +252,68 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default // CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 4 -// CHECK-NEXT: [[TMP7:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP2]], i64 [[TMP6]] -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] -// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 -// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 4 +// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP2]], i64 [[TMP10]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 +// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] +// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 +// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 +// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] +// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] +// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 +// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 4 -// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 4 +// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: -// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] +// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] // CHECK: omp.arraymap.body: // CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ] // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 -// CHECK-NEXT: [[TMP15:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP16:%.*]] = shl i64 [[TMP15]], 48 -// CHECK-NEXT: [[TMP17:%.*]] = add nuw i64 1, [[TMP16]] -// CHECK-NEXT: [[TMP18:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP19:%.*]] = icmp eq i64 [[TMP18]], 0 -// CHECK-NEXT: br i1 [[TMP19]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP24:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP25:%.*]] = shl i64 [[TMP24]], 48 +// CHECK-NEXT: [[TMP26:%.*]] = add nuw i64 1, [[TMP25]] +// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP27]], 0 +// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP20:%.*]] = and i64 [[TMP17]], -4 +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP26]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP18]], 1 -// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1 +// CHECK-NEXT: br i1 [[TMP30]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP17]], -3 +// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP26]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP23:%.*]] = icmp eq i64 [[TMP18]], 2 -// CHECK-NEXT: br i1 [[TMP23]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP27]], 2 +// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP24:%.*]] = and i64 [[TMP17]], -2 +// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP26]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP29]], [[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]], [[OMP_TYPE_FROM]] ], [ [[TMP26]], [[OMP_TYPE_TO_ELSE]] ] // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null) // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 -// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] +// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]] // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK: omp.arraymap.exit: -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0 -// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP26]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY6:%.*]] = icmp sgt i64 [[TMP10]], 1 +// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP34]], 0 +// CHECK-NEXT: [[TMP35:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY6]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP35]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4 -// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP36:%.*]] = mul nuw i64 [[TMP10]], 4 +// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP38:%.*]] = or i64 [[TMP37]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP36]], i64 [[TMP38]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 6b02de855fc66..5071fc53b43cc 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9465,6 +9465,14 @@ void OpenMPIRBuilder::emitUDMapperArrayInitOrDel( if (IsInit) { // base != begin? Value *BaseIsBegin = Builder.CreateICmpNE(Base, Begin); + // IsPtrAndObj? + Value *PtrAndObjBit = Builder.CreateAnd( + MapType, + Builder.getInt64( + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ))); + PtrAndObjBit = Builder.CreateIsNotNull(PtrAndObjBit); + BaseIsBegin = Builder.CreateAnd(BaseIsBegin, PtrAndObjBit); Cond = Builder.CreateOr(IsArray, BaseIsBegin); DeleteCond = Builder.CreateIsNull( DeleteBit, diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir index 0548c6a178d0c..0b4d63125f82f 100644 --- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir +++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir @@ -564,7 +564,10 @@ module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} { // CHECK: %[[VAL_20:.*]] = icmp sgt i64 %[[VAL_15]], 1 // CHECK: %[[VAL_21:.*]] = and i64 %[[VAL_22:.*]], 8 // CHECK: %[[VAL_23:.*]] = icmp ne ptr %[[VAL_24:.*]], %[[VAL_19]] -// CHECK: %[[VAL_28:.*]] = or i1 %[[VAL_20]], %[[VAL_23]] +// CHECK: %[[VAL_25:.*]] = and i64 %[[VAL_22]], 16 +// CHECK: %[[VAL_26:.*]] = icmp ne i64 %[[VAL_25]], 0 +// CHECK: %[[VAL_27:.*]] = and i1 %[[VAL_23]], %[[VAL_26]] +// CHECK: %[[VAL_28:.*]] = or i1 %[[VAL_20]], %[[VAL_27]] // CHECK: %[[VAL_29:.*]] = icmp eq i64 %[[VAL_21]], 0 // CHECK: %[[VAL_30:.*]] = and i1 %[[VAL_28]], %[[VAL_29]] // CHECK: br i1 %[[VAL_30]], label %[[VAL_31:.*]], label %[[VAL_32:.*]] diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index bd99edee5e1b3..1723048db7fa2 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -600,7 +600,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, // then no argument is marked as TARGET_PARAM ("omp target data map" is not // associated with a target region, so there are no target parameters). This // may be considered a hack, we could revise the scheme in the future. - bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF); + bool UpdateRef = + !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0); MappingInfoTy::HDTTMapAccessorTy HDTTMap = Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor(); @@ -1109,8 +1110,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, void *HstPtrBegin = Args[I]; int64_t DataSize = ArgSizes[I]; bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; - bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); + bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || + (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) && + !(FromMapper && I == 0); bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE; bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD; diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp deleted file mode 100644 index 1c5edae99a7c8..0000000000000 --- a/offload/test/mapping/declare_mapper_target_checks.cpp +++ /dev/null @@ -1,164 +0,0 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic -#include <omp.h> -#include <stdio.h> - -#define TRUE 1 -#define FALSE 0 - -struct TY1 { - int i1, i2, i3; - static constexpr auto name = "TY1"; -}; -struct TY2 { - int i1, i2, i3; - static constexpr auto name = "TY2"; -}; - -// TY1 is not mapped, TY2 is -#pragma omp declare mapper(TY2 t) map(to : t.i1) map(from : t.i3) - -struct TY3 { - TY2 n; - static constexpr auto name = "TY3"; -}; -struct TY4 { - int a; - TY2 n; - int b; - static constexpr auto name = "TY4"; -}; - -static TY2 q2; -struct TY5 { - TY2 &n = q2; - static constexpr auto name = "TY5"; -}; - -template <typename T> int testType() { - T t1[2], t2[3], t3[4]; - for (int i = 0; i < 2; i++) - t1[i].i1 = t3[i].i1 = 1; - -#pragma omp target map(tofrom : t1, t2, t3) - for (int i = 0; i < 2; i++) { - t1[i].i3 = t3[i].i3 = t1[i].i1; - t1[i].i1 = t3[i].i1 = 7; - } - - for (int i = 0; i < 2; i++) { - if (t1[i].i3 != 1) { - printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, - t1[i].i3, i, t1[i].i1); - return 1; - } - if (t3[i].i3 != 1) { - printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, - t3[i].i3, i, t3[i].i1); - return 1; - } - } - - int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); - int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); - int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); - - // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY3: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY4: t1 0, t2 0, t3 0, expected 3x 0 - printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, - pt0, pt1, pt2); - return pt0 + pt1 + pt2; -} - -template <typename T> int testTypeNestedPtr(T t1[2], T t2[3], T t3[4]) { - for (int i = 0; i < 2; i++) - t1[i].n.i1 = t3[i].n.i1 = 1; - -#pragma omp target map(tofrom : t1[0 : 2], t2[0 : 3], t3[0 : 4]) - for (int i = 0; i < 2; i++) { - t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; - t1[i].n.i1 = t3[i].n.i1 = 7; - } - - for (int i = 0; i < 2; i++) { - if (t1[i].n.i3 != t1[i].n.i1) { - printf("failed %s-ptr. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, - t1[i].n.i3, i, t1[i].n.i1); - return 1; - } - if (t3[i].n.i3 != t3[i].n.i1) { - printf("failed %s-ptr. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, - t3[i].n.i3, i, t3[i].n.i1); - return 1; - } - } - - int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); - int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); - int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); - - // CHECK: present check for TY3-ptr: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY4-ptr: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY5-ptr: t1 0, t2 0, t3 0, expected 3x 0 - printf("present check for %s-ptr: t1 %i, t2 %i, t3 %i, expected 3x 0\n", - T::name, pt0, pt1, pt2); - return pt0 + pt1 + pt2; -} - -template <typename T> int testTypeNested() { - T t1[2], t2[3], t3[4]; - testTypeNestedPtr(t1, t2, t3); - for (int i = 0; i < 2; i++) - t1[i].n.i1 = t3[i].n.i1 = 1; - -#pragma omp target map(tofrom : t1, t2, t3) - for (int i = 0; i < 2; i++) { - t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; - t1[i].n.i1 = t3[i].n.i1 = 7; - } - - for (int i = 0; i < 2; i++) { - if (t1[i].n.i3 != t1[i].n.i1) { - printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, - t1[i].n.i3, i, t1[i].n.i1); - return 1; - } - if (t3[i].n.i3 != t3[i].n.i1) { - printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, - t3[i].n.i3, i, t3[i].n.i1); - return 1; - } - } - - int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); - int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); - int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); - - // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0 - // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0 - printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, - pt0, pt1, pt2); - return pt0 + pt1 + pt2; -} - -int main(int argc, char **argv) { - int r = 0; - r += testType<TY1>(); - r += testType<TY2>(); - r += testTypeNested<TY3>(); - r += testTypeNested<TY4>(); - { - TY2 a[9]; - TY5 t1[2], t2[3], t3[4]; - int i = 0; - for (int j = 0; j < 2; j++) - t1[j].n = a[i++]; - for (int j = 0; j < 3; j++) - t2[j].n = a[i++]; - for (int j = 0; j < 4; j++) - t3[j].n = a[i++]; - r += testTypeNestedPtr<TY5>(t1, t2, t3); - } - return r; -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
