[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