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
|
/* { dg-do run } */
/* { dg-additional-options "-O0" } */
/* Issue showed up in the real world when large data was distributed
over multiple MPI progresses - such that for one process n == 0
happend at run time.
Before map(var[:0]) and map(var[:n]) with n > 0 was handled,
this patch now also handles map(var[:n]) with n == 0.
Failed before with "libgomp: pointer target not mapped for attach". */
/* Here, the base address is shifted - which should have no effect,
but must work as well. */
void
with_offset ()
{
struct S {
int *ptr1, *ptr2;
};
struct S s1, s2;
int *a, *b, *c, *d;
s1.ptr1 = (int *) 0L;
s1.ptr2 = (int *) 0xdeedbeef;
s2.ptr1 = (int *) 0L;
s2.ptr2 = (int *) 0xdeedbeef;
a = (int *) 0L;
b = (int *) 0xdeedbeef;
c = (int *) 0L;
d = (int *) 0xdeedbeef;
int n1, n2, n3, n4;
n1 = n2 = n3 = n4 = 0;
#pragma omp target enter data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
#pragma omp target map(s2.ptr1[4:n1], s2.ptr2[2:n2], c[6:n3], d[9:n4])
{
if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
|| c != (void *) 0L || d != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
{
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target
{
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target exit data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
}
int
main ()
{
struct S {
int *ptr1, *ptr2;
};
struct S s1, s2;
int *a, *b, *c, *d;
s1.ptr1 = (int *) 0L;
s1.ptr2 = (int *) 0xdeedbeef;
s2.ptr1 = (int *) 0L;
s2.ptr2 = (int *) 0xdeedbeef;
a = (int *) 0L;
b = (int *) 0xdeedbeef;
c = (int *) 0L;
d = (int *) 0xdeedbeef;
int n1, n2, n3, n4;
n1 = n2 = n3 = n4 = 0;
#pragma omp target enter data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
#pragma omp target map(s2.ptr1[:n1], s2.ptr2[:n2], c[:n3], d[:n4])
{
if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
|| c != (void *) 0L || d != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
{
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target
{
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
__builtin_abort ();
}
#pragma omp target exit data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
with_offset ();
}
|
