[clang] [llvm] [mlir] Revert "[OpenMP][Mappers] Fix ref-count tracking for maps inserted by mappers. (#177059)" (PR #177623)
Abhinav Gaba via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 23 09:33:40 PST 2026
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 <abhinav.gaba at intel.com>
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 <abhinav.gaba at intel.com>
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;
-}
More information about the cfe-commits
mailing list