[OPENMP]Fix PR49698: OpenMP declare mapper causes segmentation fault.

[ABataev]

Summary

The implicitly generated mappings for allocation/deallocation in mappers
runtime should be mapped as implicit, also no need to clear member_of
flag to avoid ref counter increment. Also, the ref counter should not be
incremented for the very first element that comes from the mapper
function.

[jdoerfert]

Can we use the reproducer from the bug report, I want an outermost array section with objects that have a declare mapper.

[abhinavgaba]

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp

8

Thanks for the fixes, Alexey. Are you planning on handling this case separately?

[ABataev]

In D100673#2695694, @jdoerfert wrote:

Can we use the reproducer from the bug report, I want an outermost array section with objects that have a declare mapper.

Sure, will add.

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp

8

Yes

[ABataev]

Added repro from PR49698

[jdoerfert]

LG, thanks for the quick fix.

[protze.joachim]

I tested the patch and still get wrong results (I modified the declare_mapper_nested_default_mappers_complex_structure.cpp test to use my verbose printf from bugzilla):

Modified Data Hierarchy:
Object C (0x623e50, 2) Contents:
        Object B (0x623e80, 2) Contents:
                Object A (0x6243f0) Contents:
                        data1 = 6439680  data2 = 0
                Object A (0x6243f8) Contents:
                        data1 = -784902216  data2 = 11062
        Object B (0x623e90, 2) Contents:
                Object A (0x6249e0) Contents:
                        data1 = 6441856  data2 = 0
                Object A (0x6249e8) Contents:
                        data1 = 6439904  data2 = 0
Object C (0x623e60, 2) Contents:
        Object B (0x623ef0, 2) Contents:
                Object A (0x624b90) Contents:
                        data1 = 6438736  data2 = 0
                Object A (0x624b98) Contents:
                        data1 = 6441344  data2 = 0
        Object B (0x623f00, 2) Contents:
                Object A (0x624c60) Contents:
                        data1 = 6444032  data2 = 0
                Object A (0x624c68) Contents:
                        data1 = 6441856  data2 = 0

The values inside the target region look good, but things break when mapping back.

When I remove the teams distribute for and only leave #pragma omp target, I get a different result and the test mistakenly passes:

Modified Data Hierarchy:
Object C (0x622e40, 2) Contents:
        Object B (0x622e70, 2) Contents:
                Object A (0x6233e0) Contents:
                        data1 = 6433120  data2 = 0
                Object A (0x6233e8) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622e80, 2) Contents:
                Object A (0x6239d0) Contents:
                        data1 = 6435792  data2 = 0
                Object A (0x6239d8) Contents:
                        data1 = 11  data2 = 22
Object C (0x622e50, 2) Contents:
        Object B (0x622ee0, 2) Contents:
                Object A (0x623b80) Contents:
                        data1 = 6437312  data2 = 0
                Object A (0x623b88) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622ef0, 2) Contents:
                Object A (0x623c50) Contents:
                        data1 = 6437744  data2 = 0
                Object A (0x623c58) Contents:
                        data1 = 11  data2 = 22
Testing for correctness...
outer[1].arr[1].arr[1].data2 = 22.

I'm currently only testing x86_64 offloading.

[ABataev]

In D100673#2698168, @protze.joachim wrote:

I tested the patch and still get wrong results (I modified the declare_mapper_nested_default_mappers_complex_structure.cpp test to use my verbose printf from bugzilla):

Modified Data Hierarchy:
Object C (0x623e50, 2) Contents:
        Object B (0x623e80, 2) Contents:
                Object A (0x6243f0) Contents:
                        data1 = 6439680  data2 = 0
                Object A (0x6243f8) Contents:
                        data1 = -784902216  data2 = 11062
        Object B (0x623e90, 2) Contents:
                Object A (0x6249e0) Contents:
                        data1 = 6441856  data2 = 0
                Object A (0x6249e8) Contents:
                        data1 = 6439904  data2 = 0
