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
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
|
/* { dg-do run { target { offload_device } } } */
/* OpenMP permits using the interop functions on the device,
but it is not really supported. Hence, check that the stubs
are working. */
#include <stddef.h>
#include <stdlib.h>
#include <string.h>
#include <omp.h>
/* Assert is not available with newlib on the device side, hence... */
#define assert(cond) \
do { \
if (!(cond)) \
{ \
__builtin_printf ("assert failed: " #cond "\n"); \
abort (); \
} \
} \
while (0)
#pragma omp begin declare target
void
target_test ()
{
if (omp_is_initial_device ())
return; /* Already tested in interop-routines-1.c. */
omp_interop_t interop = omp_interop_none;
assert (omp_irc_no_value == 1);
assert (omp_irc_success == 0);
assert (omp_irc_empty == -1);
assert (omp_irc_out_of_range == -2);
assert (omp_irc_type_int == -3);
assert (omp_irc_type_ptr == -4);
assert (omp_irc_type_str == -5);
assert (omp_irc_other == -6);
/* Check values, including invalid values. */
for (omp_interop_rc_t ret_code3 = (omp_interop_rc_t) ((int) omp_irc_other - 1);
ret_code3 <= omp_irc_no_value + 1;
ret_code3 = (omp_interop_rc_t) ((int) ret_code3 + 1))
{
const char *msg = omp_get_interop_rc_desc (interop, ret_code3);
if (ret_code3 < omp_irc_other || ret_code3 > omp_irc_no_value)
/* Assume NULL for an invalid value. */
assert (msg == NULL);
else if (ret_code3 == omp_irc_other)
/* In GCC, this is used on the device side, only, if the
omp_get_interop_{int,ptr,str} were invoked on the host.
see below for a check. */
assert (msg != NULL && strlen (msg) > 5); /* Some sensible message. */
else
/* Assume that omp_get_interop_rc_desc handles all of those and
not only omp_irc_empty (and possibly omp_irc_out_of_range),
which do occur for omp_interop_none. */
assert (msg != NULL && strlen (msg) > 5); /* Some sensible message. */
}
assert (omp_ifr_last >= omp_ifr_hsa);
for (omp_interop_fr_t fr = omp_ifr_cuda; fr <= omp_ifr_last;
fr = (omp_interop_fr_t) ((int) fr + 1))
{
switch (fr)
{
/* Expect the id values from the additional-definition document. */
case omp_ifr_cuda:
if (fr != 1)
abort ();
break;
case omp_ifr_cuda_driver:
if (fr != 2)
abort ();
break;
case omp_ifr_opencl:
if (fr != 3)
abort ();
break;
case omp_ifr_sycl:
if (fr != 4)
abort ();
break;
case omp_ifr_hip:
if (fr != 5)
abort ();
break;
case omp_ifr_level_zero:
if (fr != 6)
abort ();
break;
case omp_ifr_hsa:
if (fr != 7)
abort ();
break;
default:
/* Valid, but unexpected to have more interop types. */
abort ();
}
}
assert (omp_ipr_first <= omp_ipr_targetsync
&& omp_get_num_interop_properties (interop) > omp_ipr_fr_id);
for (omp_interop_property_t ipr = omp_ipr_first;
ipr < omp_get_num_interop_properties (interop);
ipr = (omp_interop_property_t) ((int) ipr + 1))
{
/* As interop == omp_interop_none, NULL is permissible;
nonetheless, require != NULL for the GCC implementation. */
const char *name = omp_get_interop_name (interop, ipr);
if (name == NULL)
abort ();
switch (ipr)
{
case omp_ipr_fr_id:
if (ipr != -1 || !!strcmp (name, "fr_id"))
abort ();
break;
case omp_ipr_fr_name:
if (ipr != -2 || !!strcmp (name, "fr_name"))
abort ();
break;
case omp_ipr_vendor:
if (ipr != -3 || !!strcmp (name, "vendor"))
abort ();
break;
case omp_ipr_vendor_name:
if (ipr != -4 || !!strcmp (name, "vendor_name"))
abort ();
break;
case omp_ipr_device_num:
if (ipr != -5 || !!strcmp (name, "device_num"))
abort ();
break;
case omp_ipr_platform:
if (ipr != -6 || !!strcmp (name, "platform"))
abort ();
break;
case omp_ipr_device:
if (ipr != -7 || !!strcmp (name, "device"))
abort ();
break;
case omp_ipr_device_context:
if (ipr != -8 || !!strcmp (name, "device_context"))
abort ();
break;
case omp_ipr_targetsync:
if (ipr != -9 || !!strcmp (name, "targetsync"))
abort ();
break;
default:
/* Valid, but unexpected to have more interop types,
especially not for interop == omp_interop_none. */
abort ();
}
/* As interop == omp_interop_none, expect NULL. */
if (omp_get_interop_type_desc (interop, ipr) != NULL)
abort ();
omp_interop_rc_t ret_code;
const char *err;
ret_code = omp_irc_success;
omp_intptr_t ival = omp_get_interop_int (interop, ipr, &ret_code);
assert (ret_code == omp_irc_empty); /* As interop == omp_interop_none. */
assert (ival == 0); /* Implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (err != NULL && strlen (err) > 5); /* Some sensible message. */
assert (!strcmp (err, "provided interoperability object is equal to "
"omp_interop_none")); /* GCC implementation choice. */
ival = omp_get_interop_int (interop, ipr, NULL);
assert (ival == 0); /* Implementation choice. */
#ifdef __cplusplus
ival = omp_get_interop_int (interop, ipr);
assert (ival == 0); /* Implementation choice. */
#endif
ret_code = omp_irc_success;
void *ptr = omp_get_interop_ptr (interop, ipr, &ret_code);
assert (ret_code == omp_irc_empty); /* As interop == omp_interop_none. */
assert (ptr == NULL); /* Obvious implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (err != NULL && strlen (err) > 5); /* Some sensible message. */
assert (!strcmp (err, "provided interoperability object is equal to "
"omp_interop_none")); /* GCC implementation choice. */
ptr = omp_get_interop_ptr (interop, ipr, NULL);
assert (ptr == NULL); /* Obvious implementation choice. */
#ifdef __cplusplus
ptr = omp_get_interop_ptr (interop, ipr);
assert (ptr == NULL); /* Obvious implementation choice. */
#endif
ret_code = omp_irc_success;
const char *str = omp_get_interop_str (interop, ipr, &ret_code);
assert (ret_code == omp_irc_empty); /* As interop == omp_interop_none. */
assert (str == NULL); /* Obvious implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (err != NULL && strlen (err) > 5); /* Some sensible message. */
assert (!strcmp (err, "provided interoperability object is equal to "
"omp_interop_none")); /* GCC implementation choice. */
str = omp_get_interop_str (interop, ipr, NULL);
assert (str == NULL); /* Obvious implementation choice. */
#ifdef __cplusplus
str = omp_get_interop_str (interop, ipr);
assert (str == NULL); /* Obvious implementation choice. */
#endif
/* Special case of GCC: For any non-'omp_interop_none' valued interop,
a device-side call to omp_get_interop_{int,ptr,src} will yield
omp_irc_other - with the error message as checked below. */
omp_interop_t interop_not_none_invalid = (omp_interop_t) 0xDEADBEEF;
ret_code = omp_irc_success;
ival = omp_get_interop_int (interop_not_none_invalid, ipr, &ret_code);
assert (ret_code == omp_irc_other);
assert (ival == 0); /* Implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (err != NULL && strlen (err) > 5); /* Some sensible message. */
ret_code = omp_irc_success;
ptr = omp_get_interop_ptr (interop_not_none_invalid, ipr, &ret_code);
assert (ret_code == omp_irc_other);
assert (ptr == NULL); /* Implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (err != NULL && strlen (err) > 5); /* Some sensible message. */
assert (!strcmp (err, "obtaining properties is only supported on the "
"initial device")); /* GCC special case. */
ret_code = omp_irc_success;
str = omp_get_interop_str (interop_not_none_invalid, ipr, &ret_code);
assert (ret_code == omp_irc_other);
assert (str == NULL); /* Implementation choice. */
err = omp_get_interop_rc_desc (interop, ret_code);
assert (!strcmp (err, "obtaining properties is only supported on the "
"initial device")); /* GCC special case. */
}
/* Invalid ipr. */
/* Valid are either omp_irc_empty (due to omp_interop_none) or
omp_irc_out_of_range; assume omp_irc_out_of_range with GCC. */
omp_interop_rc_t ret_code2;
const char *err2;
omp_intptr_t ival2;
void *ptr2;
const char *str2;
/* omp_ipr_targetsync-1, i.e < lower bound. */
ret_code2 = omp_irc_success;
ival2 = omp_get_interop_int (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), &ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (ival2 == 0); /* Implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
ival2 = omp_get_interop_int (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), NULL);
assert (ival2 == 0); /* Implementation choice. */
#ifdef __cplusplus
ival2 = omp_get_interop_int (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1));
assert (ival2 == 0); /* Implementation choice. */
#endif
ret_code2 = omp_irc_success;
ptr2 = omp_get_interop_ptr (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), &ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (ptr2 == NULL); /* Obvious implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
ptr2 = omp_get_interop_ptr (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), NULL);
assert (ptr2 == NULL); /* Obvious implementation choice. */
#ifdef __cplusplus
ptr2 = omp_get_interop_ptr (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1));
assert (ptr2 == NULL); /* Obvious implementation choice. */
#endif
ret_code2 = omp_irc_success;
str2 = omp_get_interop_str (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), &ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (str2 == NULL); /* Obvious implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
str2 = omp_get_interop_str (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1), NULL);
assert (str2 == NULL); /* Obvious implementation choice. */
#ifdef __cplusplus
str2 = omp_get_interop_str (interop,
(omp_interop_property_t) ((int) omp_ipr_targetsync-1));
assert (str2 == NULL); /* Obvious implementation choice. */
#endif
/* omp_get_num_interop_properties (), i.e > upper bound. */
ret_code2 = omp_irc_success;
ival2 = omp_get_interop_int (interop,
(omp_interop_property_t) omp_get_num_interop_properties (interop),
&ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (ival2 == 0); /* Implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
ret_code2 = omp_irc_success;
ptr2 = omp_get_interop_ptr (interop,
(omp_interop_property_t) omp_get_num_interop_properties (interop),
&ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (ptr2 == NULL); /* Obvious implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
ret_code2 = omp_irc_success;
str2 = omp_get_interop_str (interop,
(omp_interop_property_t) omp_get_num_interop_properties (interop),
&ret_code2);
assert (ret_code2 == omp_irc_out_of_range);
assert (str2 == NULL); /* Obvious implementation choice. */
err2 = omp_get_interop_rc_desc (interop, ret_code2);
assert (err2 != NULL && strlen (err2) > 5); /* Some sensible message. */
/* GCC implementation choice. */
assert (!strcmp (err2, "property ID is out of range"));
}
#pragma omp end declare target
int
main ()
{
for (int dev = 0; dev < omp_get_num_devices (); dev++)
#pragma omp target device(device_num : dev)
target_test ();
return 0;
}
|
