offload/test/mapping/target_pointers_members_map.cpp

   1 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
   2 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
   3 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
   4 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
   5 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
   6
   7 #include <cstdio>
   8 #include <cstdlib>
   9
  10 typedef struct {
  11   short *a;
  12   long d1, d2;
  13 } DV_A;
  14
  15 typedef struct {
  16   DV_A b;
  17   long d3;
  18 } C;
  19
  20 typedef struct {
  21   C *c;
  22   long d4, d5;
  23 } DV_B;
  24
  25 int main() {
  26
  27   short arr1[10] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
  28   short arr2[10] = {20, 31, 22, 23, 24, 25, 26, 27, 28, 29};
  29
  30   C c1[2];
  31   c1[0].b.a = (short *)arr1;
  32   c1[1].b.a = (short *)arr2;
  33   c1[0].b.d1 = 111;
  34
  35   DV_B dvb1;
  36   dvb1.c = (C *)&c1;
  37
  38   // CHECK: 10 111
  39   printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
  40          &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
  41 #pragma omp target map(to : dvb1, dvb1.c[0 : 2])                               \
  42     map(tofrom : dvb1.c[0].b.a[0 : 10], dvb1.c[1].b.a[0 : 10])
  43   {
  44     // CHECK: 10 111
  45     printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
  46            &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
  47     dvb1.c[0].b.a[0] = 333;
  48     dvb1.c[0].b.d1 = 444;
  49   }
  50   // CHECK: 333 111
  51   printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
  52          &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
  53 }