Object C (0x623e60, 2) Contents:
        Object B (0x623ef0, 2) Contents:
                Object A (0x624b90) Contents:
                        data1 = 6438736  data2 = 0
                Object A (0x624b98) Contents:
                        data1 = 6441344  data2 = 0
        Object B (0x623f00, 2) Contents:
                Object A (0x624c60) Contents:
                        data1 = 6444032  data2 = 0
                Object A (0x624c68) Contents:
                        data1 = 6441856  data2 = 0

The values inside the target region look good, but things break when mapping back.

When I remove the teams distribute for and only leave #pragma omp target, I get a different result and the test mistakenly passes:

Modified Data Hierarchy:
Object C (0x622e40, 2) Contents:
        Object B (0x622e70, 2) Contents:
                Object A (0x6233e0) Contents:
                        data1 = 6433120  data2 = 0
                Object A (0x6233e8) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622e80, 2) Contents:
                Object A (0x6239d0) Contents:
                        data1 = 6435792  data2 = 0
                Object A (0x6239d8) Contents:
                        data1 = 11  data2 = 22
Object C (0x622e50, 2) Contents:
        Object B (0x622ee0, 2) Contents:
                Object A (0x623b80) Contents:
                        data1 = 6437312  data2 = 0
                Object A (0x623b88) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622ef0, 2) Contents:
                Object A (0x623c50) Contents:
                        data1 = 6437744  data2 = 0
                Object A (0x623c58) Contents:
                        data1 = 11  data2 = 22
Testing for correctness...
outer[1].arr[1].arr[1].data2 = 22.

I'm currently only testing x86_64 offloading.

The problem is in transferring from the device. Looks like we need not only to invert the list of data in one data transfer block but also invert the list of blocks. Trying to implement this.

[ABataev]

Fixes and rebase

[protze.joachim]

The latest changes fix the issue. I tested with x86_64 and nvidia offloading.

The behavior and tests looks good to me, just one comment.

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp

124–127

Do these asserts make sense with the assert in line 104?

[ABataev]

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp

124–127

Ah, just forgot to remove these lines, thanks!

[jdoerfert]

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

[ABataev]

In D100673#2714087, @jdoerfert wrote:

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Will investigate this on Monday.

[ABataev]

In D100673#2714087, @jdoerfert wrote:

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

[jdoerfert]

In D100673#2716916, @ABataev wrote:

In D100673#2714087, @jdoerfert wrote:

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

I did debug something else and the clang-omp branch failed until I reverted this patch, at least that is what I thought. I will look again as soon as I find the time. For sure it run fine until very recently and we should make sure it continues to do so ;)

[ABataev]

In D100673#2716929, @jdoerfert wrote:

In D100673#2716916, @ABataev wrote:

In D100673#2714087, @jdoerfert wrote:

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

I did debug something else and the clang-omp branch failed until I reverted this patch, at least that is what I thought. I will look again as soon as I find the time. For sure it run fine until very recently and we should make sure it continues to do so ;)

Definitely! That's why I'm asking for the debug log for the investigation. It requires CUDA but I don't have an access to the machine with cuda.

[cchen]

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp

55

Do we really expect sa[1].f.b be the same as &x[0]? Shouldn't sa[1].f.b be the same as &y[0]?

69

Same here.

[ABataev]

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp

55

Yes, looks so, missed it.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
   // Size).
   llvm::Value *ArraySize = MapperCGF.Builder.CreateNUWMul(
       Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity()));
   // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves
   // memory allocation/deletion purpose only.
   llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd(
       MapType,
       MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
-                                   MappableExprsHandler::OMP_MAP_FROM |
-                                   MappableExprsHandler::OMP_MAP_MEMBER_OF)));
+                                   MappableExprsHandler::OMP_MAP_FROM)));
+  MapTypeArg = MapperCGF.Builder.CreateOr(
+      MapTypeArg,
+      MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_IMPLICIT));
 
   // Call the runtime API __tgt_push_mapper_component to fill up the runtime
   // data structure.
   llvm::Value *OffloadingArgs[] = {Handle,    Base,       Begin,
                                    ArraySize, MapTypeArg, MapName};
   MapperCGF.EmitRuntimeCall(
       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                             OMPRTL___tgt_push_mapper_component),

