1 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
2 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
3 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
4 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
5 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
6
7 #include <assert.h>
8 #include <omp.h>
9 #include <stdio.h>
10 #include <stdlib.h>
11
12 #define N 2
13
14 class MyObjectA {
15 public:
MyObjectA()16 MyObjectA() {
17 data1 = 1;
18 data2 = 2;
19 }
show()20 void show() {
21 printf("\t\tObject A Contents:\n");
22 printf("\t\t\tdata1 = %d data2 = %d\n", data1, data2);
23 }
foo()24 void foo() {
25 data1 += 10;
26 data2 += 20;
27 }
28 int data1;
29 int data2;
30 };
31
32 class MyObjectB {
33 public:
MyObjectB()34 MyObjectB() {
35 arr = new MyObjectA[N];
36 len = N;
37 }
show()38 void show() {
39 printf("\tObject B Contents:\n");
40 for (int i = 0; i < len; i++)
41 arr[i].show();
42 }
foo()43 void foo() {
44 for (int i = 0; i < len; i++)
45 arr[i].foo();
46 }
47 MyObjectA *arr;
48 int len;
49 };
50 #pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[:obj.len])
51
52 class MyObjectC {
53 public:
MyObjectC()54 MyObjectC() {
55 arr = new MyObjectB[N];
56 len = N;
57 }
show()58 void show() {
59 printf("Object C Contents:\n");
60 for (int i = 0; i < len; i++)
61 arr[i].show();
62 }
foo()63 void foo() {
64 for (int i = 0; i < len; i++)
65 arr[i].foo();
66 }
67 MyObjectB *arr;
68 int len;
69 };
70 #pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[:obj.len])
71
main(void)72 int main(void) {
73 MyObjectC *outer = new MyObjectC[N];
74
75 printf("Original data hierarchy:\n");
76 for (int i = 0; i < N; i++)
77 outer[i].show();
78
79 printf("Sending data to device...\n");
80 #pragma omp target enter data map(to : outer[:N])
81
82 printf("Calling foo()...\n");
83 #pragma omp target teams distribute parallel for
84 for (int i = 0; i < N; i++)
85 outer[i].foo();
86
87 printf("foo() complete!\n");
88
89 printf("Sending data back to host...\n");
90 #pragma omp target exit data map(from : outer[:N])
91
92 printf("Modified Data Hierarchy:\n");
93 for (int i = 0; i < N; i++)
94 outer[i].show();
95
96 printf("Testing for correctness...\n");
97 for (int i = 0; i < N; ++i)
98 for (int j = 0; j < N; ++j)
99 for (int k = 0; k < N; ++k) {
100 printf("outer[%d].arr[%d].arr[%d].data1 = %d.\n", i, j, k,
101 outer[i].arr[j].arr[k].data1);
102 printf("outer[%d].arr[%d].arr[%d].data2 = %d.\n", i, j, k,
103 outer[i].arr[j].arr[k].data2);
104 assert(outer[i].arr[j].arr[k].data1 == 11 &&
105 outer[i].arr[j].arr[k].data2 == 22);
106 }
107 // CHECK: outer[0].arr[0].arr[0].data1 = 11.
108 // CHECK: outer[0].arr[0].arr[0].data2 = 22.
109 // CHECK: outer[0].arr[0].arr[1].data1 = 11.
110 // CHECK: outer[0].arr[0].arr[1].data2 = 22.
111 // CHECK: outer[0].arr[1].arr[0].data1 = 11.
112 // CHECK: outer[0].arr[1].arr[0].data2 = 22.
113 // CHECK: outer[0].arr[1].arr[1].data1 = 11.
114 // CHECK: outer[0].arr[1].arr[1].data2 = 22.
115 // CHECK: outer[1].arr[0].arr[0].data1 = 11.
116 // CHECK: outer[1].arr[0].arr[0].data2 = 22.
117 // CHECK: outer[1].arr[0].arr[1].data1 = 11.
118 // CHECK: outer[1].arr[0].arr[1].data2 = 22.
119 // CHECK: outer[1].arr[1].arr[0].data1 = 11.
120 // CHECK: outer[1].arr[1].arr[0].data2 = 22.
121 // CHECK: outer[1].arr[1].arr[1].data1 = 11.
122 // CHECK: outer[1].arr[1].arr[1].data2 = 22.
123 assert(outer[1].arr[1].arr[0].data1 == 11 &&
124 outer[1].arr[1].arr[0].data2 == 22 &&
125 outer[1].arr[1].arr[1].data1 == 11 &&
126 outer[1].arr[1].arr[1].data2 == 22);
127
128 return 0;
129 }
130