blob: 52d828caf1cf0f4c250586116c3bbf19f59b1869 (
plain)
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
|
/* { dg-do run } */
/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
/* { dg-additional-sources reverse-offload-1-aux.c } */
/* Check that reverse offload works in particular:
- no code is generated on the device side (i.e. no
implicit declare target of called functions and no
code gen for the target-region body)
-> would otherwise fail due to 'add_3' symbol
- Plus the usual (compiles, runs, produces correct result)
Note: Running also the non-reverse-offload target regions
on the host (host fallback) is valid and will pass. */
#pragma omp requires reverse_offload
extern int add_3 (int);
static int global_var = 5;
void
check_offload (int *x, int *y)
{
*x = add_3 (*x);
*y = add_3 (*y);
}
#pragma omp declare target
void
tg_fn (int *x, int *y)
{
int x2 = *x, y2 = *y;
if (x2 != 2 || y2 != 3)
__builtin_abort ();
x2 = x2 + 2;
y2 = y2 + 7;
#pragma omp target device(ancestor : 1) map(tofrom: x2)
check_offload(&x2, &y2);
if (x2 != 2+2+3 || y2 != 3 + 7)
__builtin_abort ();
*x = x2, *y = y2;
}
#pragma omp end declare target
void
my_func (int *x, int *y)
{
if (global_var != 5)
__builtin_abort ();
global_var = 242;
*x = 2*add_3(*x);
*y = 3*add_3(*y);
}
int
main ()
{
#pragma omp target
{
int x = 2, y = 3;
tg_fn (&x, &y);
}
#pragma omp target
{
int x = -2, y = -1;
#pragma omp target device ( ancestor:1 ) firstprivate(y) map(tofrom:x)
{
if (x != -2 || y != -1)
__builtin_abort ();
my_func (&x, &y);
if (x != 2*(3-2) || y != 3*(3-1))
__builtin_abort ();
}
if (x != 2*(3-2) || y != -1)
__builtin_abort ();
}
if (global_var != 242)
__builtin_abort ();
return 0;
}
|