clang/test/OpenMP/declare_mapper_codegen.cpp

 // 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]]
 // CK0: br i1 [[CMP1]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
 // CK0: [[INIT]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK0: br label %[[LHEAD:[^,]+]]
 
 // CK0: [[LHEAD]]
 // CK0: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]]
 // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK0: [[LBODY]]
 // CK0: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0

 // CK0: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 // CK0: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 // CK0: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 // CK0: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK0: [[EVALDEL]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK0: br label %[[DONE]]
 // CK0: [[DONE]]
 // CK0: ret void
 
 
 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 void foo(int a){
   int i = a;

 // 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]]
 // CK1: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // CK1: [[INITEVALDEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK1: br label %[[LHEAD:[^,]+]]
 
 // CK1: [[LHEAD]]
 // CK1: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]]
 // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK1: [[LBODY]]
 // CK1: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0

 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // CK1: [[LEXIT]]
 // CK1: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 // CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 // CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK1: br label %[[DONE]]
 // CK1: [[DONE]]
 // CK1: ret void
 
 #endif // CK1
 
 
 ///==========================================================================///

 // 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]]
 // CK2: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // CK2: [[INITEVALDEL]]
 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK2: br label %[[LHEAD:[^,]+]]
 
 // CK2: [[LHEAD]]
 // CK2: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]]
 // CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK2: [[LBODY]]
 // CK2: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1

 // CK2: [[LEXIT]]
 // CK2: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 // CK2: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 // CK2: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 // CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK2: [[EVALDEL]]
 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK2: br label %[[DONE]]
 // CK2: [[DONE]]
 // CK2: ret void
 
 #endif // CK2
 
 
 ///==========================================================================///

 // CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 // CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
 // CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
 // CK4: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // CK4: [[INITEVALDEL]]
 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK4: br label %[[LHEAD:[^,]+]]
 
 // CK4: [[LHEAD]]
 // CK4: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]]
 // CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK4: [[LBODY]]
 // CK4: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0

 // CK4: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 // CK4: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 // CK4: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 // CK4: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK4: [[EVALDEL]]
 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK4: br label %[[DONE]]
 // CK4: [[DONE]]
 // CK4: ret void
 
 
 // CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 void foo(int a){
   int i = a;

openmp/libomptarget/src/omptarget.cpp

for (int32_t i = 0; i < arg_num; ++i) {
     // a close map modifier was associated with a map that contained a to.
     bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
     bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
     // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
     // have reached this point via __tgt_target_data_begin and not __tgt_target
     // 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 = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF);
+    bool UpdateRef =
+        !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && i == 0);
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
       // Base is address of pointer.
       //
       // Usually, the pointer is already allocated by this time.  For example:
       //
       //   #pragma omp target map(s.p[0:N])
       //
if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
     }
 
     bool IsLast, IsHostPtr;
     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 &&
