1 // RUN: %libomptarget-compile-generic -fopenmp-extensions
2 // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
3 
4 // Wrong results on amdgpu
5 // XFAIL: amdgcn-amd-amdhsa
6 // XFAIL: amdgcn-amd-amdhsa-oldDriver
7 // XFAIL: amdgcn-amd-amdhsa-LTO
8 
9 #include <omp.h>
10 #include <stdio.h>
11 
12 #define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
13   printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
14          #Var1, #Var2, #Var3,                                                  \
15          omp_target_is_present(&(Var1), omp_get_default_device()),             \
16          omp_target_is_present(&(Var2), omp_get_default_device()),             \
17          omp_target_is_present(&(Var3), omp_get_default_device()))
18 
19 #define CHECK_VALUES(Var1, Var2)                                               \
20   printf("    values of %s, %s: %d, %d\n",                                     \
21          #Var1, #Var2, (Var1), (Var2))
22 
main()23 int main() {
24   struct S { int i; int j; } s;
25   // CHECK: presence of s, s.i, s.j: 0, 0, 0
26   CHECK_PRESENCE(s, s.i, s.j);
27 
28   // =======================================================================
29   // Check that ompx_hold keeps entire struct present.
30 
31   // -----------------------------------------------------------------------
32   // CHECK-LABEL: check:{{.*}}
33   printf("check: ompx_hold only on first member\n");
34   s.i = 20;
35   s.j = 30;
36   #pragma omp target data map(tofrom: s) map(ompx_hold,tofrom: s.i) \
37                                          map(tofrom: s.j)
38   {
39     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
40     CHECK_PRESENCE(s, s.i, s.j);
41     #pragma omp target map(tofrom: s)
42     {
43       s.i = 21;
44       s.j = 31;
45     }
46     #pragma omp target exit data map(delete: s, s.i)
47     // ompx_hold on s.i applies to all of s.
48     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
49     // CHECK-NEXT: values of s.i, s.j: 20, 30
50     CHECK_PRESENCE(s, s.i, s.j);
51     CHECK_VALUES(s.i, s.j);
52   }
53   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
54   // CHECK-NEXT: values of s.i, s.j: 21, 31
55   CHECK_PRESENCE(s, s.i, s.j);
56   CHECK_VALUES(s.i, s.j);
57 
58   // -----------------------------------------------------------------------
59   // CHECK-LABEL: check:{{.*}}
60   printf("check: ompx_hold only on last member\n");
61   s.i = 20;
62   s.j = 30;
63   #pragma omp target data map(tofrom: s) map(tofrom: s.i) \
64                                          map(ompx_hold,tofrom: s.j)
65   {
66     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
67     CHECK_PRESENCE(s, s.i, s.j);
68     #pragma omp target map(tofrom: s)
69     {
70       s.i = 21;
71       s.j = 31;
72     }
73     #pragma omp target exit data map(delete: s, s.i)
74     // ompx_hold on s.j applies to all of s.
75     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
76     // CHECK-NEXT: values of s.i, s.j: 20, 30
77     CHECK_PRESENCE(s, s.i, s.j);
78     CHECK_VALUES(s.i, s.j);
79   }
80   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
81   // CHECK-NEXT: values of s.i, s.j: 21, 31
82   CHECK_PRESENCE(s, s.i, s.j);
83   CHECK_VALUES(s.i, s.j);
84 
85   // -----------------------------------------------------------------------
86   // CHECK-LABEL: check:{{.*}}
87   printf("check: ompx_hold only on struct\n");
88   s.i = 20;
89   s.j = 30;
90   #pragma omp target data map(ompx_hold,tofrom: s) map(tofrom: s.i) \
91                                               map(tofrom: s.j)
92   {
93     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
94     CHECK_PRESENCE(s, s.i, s.j);
95     #pragma omp target map(tofrom: s)
96     {
97       s.i = 21;
98       s.j = 31;
99     }
100     #pragma omp target exit data map(delete: s, s.i)
101     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
102     // CHECK-NEXT: values of s.i, s.j: 20, 30
103     CHECK_PRESENCE(s, s.i, s.j);
104     CHECK_VALUES(s.i, s.j);
105   }
106   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
107   // CHECK-NEXT: values of s.i, s.j: 21, 31
108   CHECK_PRESENCE(s, s.i, s.j);
109   CHECK_VALUES(s.i, s.j);
110 
111   // =======================================================================
112   // Check that transfer to/from host checks reference count correctly.
113 
114   // -----------------------------------------------------------------------
115   // CHECK-LABEL: check:{{.*}}
116   printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
117   s.i = 20;
118   s.j = 30;
119   #pragma omp target data map(ompx_hold, tofrom: s)
120   #pragma omp target data map(ompx_hold, tofrom: s)
121   {
122     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
123     CHECK_PRESENCE(s, s.i, s.j);
124     #pragma omp target map(from: s.i, s.j)
125     {
126       s.i = 21;
127       s.j = 31;
128     } // No transfer here even though parent's DynRefCount=1.
129     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
130     // CHECK-NEXT: values of s.i, s.j: 20, 30
131     CHECK_PRESENCE(s, s.i, s.j);
132     CHECK_VALUES(s.i, s.j);
133     #pragma omp target map(to: s.i, s.j)
134     { // No transfer here even though parent's DynRefCount=1.
135       // CHECK-NEXT: values of s.i, s.j: 21, 31
136       CHECK_VALUES(s.i, s.j);
137     }
138   }
139   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
140   // CHECK-NEXT: values of s.i, s.j: 21, 31
141   CHECK_PRESENCE(s, s.i, s.j);
142   CHECK_VALUES(s.i, s.j);
143 
144   // -----------------------------------------------------------------------
145   // CHECK-LABEL: check:{{.*}}
146   printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
147   s.i = 20;
148   s.j = 30;
149   #pragma omp target data map(tofrom: s)
150   #pragma omp target data map(tofrom: s)
151   {
152     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
153     CHECK_PRESENCE(s, s.i, s.j);
154     #pragma omp target map(ompx_hold, from: s.i, s.j)
155     {
156       s.i = 21;
157       s.j = 31;
158     } // No transfer here even though parent's HoldRefCount=1.
159     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
160     // CHECK-NEXT: values of s.i, s.j: 20, 30
161     CHECK_PRESENCE(s, s.i, s.j);
162     CHECK_VALUES(s.i, s.j);
163     #pragma omp target map(ompx_hold, to: s.i, s.j)
164     { // No transfer here even though parent's HoldRefCount=1.
165       // CHECK-NEXT: values of s.i, s.j: 21, 31
166       CHECK_VALUES(s.i, s.j);
167     }
168   }
169   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
170   // CHECK-NEXT: values of s.i, s.j: 21, 31
171   CHECK_PRESENCE(s, s.i, s.j);
172   CHECK_VALUES(s.i, s.j);
173 
174   // -----------------------------------------------------------------------
175   // CHECK-LABEL: check:{{.*}}
176   //
177   // At the beginning of a region, if the parent's TotalRefCount=1, then the
178   // transfer should happen.
179   //
180   // At the end of a region, it also must be true that the reference count being
181   // decremented is the reference count that is 1.
182   printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
183   s.i = 20;
184   s.j = 30;
185   #pragma omp target data map(ompx_hold, tofrom: s)
186   {
187     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
188     CHECK_PRESENCE(s, s.i, s.j);
189     #pragma omp target map(ompx_hold, tofrom: s.i, s.j)
190     {
191       s.i = 21;
192       s.j = 31;
193     }
194     #pragma omp target exit data map(from: s.i, s.j)
195     // No transfer here even though parent's TotalRefCount=1.
196     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
197     // CHECK-NEXT: values of s.i, s.j: 20, 30
198     CHECK_PRESENCE(s, s.i, s.j);
199     CHECK_VALUES(s.i, s.j);
200   }
201   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
202   // CHECK-NEXT: values of s.i, s.j: 21, 31
203   CHECK_PRESENCE(s, s.i, s.j);
204   CHECK_VALUES(s.i, s.j);
205 
206   return 0;
207 }
208