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 
enter_data(int * X)16 void enter_data (int *X)
17 {
18   #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
19 }
20 
exit_data_0(int * D)21 void exit_data_0 (int *D)
22 {
23   #pragma omp target exit data map(delete: D[:N])
24 }
25 
exit_data_1()26 void exit_data_1 ()
27 {
28   #pragma omp target exit data map(from: var1)
29 }
30 
exit_data_2(int * X)31 void exit_data_2 (int *X)
32 {
33   #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
34 }
35 
exit_data_3(int * p)36 void exit_data_3 (int *p)
37 {
38   #pragma omp target exit data map(from: p[:0])
39 }
40 
test_nested()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 
main()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 }
121