[clang] 187c39b - [OpenMP] Propagate ALWAYS/DELETE/CLOSE map-type modifiers to mapper entries (#210210)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Aug 2 23:07:41 PDT 2026
Author: Abhinav Gaba
Date: 2026-08-02T23:07:32-07:00
New Revision: 187c39ba03cfed96c8aab9d0c951d0bea6f27b82
URL: https://github.com/llvm/llvm-project/commit/187c39ba03cfed96c8aab9d0c951d0bea6f27b82
DIFF: https://github.com/llvm/llvm-project/commit/187c39ba03cfed96c8aab9d0c951d0bea6f27b82.diff
LOG: [OpenMP] Propagate ALWAYS/DELETE/CLOSE map-type modifiers to mapper entries (#210210)
Per OpenMP, when a map/motion clause uses a mapper, any
map-type-modifying modifier on that clause applies to each map the
declared mapper specifies.
This change propagates the `ALWAYS`, `DELETE`, and `CLOSE` bits from the
outer clause's map type into every entry emitted by
emitUserDefinedMapper, except `ATTACH` entries (`ATTACH`|`ALWAYS` is
reserved for `attach(always)`, and the other bits have no meaning for an
`ATTACH` entry).
`PRESENT` is intentionally NOT propagated here: it requires
distinguishing pointee entries from the struct's own storage and is
handled in a follow-up.
Co-authored-by: Claude Opus 4.8 <noreply at anthropic.com>
Added:
Modified:
clang/docs/ReleaseNotes.md
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
mlir/test/Target/LLVMIR/omptarget-llvm.mlir
offload/test/mapping/mapper_map_always_from.c
offload/test/mapping/mapper_map_always_to_enter_data.c
offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
offload/test/mapping/mapper_map_ptee_only_always_array.c
offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
Removed:
################################################################################
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index a38b99ff8e075..7108392abbaa1 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -546,6 +546,8 @@ features cannot lower the translation-unit ABI level;
- Added parsing and semantic support for `dims` modifier in `num_teams` and
`thread_limit` clauses for OpenMP 6.1 or later.
+- Map-type-modifying modifiers applied to a list item with a user-defined mapper
+ are now propagated onto the maps the mapper expands to.
### SYCL Support
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index ff891396fe72e..eb90f218f5ca8 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -148,7 +148,9 @@ class C {
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
+// CK0-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036
+// CK0-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}})
// 281474976710659 == 0x1,000,000,003
// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
@@ -171,7 +173,9 @@ class C {
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
+// CK0-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036
+// CK0-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
// 281474976710675 == 0x1,000,000,013
// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
@@ -194,7 +198,9 @@ class C {
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
+// CK0-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036
+// CK0-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}})
// CK0: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK0: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -636,7 +642,9 @@ class C {
// CK1-DAG: br label %[[TYEND]]
// CK1-DAG: [[TYEND]]
// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
+// CK1-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036
+// CK1-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
// CK1: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK1: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -743,7 +751,9 @@ class C {
// CK2-DAG: br label %[[TYEND]]
// CK2-DAG: [[TYEND]]
// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1]], {{.*}})
+// CK2-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036
+// CK2-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1_MOD]], {{.*}})
// CK2: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK2: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -948,7 +958,9 @@ class C {
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
+// CK4-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036
+// CK4-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
+// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}})
// 281474976710659 == 0x1,000,000,003
// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
@@ -971,7 +983,9 @@ class C {
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
+// CK4-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036
+// CK4-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
+// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
// 281474976710675 == 0x1,000,000,013
// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
@@ -994,7 +1008,9 @@ class C {
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
+// CK4-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036
+// CK4-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]]
+// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}})
// CK4: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK4: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -1126,7 +1142,9 @@ void foo(){
// CK5-DAG: br label %[[TYEND]]
// CK5-DAG: [[TYEND]]
// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}})
+// CK5-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036
+// CK5-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1_MOD]], {{.*}})
// CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 1
// CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK5: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
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 c9bc3a54c38f2..d3270669da3da 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
@@ -139,7 +139,7 @@ void foo() {
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// 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_END22:%.*]] ]
// 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
@@ -170,87 +170,95 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0
-// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// 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: omp.type.alloc1:
-// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP29]], -4
+// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
-// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1
-// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
+// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1
+// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
// CHECK: omp.type.to3:
-// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP29]], -3
+// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.to.else4:
-// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2
-// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
+// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2
+// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
// CHECK: omp.type.from5:
-// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP29]], -2
+// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
-// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[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: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0
-// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]]
-// CHECK: omp.type.alloc8:
-// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
-// CHECK: omp.type.alloc.else9:
-// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1
-// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]]
-// CHECK: omp.type.to10:
-// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.to.else11:
-// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]]
-// CHECK: omp.type.from12:
-// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.end13:
-// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[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: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0
-// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]]
-// CHECK: omp.type.alloc15:
-// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.alloc.else16:
-// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1
-// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]]
-// CHECK: omp.type.to17:
-// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.to.else18:
-// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2
-// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]]
-// CHECK: omp.type.from19:
-// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.end20:
-// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[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_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ]
+// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null)
+// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
+// CHECK: omp.type.alloc9:
+// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]]
+// CHECK: omp.type.alloc.else10:
+// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
+// CHECK: omp.type.to11:
+// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.to.else12:
+// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
+// CHECK: omp.type.from13:
+// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.end14:
+// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ]
+// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]]
+// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]]
+// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]]
+// CHECK: omp.type.alloc17:
+// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.alloc.else18:
+// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]]
+// CHECK: omp.type.to19:
+// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.to.else20:
+// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]]
+// CHECK: omp.type.from21:
+// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.end22:
+// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ]
+// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], 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: [[TMP53:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0
-// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
+// 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_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12
-// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], 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
@@ -303,21 +311,23 @@ void foo() {
// 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: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], 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: [[TMP26:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0
+// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP27]], 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: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], 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 f530c92c8f1fe..8723ad765c5fb 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
@@ -136,7 +136,7 @@ void foo() {
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// 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_END22:%.*]] ]
// 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
@@ -167,87 +167,95 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0
-// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// 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: omp.type.alloc1:
-// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP29]], -4
+// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
-// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1
-// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
+// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1
+// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
// CHECK: omp.type.to3:
-// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP29]], -3
+// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.to.else4:
-// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2
-// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
+// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2
+// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
// CHECK: omp.type.from5:
-// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP29]], -2
+// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
-// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[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: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0
-// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]]
-// CHECK: omp.type.alloc8:
-// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
-// CHECK: omp.type.alloc.else9:
-// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1
-// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]]
-// CHECK: omp.type.to10:
-// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.to.else11:
-// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]]
-// CHECK: omp.type.from12:
-// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.end13:
-// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[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: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0
-// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]]
-// CHECK: omp.type.alloc15:
-// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.alloc.else16:
-// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1
-// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]]
-// CHECK: omp.type.to17:
-// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.to.else18:
-// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2
-// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]]
-// CHECK: omp.type.from19:
-// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.end20:
-// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[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_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ]
+// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null)
+// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
+// CHECK: omp.type.alloc9:
+// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]]
+// CHECK: omp.type.alloc.else10:
+// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
+// CHECK: omp.type.to11:
+// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.to.else12:
+// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
+// CHECK: omp.type.from13:
+// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.end14:
+// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ]
+// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]]
+// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]]
+// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]]
+// CHECK: omp.type.alloc17:
+// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.alloc.else18:
+// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]]
+// CHECK: omp.type.to19:
+// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.to.else20:
+// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]]
+// CHECK: omp.type.from21:
+// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.end22:
+// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ]
+// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], 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: [[TMP53:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0
-// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
+// 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_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12
-// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], 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
@@ -300,21 +308,23 @@ void foo() {
// 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: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], 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: [[TMP26:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0
+// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP27]], 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: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
index cfdc9b861090c..bd9189dce7c9f 100644
--- a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
@@ -98,7 +98,7 @@ void foo(S2 *arr) {
// CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK: omp.arraymap.body:
-// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ]
+// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
// CHECK: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
@@ -135,87 +135,95 @@ void foo(S2 *arr) {
// CHECK: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK: [[TMP31:%.*]] = add nuw i64 281474976710659, [[TMP22]]
-// CHECK: [[TMP32:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0
-// CHECK: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
+// CHECK: [[TMP31:%.*]] = and i64 [[TMP4]], 1036
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]]
+// CHECK: [[TMP33:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0
+// CHECK: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
// CHECK: omp.type.alloc4:
-// CHECK: [[TMP34:%.*]] = and i64 [[TMP31]], -4
+// CHECK: [[TMP35:%.*]] = and i64 [[TMP32]], -4
// CHECK: br label [[OMP_TYPE_END9:%.*]]
// CHECK: omp.type.alloc.else5:
-// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1
-// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
+// CHECK: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1
+// CHECK: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
// CHECK: omp.type.to6:
-// CHECK: [[TMP36:%.*]] = and i64 [[TMP31]], -3
+// CHECK: [[TMP37:%.*]] = and i64 [[TMP32]], -3
// CHECK: br label [[OMP_TYPE_END9]]
// CHECK: omp.type.to.else7:
-// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2
-// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
+// CHECK: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2
+// CHECK: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
// CHECK: omp.type.from8:
-// CHECK: [[TMP38:%.*]] = and i64 [[TMP31]], -2
+// CHECK: [[TMP39:%.*]] = and i64 [[TMP32]], -2
// CHECK: br label [[OMP_TYPE_END9]]
// CHECK: omp.type.end9:
-// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP36]], [[OMP_TYPE_TO6]] ], [ [[TMP38]], [[OMP_TYPE_FROM8]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE7]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE10]], ptr null)
-// CHECK: [[TMP39:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
-// CHECK: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC11:%.*]], label [[OMP_TYPE_ALLOC_ELSE12:%.*]]
-// CHECK: omp.type.alloc11:
-// CHECK: [[TMP42:%.*]] = and i64 [[TMP39]], -4
-// CHECK: br label [[OMP_TYPE_END16:%.*]]
-// CHECK: omp.type.alloc.else12:
-// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
-// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_TO13:%.*]], label [[OMP_TYPE_TO_ELSE14:%.*]]
-// CHECK: omp.type.to13:
-// CHECK: [[TMP44:%.*]] = and i64 [[TMP39]], -3
-// CHECK: br label [[OMP_TYPE_END16]]
-// CHECK: omp.type.to.else14:
-// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
-// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_FROM15:%.*]], label [[OMP_TYPE_END16]]
-// CHECK: omp.type.from15:
-// CHECK: [[TMP46:%.*]] = and i64 [[TMP39]], -2
-// CHECK: br label [[OMP_TYPE_END16]]
-// CHECK: omp.type.end16:
-// CHECK: [[OMP_MAPTYPE17:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC11]] ], [ [[TMP44]], [[OMP_TYPE_TO13]] ], [ [[TMP46]], [[OMP_TYPE_FROM15]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE14]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE17]], ptr null)
-// CHECK: [[TMP47:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK: [[TMP48:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0
-// CHECK: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC18:%.*]], label [[OMP_TYPE_ALLOC_ELSE19:%.*]]
-// CHECK: omp.type.alloc18:
-// CHECK: [[TMP50:%.*]] = and i64 [[TMP47]], -4
-// CHECK: br label [[OMP_TYPE_END23]]
-// CHECK: omp.type.alloc.else19:
-// CHECK: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1
-// CHECK: br i1 [[TMP51]], label [[OMP_TYPE_TO20:%.*]], label [[OMP_TYPE_TO_ELSE21:%.*]]
-// CHECK: omp.type.to20:
-// CHECK: [[TMP52:%.*]] = and i64 [[TMP47]], -3
-// CHECK: br label [[OMP_TYPE_END23]]
-// CHECK: omp.type.to.else21:
-// CHECK: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2
-// CHECK: br i1 [[TMP53]], label [[OMP_TYPE_FROM22:%.*]], label [[OMP_TYPE_END23]]
-// CHECK: omp.type.from22:
-// CHECK: [[TMP54:%.*]] = and i64 [[TMP47]], -2
-// CHECK: br label [[OMP_TYPE_END23]]
-// CHECK: omp.type.end23:
-// CHECK: [[OMP_MAPTYPE24:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC18]] ], [ [[TMP52]], [[OMP_TYPE_TO20]] ], [ [[TMP54]], [[OMP_TYPE_FROM22]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE21]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE24]], ptr null)
+// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ]
+// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 1036
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
+// CHECK: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]]
+// CHECK: [[TMP42:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
+// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
+// CHECK: omp.type.alloc12:
+// CHECK: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK: br label [[OMP_TYPE_END17:%.*]]
+// CHECK: omp.type.alloc.else13:
+// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
+// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
+// CHECK: omp.type.to14:
+// CHECK: [[TMP46:%.*]] = and i64 [[TMP41]], -3
+// CHECK: br label [[OMP_TYPE_END17]]
+// CHECK: omp.type.to.else15:
+// CHECK: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
+// CHECK: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
+// CHECK: omp.type.from16:
+// CHECK: [[TMP48:%.*]] = and i64 [[TMP41]], -2
+// CHECK: br label [[OMP_TYPE_END17]]
+// CHECK: omp.type.end17:
+// CHECK: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ]
+// CHECK: [[TMP49:%.*]] = and i64 [[TMP4]], 1036
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
+// CHECK: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]]
+// CHECK: [[TMP51:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0
+// CHECK: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// CHECK: omp.type.alloc20:
+// CHECK: [[TMP53:%.*]] = and i64 [[TMP50]], -4
+// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.alloc.else21:
+// CHECK: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1
+// CHECK: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// CHECK: omp.type.to22:
+// CHECK: [[TMP55:%.*]] = and i64 [[TMP50]], -3
+// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.to.else23:
+// CHECK: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2
+// CHECK: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK: omp.type.from24:
+// CHECK: [[TMP57:%.*]] = and i64 [[TMP50]], -2
+// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.end25:
+// CHECK: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ]
+// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], 1036
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
// CHECK: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
-// CHECK: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK: [[TMP55:%.*]] = and i64 [[TMP4]], 8
-// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP55]], 0
-// CHECK: [[TMP56:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK: br i1 [[TMP56]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK: [[TMP59:%.*]] = and i64 [[TMP4]], 8
+// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0
+// CHECK: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK: br i1 [[TMP60]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK: [[TMP57:%.*]] = mul nuw i64 [[TMP6]], 16
-// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], -4
-// CHECK: [[TMP59:%.*]] = or i64 [[TMP58]], 512
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP57]], i64 [[TMP59]], ptr [[TMP5]])
+// CHECK: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK: [[TMP62:%.*]] = and i64 [[TMP4]], -4
+// CHECK: [[TMP63:%.*]] = or i64 [[TMP62]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]])
// CHECK: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK: ret void
@@ -268,7 +276,7 @@ void foo(S2 *arr) {
// CHECK-60: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK-60: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK-60: omp.arraymap.body:
-// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ]
+// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
// CHECK-60: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK-60: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK-60: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
@@ -305,87 +313,95 @@ void foo(S2 *arr) {
// CHECK-60: br label [[OMP_TYPE_END]]
// CHECK-60: omp.type.end:
// CHECK-60: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK-60: [[TMP31:%.*]] = add nuw i64 281474976710659, [[TMP22]]
-// CHECK-60: [[TMP32:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0
-// CHECK-60: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
+// CHECK-60: [[TMP31:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-60: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]]
+// CHECK-60: [[TMP33:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0
+// CHECK-60: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
// CHECK-60: omp.type.alloc4:
-// CHECK-60: [[TMP34:%.*]] = and i64 [[TMP31]], -4
+// CHECK-60: [[TMP35:%.*]] = and i64 [[TMP32]], -4
// CHECK-60: br label [[OMP_TYPE_END9:%.*]]
// CHECK-60: omp.type.alloc.else5:
-// CHECK-60: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1
-// CHECK-60: br i1 [[TMP35]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
+// CHECK-60: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1
+// CHECK-60: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
// CHECK-60: omp.type.to6:
-// CHECK-60: [[TMP36:%.*]] = and i64 [[TMP31]], -3
+// CHECK-60: [[TMP37:%.*]] = and i64 [[TMP32]], -3
// CHECK-60: br label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.to.else7:
-// CHECK-60: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2
-// CHECK-60: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
+// CHECK-60: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2
+// CHECK-60: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.from8:
-// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP31]], -2
+// CHECK-60: [[TMP39:%.*]] = and i64 [[TMP32]], -2
// CHECK-60: br label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.end9:
-// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP36]], [[OMP_TYPE_TO6]] ], [ [[TMP38]], [[OMP_TYPE_FROM8]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE7]] ]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE10]], ptr null)
-// CHECK-60: [[TMP39:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
-// CHECK-60: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC11:%.*]], label [[OMP_TYPE_ALLOC_ELSE12:%.*]]
-// CHECK-60: omp.type.alloc11:
-// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP39]], -4
-// CHECK-60: br label [[OMP_TYPE_END16:%.*]]
-// CHECK-60: omp.type.alloc.else12:
-// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
-// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_TO13:%.*]], label [[OMP_TYPE_TO_ELSE14:%.*]]
-// CHECK-60: omp.type.to13:
-// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP39]], -3
-// CHECK-60: br label [[OMP_TYPE_END16]]
-// CHECK-60: omp.type.to.else14:
-// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
-// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_FROM15:%.*]], label [[OMP_TYPE_END16]]
-// CHECK-60: omp.type.from15:
-// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP39]], -2
-// CHECK-60: br label [[OMP_TYPE_END16]]
-// CHECK-60: omp.type.end16:
-// CHECK-60: [[OMP_MAPTYPE17:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC11]] ], [ [[TMP44]], [[OMP_TYPE_TO13]] ], [ [[TMP46]], [[OMP_TYPE_FROM15]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE14]] ]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE17]], ptr null)
-// CHECK-60: [[TMP47:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK-60: [[TMP48:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0
-// CHECK-60: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC18:%.*]], label [[OMP_TYPE_ALLOC_ELSE19:%.*]]
-// CHECK-60: omp.type.alloc18:
-// CHECK-60: [[TMP50:%.*]] = and i64 [[TMP47]], -4
-// CHECK-60: br label [[OMP_TYPE_END23]]
-// CHECK-60: omp.type.alloc.else19:
-// CHECK-60: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1
-// CHECK-60: br i1 [[TMP51]], label [[OMP_TYPE_TO20:%.*]], label [[OMP_TYPE_TO_ELSE21:%.*]]
-// CHECK-60: omp.type.to20:
-// CHECK-60: [[TMP52:%.*]] = and i64 [[TMP47]], -3
-// CHECK-60: br label [[OMP_TYPE_END23]]
-// CHECK-60: omp.type.to.else21:
-// CHECK-60: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2
-// CHECK-60: br i1 [[TMP53]], label [[OMP_TYPE_FROM22:%.*]], label [[OMP_TYPE_END23]]
-// CHECK-60: omp.type.from22:
-// CHECK-60: [[TMP54:%.*]] = and i64 [[TMP47]], -2
-// CHECK-60: br label [[OMP_TYPE_END23]]
-// CHECK-60: omp.type.end23:
-// CHECK-60: [[OMP_MAPTYPE24:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC18]] ], [ [[TMP52]], [[OMP_TYPE_TO20]] ], [ [[TMP54]], [[OMP_TYPE_FROM22]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE21]] ]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE24]], ptr null)
+// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ]
+// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
+// CHECK-60: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]]
+// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
+// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
+// CHECK-60: omp.type.alloc12:
+// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK-60: br label [[OMP_TYPE_END17:%.*]]
+// CHECK-60: omp.type.alloc.else13:
+// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
+// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
+// CHECK-60: omp.type.to14:
+// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP41]], -3
+// CHECK-60: br label [[OMP_TYPE_END17]]
+// CHECK-60: omp.type.to.else15:
+// CHECK-60: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
+// CHECK-60: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
+// CHECK-60: omp.type.from16:
+// CHECK-60: [[TMP48:%.*]] = and i64 [[TMP41]], -2
+// CHECK-60: br label [[OMP_TYPE_END17]]
+// CHECK-60: omp.type.end17:
+// CHECK-60: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ]
+// CHECK-60: [[TMP49:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
+// CHECK-60: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]]
+// CHECK-60: [[TMP51:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0
+// CHECK-60: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// CHECK-60: omp.type.alloc20:
+// CHECK-60: [[TMP53:%.*]] = and i64 [[TMP50]], -4
+// CHECK-60: br label [[OMP_TYPE_END25]]
+// CHECK-60: omp.type.alloc.else21:
+// CHECK-60: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1
+// CHECK-60: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// CHECK-60: omp.type.to22:
+// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP50]], -3
+// CHECK-60: br label [[OMP_TYPE_END25]]
+// CHECK-60: omp.type.to.else23:
+// CHECK-60: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2
+// CHECK-60: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK-60: omp.type.from24:
+// CHECK-60: [[TMP57:%.*]] = and i64 [[TMP50]], -2
+// CHECK-60: br label [[OMP_TYPE_END25]]
+// CHECK-60: omp.type.end25:
+// CHECK-60: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ]
+// CHECK-60: [[TMP58:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
// CHECK-60: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-60: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-60: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK-60: omp.arraymap.exit:
-// CHECK-60: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP4]], 8
-// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP55]], 0
-// CHECK-60: [[TMP56:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-60: br i1 [[TMP56]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-60: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK-60: [[TMP59:%.*]] = and i64 [[TMP4]], 8
+// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0
+// CHECK-60: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-60: br i1 [[TMP60]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK-60: .omp.array..del:
-// CHECK-60: [[TMP57:%.*]] = mul nuw i64 [[TMP6]], 16
-// CHECK-60: [[TMP58:%.*]] = and i64 [[TMP4]], -4
-// CHECK-60: [[TMP59:%.*]] = or i64 [[TMP58]], 512
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP57]], i64 [[TMP59]], ptr [[TMP5]])
+// CHECK-60: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK-60: [[TMP62:%.*]] = and i64 [[TMP4]], -4
+// CHECK-60: [[TMP63:%.*]] = or i64 [[TMP62]], 512
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]])
// CHECK-60: br label [[OMP_DONE]]
// CHECK-60: omp.done:
// CHECK-60: ret void
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 05307414c94d0..d4798f5a92dc6 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3627,14 +3627,20 @@ class OpenMPIRBuilder {
/// for (unsigned i = 0; i < size; i++) {
/// // For each component specified by this mapper:
/// for (auto c : begin[i]->all_components) {
+ /// // Map-type-modifying bits (ALWAYS, DELETE, CLOSE) from the outer
+ /// // map clause are propagated to each component, except ATTACH
+ /// // entries (ATTACH|ALWAYS is reserved for attach(always), and other
+ /// // modifier bits have no meaning for ATTACH).
+ /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE);
+ /// effective_type = c.isAttach() ? c.arg_type
+ /// : c.arg_type | imported_modifier_bits;
/// if (c.hasMapper())
/// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin,
- /// c.arg_size,
- /// c.arg_type, c.arg_name);
+ /// c.arg_size, effective_type, c.arg_name);
/// else
/// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
- /// c.arg_begin, c.arg_size, c.arg_type,
- /// c.arg_name);
+ /// c.arg_begin, c.arg_size,
+ /// effective_type, c.arg_name);
/// }
/// }
/// // Delete the array section.
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index a227397b06f4e..9704f68df4ec7 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10615,8 +10615,44 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
CurMapType->addIncoming(FromMapType, FromBB);
CurMapType->addIncoming(MemberMapType, ToElseBB);
- Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
- CurSizeArg, CurMapType, CurNameArg};
+ // Propagate map-type-modifying bits from the outer map clause to each map
+ // inserted by the mapper.
+ //
+ // OpenMP 6.0:281:34: The effect of the mapper modifier is to remove the
+ // list item from the map clause and to apply the clauses specified in the
+ // declared mapper to the construct on which the map clause appears...
+ // If any modifier with the map-type-modifying property appears in the map
+ // clause then the effect is as if that modifier appears in each map clause
+ // specified in the declared mapper.
+ //
+ // Map-type-modifying bits: ALWAYS, DELETE, CLOSE, PRESENT.
+ // TODO: PRESENT is not propagated here yet. Doing so requires
+ // distinguishing pointee entries from the struct's own storage; it is
+ // handled in a follow-on.
+ Value *ImportedModifierBits = Builder.CreateAnd(
+ MapType,
+ Builder.getInt64(
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS |
+ OpenMPOffloadMappingFlags::OMP_MAP_DELETE |
+ OpenMPOffloadMappingFlags::OMP_MAP_CLOSE)));
+ Value *CurMapTypeWithModifiers = Builder.CreateOr(
+ CurMapType, ImportedModifierBits, "omp.maptype.with.modifiers");
+
+ // ATTACH entries must not receive map-type-modifying bits: ATTACH|ALWAYS is
+ // reserved for the attach(always) map-type modifier, and other modifier
+ // bits (DELETE, CLOSE) have no meaning for an ATTACH entry.
+ auto RawType =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ Info->Types[I]);
+ constexpr uint64_t AttachBit =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+ Value *FinalMapType =
+ (RawType & AttachBit) ? CurMapType : CurMapTypeWithModifiers;
+
+ Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
+ CurSizeArg, FinalMapType, CurNameArg};
auto ChildMapperFn = CustomMapperCB(I);
if (!ChildMapperFn)
diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
index bb6085840cf07..baf3404e468ab 100644
--- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
@@ -599,7 +599,11 @@ module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
// CHECK: br label %[[VAL_42]]
// CHECK: omp.type.end: ; preds = %[[VAL_59]], %[[VAL_56]], %[[VAL_55]], %[[VAL_51]]
// CHECK: %[[VAL_61:.*]] = phi i64 [ 0, %[[VAL_51]] ], [ 1, %[[VAL_55]] ], [ 2, %[[VAL_59]] ], [ 3, %[[VAL_56]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_45]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_61]], ptr @2)
+// Map-type modifiers (ALWAYS|DELETE|CLOSE = 1036) from the outer clause are
+// propagated to each entry the mapper pushes.
+// CHECK: %[[VAL_MODMASK:.*]] = and i64 %{{.*}}, 1036
+// CHECK: %[[VAL_MTYPEMOD:.*]] = or i64 %[[VAL_61]], %[[VAL_MODMASK]]
+// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_45]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_MTYPEMOD]], ptr @2)
// CHECK: %[[VAL_44]] = getelementptr %[[VAL_18]], ptr %[[VAL_43]], i32 1
// CHECK: %[[VAL_62:.*]] = icmp eq ptr %[[VAL_44]], %[[VAL_17]]
// CHECK: br i1 %[[VAL_62]], label %[[VAL_63:.*]], label %[[VAL_41]]
diff --git a/offload/test/mapping/mapper_map_always_from.c b/offload/test/mapping/mapper_map_always_from.c
index 848bdbb448458..1048b32e990ef 100644
--- a/offload/test/mapping/mapper_map_always_from.c
+++ b/offload/test/mapping/mapper_map_always_from.c
@@ -1,17 +1,12 @@
-// Show that the ALWAYS map-type modifier on the outer map clause should be
-// propagated to the entries pushed by a user-defined mapper.
+// Show that the ALWAYS map-type modifier on the outer map clause is propagated
+// to the entries pushed by a user-defined mapper.
//
// The mapper transfers s.y. We pre-map s.y so that on the target region below
// it already has a device copy with a nonzero reference count. Without ALWAYS,
-// the `from` at the end of the target region is suppressed (a present,
-// ref-counted entry is not copied back), so the write of 111 is lost. ALWAYS
-// must force the copy-back -- but only if ALWAYS is propagated from the outer
-// clause to the mapper's s.y entry.
-//
-// FIXME: ALWAYS is not propagated to the mapper's entries yet, so the copy-back
-// is currently suppressed and s.y reads back as 0 (its pre-map value). Once
-// ALWAYS is propagated:
-// EXPECTED: s.y = 111
+// the `from` at the end of the target region would be suppressed (a present,
+// ref-counted entry is not copied back), so the write of 111 would be lost.
+// ALWAYS forces the copy-back, which only happens if ALWAYS is propagated from
+// the outer clause to the mapper's s.y entry.
// RUN: %libomptarget-compile-run-and-check-generic
@@ -37,10 +32,8 @@ int main() {
s.y = 111;
}
- // ALWAYS should force s.y back even though it is still mapped (ref count >
- // 0), but the modifier is not propagated yet, so the copy-back does not
- // happen.
- printf("s.y = %d\n", s.y); // CHECK: s.y = 0
+ // ALWAYS forces s.y back even though it is still mapped (ref count > 0).
+ printf("s.y = %d\n", s.y); // CHECK: s.y = 111
#pragma omp target exit data map(delete : s.y)
}
diff --git a/offload/test/mapping/mapper_map_always_to_enter_data.c b/offload/test/mapping/mapper_map_always_to_enter_data.c
index 185cc76a75246..b1742d9d14ff5 100644
--- a/offload/test/mapping/mapper_map_always_to_enter_data.c
+++ b/offload/test/mapping/mapper_map_always_to_enter_data.c
@@ -1,20 +1,15 @@
-// Show that the ALWAYS map-type modifier on the outer map clause should be
-// propagated to the entries pushed by a user-defined mapper, using no target
-// construct at all: the device data is inspected with omp_get_mapped_ptr() and
+// Show that the ALWAYS map-type modifier on the outer map clause is propagated
+// to the entries pushed by a user-defined mapper, using no target construct at
+// all: the device data is inspected with omp_get_mapped_ptr() and
// omp_target_memcpy(), so what is checked is purely the data-motion done by
// `target enter data`.
//
// The mapper transfers s.y. We pre-map s.y so that it already has a device copy
// with a nonzero reference count, and set that copy to a known value. Without
-// ALWAYS, the `to` of the second `enter data` is suppressed (a present,
-// ref-counted entry is not copied), so the host's 111 does not reach the
-// device. ALWAYS must force the copy -- but only if ALWAYS is propagated from
-// the outer clause to the mapper's s.y entry.
-//
-// FIXME: ALWAYS is not propagated to the mapper's entries yet, so the transfer
-// is currently suppressed and the device copy still reads as 0 (its pre-set
-// value). Once ALWAYS is propagated:
-// EXPECTED: device s.y = 111
+// ALWAYS, the `to` of the second `enter data` would be suppressed (a present,
+// ref-counted entry is not copied), so the host's 111 would not reach the
+// device. ALWAYS forces the copy, which only happens if ALWAYS is propagated
+// from the outer clause to the mapper's s.y entry.
// RUN: %libomptarget-compile-run-and-check-generic
@@ -48,13 +43,12 @@ int main() {
#pragma omp target enter data map(always, to : s)
- // ALWAYS should force s.y to the device even though it is already mapped
- // (ref count > 0), but the modifier is not propagated yet, so the transfer
- // does not happen.
+ // ALWAYS forces s.y to the device even though it is already mapped
+ // (ref count > 0).
int dev_y_val = -1;
omp_target_memcpy(&dev_y_val, dev_y, sizeof(int), 0, 0, host, dev);
- printf("device s.y = %d\n", dev_y_val); // CHECK: device s.y = 0
+ printf("device s.y = %d\n", dev_y_val); // CHECK: device s.y = 111
#pragma omp target exit data map(delete : s.y)
}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
index fb03accec0083..1b91d87071cbd 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
@@ -54,8 +54,6 @@ int main() {
print_status(&s2.s1p->y, "y"); // CHECK: y is not present
print_status(&s2.z, "z"); // CHECK: z is not present
print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present
- print_status(&s2.s1p->p, "p"); // CHECK: p is present
- // FIXME: the DELETE modifier is not propagated to mapper entries yet.
- // EXPECTED: p is not present
- print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is not present
+ print_status(&s2.s1p->p, "p"); // CHECK: p is not present
+ print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is not present
}
diff --git a/offload/test/mapping/mapper_map_ptee_only_always_array.c b/offload/test/mapping/mapper_map_ptee_only_always_array.c
index 64aaabd80dfe2..48e2a492131d3 100644
--- a/offload/test/mapping/mapper_map_ptee_only_always_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_always_array.c
@@ -49,18 +49,10 @@ int main() {
printf("\n");
printf("After map(always,from)\n");
- printf("s[0].x = %d\n", s1[0].x); // CHECK: s[0].x = 111
- // FIXME: modifier-bits are not propagated to mapper entries yet.
- // EXPECTED: s[0].x = 222
- printf("s[1].x = %d\n", s1[1].x); // CHECK: s[1].x = 111
- // FIXME: modifier-bits are not propagated to mapper entries yet.
- // EXPECTED: s[1].x = 222
- printf("s[0].p[0] = %d\n", s1[0].p[0]); // CHECK: s[0].p[0] = 111
- // FIXME: modifier-bits are not propagated to mapper entries yet.
- // EXPECTED: s[0].p[0] = 222
- printf("s[1].p[0] = %d\n", s1[1].p[0]); // CHECK: s[1].p[0] = 0
- // FIXME: modifier-bits are not propagated to mapper entries yet.
- // EXPECTED: s[1].p[0] = 222
+ printf("s[0].x = %d\n", s1[0].x); // CHECK: s[0].x = 222
+ printf("s[1].x = %d\n", s1[1].x); // CHECK: s[1].x = 222
+ printf("s[0].p[0] = %d\n", s1[0].p[0]); // CHECK: s[0].p[0] = 222
+ printf("s[1].p[0] = %d\n", s1[1].p[0]); // CHECK: s[1].p[0] = 222
printf("\n");
#pragma omp target exit data map(delete : s1)
diff --git a/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 b/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
index 52e32b1d6ccad..1c2a67ad92d1a 100644
--- a/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
+++ b/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
@@ -1,20 +1,15 @@
-! Show that the ALWAYS map-type modifier on the outer map clause should be
-! propagated to the entries pushed by a user-defined mapper. The test uses no
-! target region at all: the device data is inspected with omp_get_mapped_ptr()
-! and omp_target_memcpy(), so what is checked is purely the data-motion done by
+! Show that the ALWAYS map-type modifier on the outer map clause is propagated
+! to the entries pushed by a user-defined mapper. The test uses no target
+! region at all: the device data is inspected with omp_get_mapped_ptr() and
+! omp_target_memcpy(), so what is checked is purely the data-motion done by
! `target enter data`.
!
! The mapper transfers s%y. We pre-map s%y so that it already has a device copy
! with a nonzero reference count, and set that copy to a known value. Without
-! ALWAYS, the `to` of the second `enter data` is suppressed (a present,
-! ref-counted entry is not copied), so the host's 111 does not reach the device.
-! ALWAYS must force the copy -- but only if ALWAYS is propagated from the outer
-! clause to the mapper's s%y entry.
-!
-! FIXME: ALWAYS is not propagated to the mapper's entries yet, so the transfer
-! is currently suppressed and the device copy still reads as 0 (its pre-set
-! value). Once ALWAYS is propagated:
-! EXPECTED: device s%y = 111
+! ALWAYS, the `to` of the second `enter data` would be suppressed (a present,
+! ref-counted entry is not copied), so the host's 111 would not reach the
+! device. ALWAYS forces the copy, which only happens if ALWAYS is propagated
+! from the outer clause to the mapper's s%y entry.
! REQUIRES: flang
@@ -56,16 +51,15 @@ program main
!$omp target enter data map(always, to: s)
- ! ALWAYS should force s%y to the device even though it is already mapped
- ! (ref count > 0), but the modifier is not propagated yet, so the transfer
- ! does not happen.
+ ! ALWAYS forces s%y to the device even though it is already mapped
+ ! (ref count > 0).
dev_y_val = -1
rc = omp_target_memcpy(c_loc(dev_y_val), dev_y, &
int(c_sizeof(dev_y_val), c_size_t), &
0_c_size_t, 0_c_size_t, host, dev)
print *, "device s%y =", dev_y_val
- ! CHECK: device s%y = 0
+ ! CHECK: device s%y = 111
!$omp target exit data map(delete: s%y)
end program main
More information about the cfe-commits
mailing list