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

Alexey Bataev via cfe-commits cfe-commits at lists.llvm.org
Wed Apr 21 10:39:24 PDT 2021


Author: Alexey Bataev
Date: 2021-04-21T10:38:31-07:00
New Revision: 079884225a5571f4caf1a9c71b6748db8192e383

URL: https://github.com/llvm/llvm-project/commit/079884225a5571f4caf1a9c71b6748db8192e383
DIFF: https://github.com/llvm/llvm-project/commit/079884225a5571f4caf1a9c71b6748db8192e383.diff

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

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.

Differential Revision: https://reviews.llvm.org/D100673

Added: 
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp
    clang/test/OpenMP/declare_mapper_codegen.cpp
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 18ad5066fbba2..0a408837d1c6a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10005,8 +10005,10 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
   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.

diff  --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index 2c488c556a5e3..61eec7a017323 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -118,8 +118,11 @@ class C {
 // 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]]
@@ -228,8 +231,11 @@ class C {
 // 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
@@ -672,8 +678,11 @@ class C {
 
 // 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]]
@@ -718,8 +727,11 @@ class C {
 // 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
@@ -793,8 +805,11 @@ class C {
 
 // 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]]
@@ -841,8 +856,11 @@ class C {
 // 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
@@ -998,8 +1016,11 @@ class C {
 // 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]]
@@ -1108,8 +1129,11 @@ class C {
 // 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

diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 34645535a5960..dfa1686e84e75 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -383,9 +383,7 @@ int targetDataMapper(ident_t *loc, DeviceTy &Device, void *arg_base, void *arg,
   std::vector<void *> MapperArgNames(MapperComponents.Components.size());
 
   for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {
-    auto &C =
-        MapperComponents
-            .Components[target_data_function == targetDataEnd ? E - I - 1 : I];
+    auto &C = MapperComponents.Components[I];
     MapperArgsBase[I] = C.Base;
     MapperArgs[I] = C.Begin;
     MapperArgSizes[I] = C.Size;
@@ -472,7 +470,8 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
     // 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.
@@ -615,6 +614,7 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
                   void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) {
   int Ret;
   std::vector<DeallocTgtPtrInfo> DeallocTgtPtrs;
+  void *FromMapperBase = nullptr;
   // process each input.
   for (int32_t I = ArgNum - 1; I >= 0; --I) {
     // Ignore private variables and arrays - there is no mapping for them.
@@ -664,9 +664,9 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
 
     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));
+    bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
+                      (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) &&
+                     !(FromMapper && I == 0);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
@@ -717,10 +717,8 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
     // 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)) {
+    if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
+        !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
@@ -755,6 +753,10 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
           }
         }
       }
+      if (DelEntry && FromMapper && I == 0) {
+        DelEntry = false;
+        FromMapperBase = HstPtrBegin;
+      }
 
       // If we copied back to the host a struct/array containing pointers, we
       // need to restore the original host pointer values from their shadow
@@ -810,6 +812,8 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
 
   // Deallocate target pointer
   for (DeallocTgtPtrInfo &Info : DeallocTgtPtrs) {
+    if (FromMapperBase && FromMapperBase == Info.HstPtrBegin)
+      continue;
     Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
                                Info.ForceDelete, Info.HasCloseModifier);
     if (Ret != OFFLOAD_SUCCESS) {

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
new file mode 100644
index 0000000000000..9274c493e6d0d
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
@@ -0,0 +1,70 @@
+// 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
+}

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp
new file mode 100644
index 0000000000000..e7242aceb1d58
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp
@@ -0,0 +1,60 @@
+// 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
+}

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
new file mode 100644
index 0000000000000..968e2dfb2365e
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
@@ -0,0 +1,129 @@
+// 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");
+  for (int i = 0; i < N; ++i)
+    for (int j = 0; j < N; ++j)
+      for (int k = 0; k < N; ++k) {
+        printf("outer[%d].arr[%d].arr[%d].data1 = %d.\n", i, j, k,
+               outer[i].arr[j].arr[k].data1);
+        printf("outer[%d].arr[%d].arr[%d].data2 = %d.\n", i, j, k,
+               outer[i].arr[j].arr[k].data2);
+        assert(outer[i].arr[j].arr[k].data1 == 11 &&
+               outer[i].arr[j].arr[k].data2 == 22);
+      }
+  // CHECK: outer[0].arr[0].arr[0].data1 = 11.
+  // CHECK: outer[0].arr[0].arr[0].data2 = 22.
+  // CHECK: outer[0].arr[0].arr[1].data1 = 11.
+  // CHECK: outer[0].arr[0].arr[1].data2 = 22.
+  // CHECK: outer[0].arr[1].arr[0].data1 = 11.
+  // CHECK: outer[0].arr[1].arr[0].data2 = 22.
+  // CHECK: outer[0].arr[1].arr[1].data1 = 11.
+  // CHECK: outer[0].arr[1].arr[1].data2 = 22.
+  // CHECK: outer[1].arr[0].arr[0].data1 = 11.
+  // CHECK: outer[1].arr[0].arr[0].data2 = 22.
+  // CHECK: outer[1].arr[0].arr[1].data1 = 11.
+  // CHECK: outer[1].arr[0].arr[1].data2 = 22.
+  // CHECK: outer[1].arr[1].arr[0].data1 = 11.
+  // CHECK: outer[1].arr[1].arr[0].data2 = 22.
+  // CHECK: outer[1].arr[1].arr[1].data1 = 11.
+  // CHECK: outer[1].arr[1].arr[1].data2 = 22.
+  assert(outer[1].arr[1].arr[0].data1 == 11 &&
+         outer[1].arr[1].arr[0].data2 == 22 &&
+         outer[1].arr[1].arr[1].data1 == 11 &&
+         outer[1].arr[1].arr[1].data2 == 22);
+
+  return 0;
+}

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp
new file mode 100644
index 0000000000000..8847919a62354
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp
@@ -0,0 +1,62 @@
+// 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
+}

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp
new file mode 100644
index 0000000000000..20a907ea95888
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp
@@ -0,0 +1,62 @@
+// 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
+}


        


More information about the cfe-commits mailing list