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
|
// RUN: %libomptarget-compilexx-run-and-check-generic
// Wrong results on amdgpu
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-oldDriver
// XFAIL: amdgcn-amd-amdhsa-LTO
#include <cstdio>
#include <cstdlib>
typedef struct {
int a;
double *b;
} C;
#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2])
typedef struct {
int e;
C f;
int h;
short *g;
} D;
#pragma omp declare mapper(default \
: D r) map(from \
: r.e) map(mapper(id1), tofrom \
: r.f) map(tofrom \
: r.g [0:r.h])
int main() {
constexpr int N = 10;
D s;
s.e = 111;
s.f.a = 222;
double x[2];
x[1] = 20;
short y[N];
y[1] = 30;
s.f.b = &x[0];
s.g = &y[0];
s.h = N;
D *sp = &s;
D **spp = &sp;
printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
spp[0][0].g == &y[0] ? 1 : 0);
// CHECK: 111 222 20.00000 1 30 1
__intptr_t p = reinterpret_cast<__intptr_t>(&x[0]),
p1 = reinterpret_cast<__intptr_t>(&y[0]);
#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)
{
printf("%d %d %d %d\n", spp[0][0].f.a,
spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0, spp[0][0].g[1],
spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0);
// CHECK: 222 0 30 0
spp[0][0].e = 333;
spp[0][0].f.a = 444;
spp[0][0].f.b[1] = 40;
spp[0][0].g[1] = 50;
}
printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
spp[0][0].g == &y[0] ? 1 : 0);
// CHECK: 333 222 40.00000 1 50 1
}
|
