blob: b2a593d03afaa28815ab6a642fb4b4d98fc709b1 (
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
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
|
// We use 'auto' without a function return type, so specify dialect here
// { dg-additional-options "-std=c++14" }
// { dg-do run { target offload_device_nonshared_as } }
#include <cstdlib>
#include <cstring>
#include <cstdint>
struct T
{
int *ptr;
int ptr_len;
int *&refptr;
int refptr_len;
auto set_ptr_func (int n)
{
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
mapped = (ptr != (int *) hostptr);
}
return mapped;
};
return fn;
}
auto set_refptr_func (int n)
{
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
mapped = (refptr != (int *) hostrefptr);
}
return mapped;
};
return fn;
}
};
int main (void)
{
#define N 10
int *ptr1 = new int[N];
int *ptr2 = new int[N];
memset (ptr1, 0, sizeof (int) * N);
memset (ptr2, 0, sizeof (int) * N);
T a = { ptr1, N, ptr2, N };
auto p1 = a.set_ptr_func (1);
auto r2 = a.set_refptr_func (2);
if (p1 ())
abort ();
if (r2 ())
abort ();
if (a.ptr != ptr1)
abort ();
if (a.refptr != ptr2)
abort ();
for (int i = 0; i < N; i++)
if (ptr1[i] != 0)
abort ();
for (int i = 0; i < N; i++)
if (ptr2[i] != 0)
abort ();
#pragma omp target data map(ptr1[:N], ptr2[:N])
{
if (!p1 ())
abort ();
if (!r2 ())
abort ();
}
if (a.ptr != ptr1)
abort ();
if (a.refptr != ptr2)
abort ();
for (int i = 0; i < N; i++)
if (ptr1[i] != 1)
abort ();
for (int i = 0; i < N; i++)
if (ptr2[i] != 2)
abort ();
return 0;
}
|
