1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
|
// RUN: %libomptarget-compilexx-generic -fopenmp-version=51
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// FIXME: This is currently broken on all GPU targets
// UNSUPPORTED: amdgcn-amd-amdhsa
// UNSUPPORTED: nvptx64-nvidia-cuda
// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
#include <stdio.h>
#include <stdlib.h>
void foo(int **t1d) {
int ***t2d = &t1d;
int ****t3d = &t2d;
*t1d = (int *)malloc(3 * sizeof(int));
int j, a = 0, b = 0;
for (j = 0; j < 3; j++)
(*t1d)[j] = 0;
#pragma omp target map(tofrom : (*t1d)[0 : 3]) map(alloc : *t1d)
{ (*t1d)[1] = 1; }
// CHECK: 1
printf("%d\n", (*t1d)[1]);
#pragma omp target map(tofrom : (**t2d)[0 : 3]) map(alloc : **t2d, *t2d)
{ (**t2d)[1] = 2; }
// CHECK: 2
printf("%d\n", (**t2d)[1]);
#pragma omp target map(tofrom : (***t3d)[0 : 3]) \
map(alloc : ***t3d, **t3d, *t3d)
{ (***t3d)[1] = 3; }
// CHECK: 3
printf("%d\n", (***t3d)[1]);
#pragma omp target map(tofrom : (**t1d)) map(alloc : *t1d)
{ (*t1d)[0] = 4; }
// CHECK: 4
printf("%d\n", (*t1d)[0]);
#pragma omp target map(tofrom : (*(*(t1d + a) + b))) map(to : *(t1d + a))
{ *(*(t1d + a) + b) = 5; }
// CHECK: 5
printf("%d\n", *(*(t1d + a) + b));
}
typedef int(T)[3];
void bar() {
T **a;
int b[2][3];
int(*p)[3] = b;
a = &p;
for (int i = 0; i < 3; i++) {
(**a)[1] = i;
}
#pragma omp target map((**a)[ : 3]) map(alloc : **a, *a)
{
(**a)[1] = 6;
// CHECK: 6
printf("%d\n", (**a)[1]);
}
}
struct SSA {
int i;
SSA *sa;
SSA() {
i = 1;
sa = this;
}
};
void zoo(int **f, SSA *sa) {
int *t = *f;
f = (int **)malloc(sa->i * 4 * sizeof(int));
t = (int *)malloc(sa->i * sizeof(int));
*(f + sa->i + 1) = t;
*(sa->sa->i + *(f + sa->i + 1)) = 4;
printf("%d\n", *(sa->sa->i + *(1 + sa->i + f)));
#pragma omp target map(*(sa->sa->i + *(1 + sa->i + f))) map(alloc : sa->sa) \
map(to : sa->i) map(to : sa->sa->i) map(to : *(1 + sa->i + f))
{ *(sa->sa->i + *(1 + sa->i + f)) = 7; }
// CHECK: 7
printf("%d\n", *(sa->sa->i + *(1 + sa->i + f)));
}
void xoo() {
int *x = 0;
SSA *sa = new SSA();
zoo(&x, sa);
}
void yoo(int **x) {
*x = (int *)malloc(2 * sizeof(int));
#pragma omp target map(**x) map(alloc : *x)
{
**x = 8;
// CHECK: 8
printf("%d\n", **x);
}
#pragma omp target map(*(*x + 1)) map(alloc : *x)
{
*(*x + 1) = 9;
// CHECK: 9
printf("%d\n", *(*x + 1));
}
}
int main() {
int *data = 0;
foo(&data);
bar();
xoo();
yoo(&data);
}
|
