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
21 printf("\t\tObject A Contents:\n");
22 printf("\t\t\tdata1 = %d data2 = %d\n", data1
, data2
);
35 arr
= new MyObjectA
[N
];
39 printf("\tObject B Contents:\n");
40 for (int i
= 0; i
< len
; i
++)
44 for (int i
= 0; i
< len
; i
++)
50 #pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[ : obj.len])
55 arr
= new MyObjectB
[N
];
59 printf("Object C Contents:\n");
60 for (int i
= 0; i
< len
; i
++)
64 for (int i
= 0; i
< len
; i
++)
70 #pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[ : obj.len])
73 MyObjectC
*outer
= new MyObjectC
[N
];
75 printf("Original data hierarchy:\n");
76 for (int i
= 0; i
< N
; i
++)
79 printf("Sending data to device...\n");
80 #pragma omp target enter data map(to : outer[ : N])
82 printf("Calling foo()...\n");
83 #pragma omp target teams distribute parallel for
84 for (int i
= 0; i
< N
; i
++)
87 printf("foo() complete!\n");
89 printf("Sending data back to host...\n");
90 #pragma omp target exit data map(from : outer[ : N])
92 printf("Modified Data Hierarchy:\n");
93 for (int i
= 0; i
< N
; i
++)
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);
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);