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