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
|
// { dg-do run }
#include <cstdlib>
#include <cstring>
#include <cstdint>
template <typename L>
void
omp_target_loop (int begin, int end, L loop)
{
#pragma omp target teams distribute parallel for
for (int i = begin; i < end; i++)
loop (i);
}
struct S
{
int a, len;
int *ptr;
auto merge_data_func (int *iptr, int &b)
{
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostptr = (uintptr_t) ptr;
uintptr_t hostiptr = (uintptr_t) iptr;
#pragma omp target map(from:mapped)
{
mapped = (ptr != (int*) hostptr && iptr != (int*) hostiptr);
if (mapped)
{
for (int i = 0; i < len; i++)
ptr[i] += a + b + iptr[i];
}
}
return mapped;
};
return fn;
}
};
int x = 1;
int main (void)
{
const int N = 10;
int *data1 = new int[N];
int *data2 = new int[N];
memset (data1, 0xab, sizeof (int) * N);
memset (data2, 0xcd, sizeof (int) * N);
bool shared_mem = false;
#pragma omp target map(to: shared_mem)
shared_mem = true;
int val = 1;
int &valref = val;
#pragma omp target enter data map(alloc: data1[:N], data2[:N])
omp_target_loop (0, N, [=](int i) { data1[i] = val; });
omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
#pragma omp target update from(data1[:N], data2[:N])
for (int i = 0; i < N; i++)
{
if (data1[i] != 1) abort ();
if (data2[i] != 2) abort ();
}
#pragma omp target exit data map(delete: data1[:N], data2[:N])
int b = 8;
S s = { 4, N, data1 };
auto f = s.merge_data_func (data2, b);
if (f ()) abort ();
#pragma omp target enter data map(to: data1[:N])
if (f ()) abort ();
#pragma omp target enter data map(to: data2[:N])
if (!f () && !shared_mem) abort ();
#pragma omp target exit data map(from: data1[:N], data2[:N])
if (!shared_mem)
for (int i = 0; i < N; i++)
{
/* With shared memory, data1 is not modified inside 'f'
as mapped = false. */
if (!shared_mem && data1[i] != 0xf) abort ();
if (data2[i] != 2) abort ();
}
return 0;
}
|