-                      (!FromMapper || I != ArgNum - 1));
+                     (!FromMapper || I != ArgNum - 1);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(
         HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit);
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
for (int32_t I = ArgNum - 1; I >= 0; --I) {
 
     bool DelEntry = IsLast || ForceDelete;
 
     // If the last element from the mapper (for end transfer args comes in
     // reverse order), do not remove the partial entry, the parent struct still
     // exists.
     if (((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
          !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) ||
-        (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && FromMapper &&
-         I == ArgNum - 1)) {
+        (FromMapper && I == ArgNum - 1)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       // Move data back to the host
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) {
         bool Always = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
         bool CopyMember = false;

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+// XFAIL: clang
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D sa[2];
+  double x[2], y[2];
+  double x1[2], y1[2];
+  y[1] = x[1] = 20;
+
+  sa[0].e = 111;
+  sa[0].f.a = 222;
+  sa[0].f.c.a = 777;
+  sa[0].f.b = &x[0];
+  sa[0].f.c.b = &x1[0];
+  sa[0].h = N;
+
+  sa[1].e = 111;
+  sa[1].f.a = 222;
+  sa[1].f.c.a = 777;
+  sa[1].f.b = &y[0];
+  sa[1].f.c.b = &y1[0];
+  sa[1].h = N;
+
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&y[0]);
+#pragma omp target map(tofrom : sa) firstprivate(p)
+  {
+    printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a,
+           sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sa[1].e = 333;
+    sa[1].f.a = 444;
+    sa[1].f.c.a = 555;
+    sa[1].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D sa[10];
+  sa[1].e = 111;
+  sa[1].f.a = 222;
+  sa[1].f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  sa[1].f.b = &x[0];
+  sa[1].f.c.b = &x1[0];
+  sa[1].h = N;
+
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+#pragma omp target map(tofrom : sa[1]) firstprivate(p)
+  {
+    printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a,
+           sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sa[1].e = 333;
+    sa[1].f.a = 444;
+    sa[1].f.c.a = 555;
+    sa[1].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 2
+
+class MyObjectA {
+public:
+  MyObjectA() {
+    data1 = 1;
+    data2 = 2;
+  }
+  void show() {
+    printf("\t\tObject A Contents:\n");
+    printf("\t\t\tdata1 = %d  data2 = %d\n", data1, data2);
+  }
+  void foo() {
+    data1 += 10;
+    data2 += 20;
+  }
+  int data1;
+  int data2;
+};
+
+class MyObjectB {
+public:
+  MyObjectB() {
+    arr = new MyObjectA[N];
+    len = N;
+  }
+  void show() {
+    printf("\tObject B Contents:\n");
+    for (int i = 0; i < len; i++)
+      arr[i].show();
+  }
+  void foo() {
+    for (int i = 0; i < len; i++)
+      arr[i].foo();
+  }
+  MyObjectA *arr;
+  int len;
+};
+#pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[:obj.len])
+
+class MyObjectC {
+public:
+  MyObjectC() {
+    arr = new MyObjectB[N];
+    len = N;
+  }
+  void show() {
+    printf("Object C Contents:\n");
+    for (int i = 0; i < len; i++)
+      arr[i].show();
+  }
+  void foo() {
+    for (int i = 0; i < len; i++)
+      arr[i].foo();
+  }
+  MyObjectB *arr;
+  int len;
+};
+#pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[:obj.len])
+
+int main(void) {
+  MyObjectC *outer = new MyObjectC[N];
+
+  printf("Original data hierarchy:\n");
+  for (int i = 0; i < N; i++)
+    outer[i].show();
+
+  printf("Sending data to device...\n");
+#pragma omp target enter data map(to : outer[:N])
+
+  printf("Calling foo()...\n");
+#pragma omp target teams distribute parallel for
+  for (int i = 0; i < N; i++)
+    outer[i].foo();
+
+  printf("foo() complete!\n");
+
+  printf("Sending data back to host...\n");
+#pragma omp target exit data map(from : outer[:N])
+
+  printf("Modified Data Hierarchy:\n");
+  for (int i = 0; i < N; i++)
+    outer[i].show();
+
+  printf("Testing for correctness...\n");
+  printf("outer[1].arr[1].arr[1].data2 = %d.\n", outer[1].arr[1].arr[1].data2);
+  // CHECK: outer[1].arr[1].arr[1].data2 = 22.
+  assert(outer[1].arr[1].arr[1].data2 == 22);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D s;
+  s.e = 111;
+  s.f.a = 222;
+  s.f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  s.f.b = &x[0];
+  s.f.c.b = &x1[0];
+  s.h = N;
+
+  D *sp = &s;
+
+  printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1],
+         sp[0].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+#pragma omp target map(tofrom : sp[0]) firstprivate(p)
+  {
+    printf("%d %d %d\n", sp[0].f.a, sp[0].f.c.a,
+           sp[0].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sp[0].e = 333;
+    sp[0].f.a = 444;
+    sp[0].f.c.a = 555;
+    sp[0].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1],
+         sp[0].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D s;
+  s.e = 111;
+  s.f.a = 222;
+  s.f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  s.f.b = &x[0];
+  s.f.c.b = &x1[0];
+  s.h = N;
+
+  printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1],
+         s.f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+
+#pragma omp target map(tofrom : s) firstprivate(p)
+  {
+    printf("%d %d %d\n", s.f.a, s.f.c.a,
+           s.f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    s.e = 333;
+    s.f.a = 444;
+    s.f.c.a = 555;
+    s.f.b[1] = 40;
+  }
+
+  printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1],
+         s.f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}