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
|
// We use 'auto' without a function return type, so specify dialect here
// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
#include <cstdlib>
#include <cstring>
#include <omp.h>
template <typename L>
void
omp_target_loop (int begin, int end, L loop, int dev)
{
#pragma omp target teams distribute parallel for device(dev)
for (int i = begin; i < end; i++)
loop (i);
}
struct S
{
int a, len;
int *ptr;
auto merge_data_func (int *iptr, int &b, int dev)
{
auto fn = [=](void) -> bool
{
bool mapped = (omp_target_is_present (iptr, dev)
&& omp_target_is_present (ptr, dev));
#pragma omp target device(dev)
{
if (mapped)
{
for (int i = 0; i < len; i++)
ptr[i] += a + b + iptr[i];
}
}
return mapped;
};
return fn;
}
};
int x = 1;
void run (int dev)
{
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 = (omp_target_is_present (data1, dev)
&& omp_target_is_present (data2, dev));
int val = 1;
int &valref = val;
#pragma omp target enter data map(alloc: data1[ :N], data2[ :N]) device(dev)
omp_target_loop (0, N, [=](int i) { data1[i] = val; }, dev);
omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }, dev);
#pragma omp target update from(data1[ :N], data2[ :N]) device(dev)
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]) device(dev)
int b = 8;
S s = { 4, N, data1 };
auto f = s.merge_data_func (data2, b, dev);
if (f () ^ shared_mem) abort ();
#pragma omp target enter data map(to: data1[ :N]) device(dev)
if (f () ^ shared_mem) abort ();
#pragma omp target enter data map(to: data2[ :N]) device(dev)
if (!f ()) abort ();
#pragma omp target exit data map(from: data1[ :N], data2[ :N]) device(dev)
for (int i = 0; i < N; i++)
{
if ((!shared_mem && data1[i] != 0xf)
|| (shared_mem && data1[i] != 0x2b))
abort ();
if (data2[i] != 2) abort ();
}
delete [] data1;
delete [] data2;
}
int main ()
{
for (int dev = 0; dev <= omp_get_num_devices (); dev++)
run (dev);
}
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) firstprivate\(mapped\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(_[0-9]+\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */
|
