|
111
|
1 /* { dg-require-effective-target offload_device_nonshared_as } */
|
|
|
2
|
|
|
3 #include <stdlib.h>
|
|
|
4 #include <assert.h>
|
|
|
5
|
|
|
6 #define N 40
|
|
|
7
|
|
|
8 int sum;
|
|
|
9 int var1 = 1;
|
|
|
10 int var2 = 2;
|
|
|
11
|
|
|
12 #pragma omp declare target
|
|
|
13 int D[N];
|
|
|
14 #pragma omp end declare target
|
|
|
15
|
|
|
16 void enter_data (int *X)
|
|
|
17 {
|
|
|
18 #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
|
|
|
19 }
|
|
|
20
|
|
|
21 void exit_data_0 (int *D)
|
|
|
22 {
|
|
|
23 #pragma omp target exit data map(delete: D[:N])
|
|
|
24 }
|
|
|
25
|
|
|
26 void exit_data_1 ()
|
|
|
27 {
|
|
|
28 #pragma omp target exit data map(from: var1)
|
|
|
29 }
|
|
|
30
|
|
|
31 void exit_data_2 (int *X)
|
|
|
32 {
|
|
|
33 #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
|
|
|
34 }
|
|
|
35
|
|
|
36 void exit_data_3 (int *p)
|
|
|
37 {
|
|
|
38 #pragma omp target exit data map(from: p[:0])
|
|
|
39 }
|
|
|
40
|
|
|
41 void test_nested ()
|
|
|
42 {
|
|
|
43 int X = 0, Y = 0, Z = 0;
|
|
|
44
|
|
|
45 #pragma omp target data map(from: X, Y, Z)
|
|
|
46 {
|
|
|
47 #pragma omp target data map(from: X, Y, Z)
|
|
|
48 {
|
|
|
49 #pragma omp target map(from: X, Y, Z)
|
|
|
50 X = Y = Z = 1337;
|
|
|
51 assert (X == 0);
|
|
|
52 assert (Y == 0);
|
|
|
53 assert (Z == 0);
|
|
|
54
|
|
|
55 #pragma omp target exit data map(from: X) map(release: Y)
|
|
|
56 assert (X == 0);
|
|
|
57 assert (Y == 0);
|
|
|
58
|
|
|
59 #pragma omp target exit data map(release: Y) map(delete: Z)
|
|
|
60 assert (Y == 0);
|
|
|
61 assert (Z == 0);
|
|
|
62 }
|
|
|
63 assert (X == 1337);
|
|
|
64 assert (Y == 0);
|
|
|
65 assert (Z == 0);
|
|
|
66
|
|
|
67 #pragma omp target map(from: X)
|
|
|
68 X = 2448;
|
|
|
69 assert (X == 2448);
|
|
|
70 assert (Y == 0);
|
|
|
71 assert (Z == 0);
|
|
|
72
|
|
|
73 X = 4896;
|
|
|
74 }
|
|
|
75 assert (X == 4896);
|
|
|
76 assert (Y == 0);
|
|
|
77 assert (Z == 0);
|
|
|
78 }
|
|
|
79
|
|
|
80 int main ()
|
|
|
81 {
|
|
|
82 int *X = malloc (N * sizeof (int));
|
|
|
83 int *Y = malloc (N * sizeof (int));
|
|
|
84 X[10] = 10;
|
|
|
85 Y[20] = 20;
|
|
|
86 enter_data (X);
|
|
|
87
|
|
|
88 exit_data_0 (D); /* This should have no effect on D. */
|
|
|
89
|
|
|
90 #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) map(always from: sum)
|
|
|
91 {
|
|
|
92 var1 += X[10];
|
|
|
93 var2 += Y[20];
|
|
|
94 sum = var1 + var2;
|
|
|
95 D[sum]++;
|
|
|
96 }
|
|
|
97
|
|
|
98 assert (var1 == 1);
|
|
|
99 assert (var2 == 2);
|
|
|
100 assert (sum == 33);
|
|
|
101
|
|
|
102 exit_data_1 ();
|
|
|
103 assert (var1 == 11);
|
|
|
104 assert (var2 == 2);
|
|
|
105
|
|
|
106 /* Increase refcount of already mapped X[0:N]. */
|
|
|
107 #pragma omp target enter data map(alloc: X[16:1])
|
|
|
108
|
|
|
109 exit_data_2 (X);
|
|
|
110 assert (var2 == 22);
|
|
|
111
|
|
|
112 exit_data_3 (X + 5); /* Unmap X[0:N]. */
|
|
|
113
|
|
|
114 free (X);
|
|
|
115 free (Y);
|
|
|
116
|
|
|
117 test_nested ();
|
|
|
118
|
|
|
119 return 0;
|
|
|
120 }
|
