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
|
/* { dg-additional-options "-fdump-tree-gimple" } */
// Do diagnostic check / dump check only;
// Note: this test should work as run-test as well.
#if 0
#include <omp.h>
#else
#ifdef __cplusplus
extern "C" {
#endif
extern int omp_get_default_device ();
extern int omp_get_num_devices ();
#ifdef __cplusplus
}
#endif
#endif
void f(int *x, int *y);
#pragma omp declare variant(f) adjust_args(need_device_ptr: x, y) match(construct={dispatch})
void g(int *x, int *y);
void
sub (int *a, int *b)
{
// The has_device_addr is a bit questionable as the caller is not actually
// passing a device address - but we cannot pass one because of the
// following:
//
// As for 'b' need_device_ptr has been specified and 'b' is not
// in the semantic requirement set 'is_device_ptr' (and only in 'has_device_addr')
// "the argument is converted in the same manner that a use_device_ptr clause
// on a target_data construct converts its pointer"
#pragma omp dispatch is_device_ptr(a), has_device_addr(b) /* { dg-warning "'has_device_addr' for 'b' does not imply 'is_device_ptr' required for 'need_device_ptr' \\\[-Wopenmp\\\]" } */
g(a, b);
}
void
f(int *from, int *to)
{
static int cnt = 0;
cnt++;
if (cnt >= 3)
{
if (omp_get_default_device () != -1
&& omp_get_default_device () < omp_get_num_devices ())
{
// On offload device but not mapped
if (from != (void *)0L) // Not mapped
__builtin_abort ();
}
else if (from[0] != 5)
__builtin_abort ();
return;
}
#pragma omp target is_device_ptr(from, to)
{
to[0] = from[0] * 10;
to[1] = from[1] * 10;
}
}
int
main ()
{
int A[2], B[2] = {123, 456}, C[1] = {5};
int *p = A;
#pragma omp target enter data map(A, B)
/* Note: We don't add 'use_device_addr(B)' here;
if we do, it will fail with an illegal memory access (why?). */
#pragma omp target data use_device_ptr(p)
{
sub(p, B);
sub(C, B); /* C is not mapped -> 'from' ptr == NULL */
}
#pragma omp target exit data map(A, B)
}
// { dg-final { scan-tree-dump-times "#pragma omp dispatch has_device_addr\\(b\\) is_device_ptr\\(a\\)" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(b" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "f \\(a, D\\.\[0-9\]+\\);" 1 "gimple" } }
|
