1 // RUN: %libomptarget-compilexx-run-and-check-generic 2 3 // Wrong results on amdgpu 4 // XFAIL: amdgcn-amd-amdhsa 5 // XFAIL: amdgcn-amd-amdhsa-oldDriver 6 // XFAIL: amdgcn-amd-amdhsa-LTO 7 8 #include <cstdio> 9 #include <cstdlib> 10 11 typedef struct { 12 int a; 13 double *b; 14 } C; 15 #pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2]) 16 17 typedef struct { 18 int e; 19 C f; 20 int h; 21 short *g; 22 } D; 23 #pragma omp declare mapper(default \ 24 : D r) map(from \ 25 : r.e) map(mapper(id1), tofrom \ 26 : r.f) map(tofrom \ 27 : r.g [0:r.h]) 28 29 int main() { 30 constexpr int N = 10; 31 D s; 32 s.e = 111; 33 s.f.a = 222; 34 double x[2]; 35 x[1] = 20; 36 short y[N]; 37 y[1] = 30; 38 s.f.b = &x[0]; 39 s.g = &y[0]; 40 s.h = N; 41 42 D *sp = &s; 43 D **spp = &sp; 44 45 printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], 46 spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], 47 spp[0][0].g == &y[0] ? 1 : 0); 48 // CHECK: 111 222 20.00000 1 30 1 49 50 __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]), 51 p1 = reinterpret_cast<__intptr_t>(&y[0]); 52 #pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1) 53 { 54 printf("%d %d %d %d\n", spp[0][0].f.a, 55 spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0, spp[0][0].g[1], 56 spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0); 57 // CHECK: 222 0 30 0 58 spp[0][0].e = 333; 59 spp[0][0].f.a = 444; 60 spp[0][0].f.b[1] = 40; 61 spp[0][0].g[1] = 50; 62 } 63 printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], 64 spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], 65 spp[0][0].g == &y[0] ? 1 : 0); 66 // CHECK: 333 222 40.00000 1 50 1 67 } 68