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: |
16 | MyObjectA() { |
17 | data1 = 1; |
18 | data2 = 2; |
19 | } |
20 | void show() { |
21 | printf(format: "\t\tObject A Contents:\n" ); |
22 | printf(format: "\t\t\tdata1 = %d data2 = %d\n" , data1, data2); |
23 | } |
24 | void foo() { |
25 | data1 += 10; |
26 | data2 += 20; |
27 | } |
28 | int data1; |
29 | int data2; |
30 | }; |
31 | |
32 | class MyObjectB { |
33 | public: |
34 | MyObjectB() { |
35 | arr = new MyObjectA[N]; |
36 | len = N; |
37 | } |
38 | void show() { |
39 | printf(format: "\tObject B Contents:\n" ); |
40 | for (int i = 0; i < len; i++) |
41 | arr[i].show(); |
42 | } |
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: |
54 | MyObjectC() { |
55 | arr = new MyObjectB[N]; |
56 | len = N; |
57 | } |
58 | void show() { |
59 | printf(format: "Object C Contents:\n" ); |
60 | for (int i = 0; i < len; i++) |
61 | arr[i].show(); |
62 | } |
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 | |
72 | int main(void) { |
73 | MyObjectC *outer = new MyObjectC[N]; |
74 | |
75 | printf(format: "Original data hierarchy:\n" ); |
76 | for (int i = 0; i < N; i++) |
77 | outer[i].show(); |
78 | |
79 | printf(format: "Sending data to device...\n" ); |
80 | #pragma omp target enter data map(to : outer[ : N]) |
81 | |
82 | printf(format: "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(format: "foo() complete!\n" ); |
88 | |
89 | printf(format: "Sending data back to host...\n" ); |
90 | #pragma omp target exit data map(from : outer[ : N]) |
91 | |
92 | printf(format: "Modified Data Hierarchy:\n" ); |
93 | for (int i = 0; i < N; i++) |
94 | outer[i].show(); |
95 | |
96 | printf(format: "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(format: "outer[%d].arr[%d].arr[%d].data1 = %d.\n" , i, j, k, |
101 | outer[i].arr[j].arr[k].data1); |
102 | printf(format: "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 | |