offload/test/mapping/map_both_pointer_pointee.c

   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 // REQUIRES: unified_shared_memory
   8 // UNSUPPORTED: amdgcn-amd-amdhsa
   9
  10 #pragma omp declare target
  11 int *ptr1;
  12 #pragma omp end declare target
  13 int a[10];
  14
  15 #include <stdio.h>
  16 #include <stdlib.h>
  17 int main() {
  18   ptr1 = (int *)malloc(sizeof(int) * 100);
  19   int *ptr2;
  20   ptr2 = (int *)malloc(sizeof(int) * 100);
  21 #pragma omp target map(ptr1, ptr1[ : 100])
  22   { ptr1[1] = 6; }
  23   // CHECK: 6
  24   printf(" %d \n", ptr1[1]);
  25 #pragma omp target data map(ptr1[ : 5])
  26   {
  27 #pragma omp target map(ptr1[2], ptr1, ptr1[3]) map(ptr2, ptr2[2])
  28     {
  29       ptr1[2] = 7;
  30       ptr1[3] = 9;
  31       ptr2[2] = 7;
  32     }
  33   }
  34   // CHECK: 7 7 9
  35   printf(" %d %d %d \n", ptr2[2], ptr1[2], ptr1[3]);
  36   free(ptr1);
  37 #pragma omp target map(ptr2, ptr2[ : 100])
  38   { ptr2[1] = 6; }
  39   // CHECK: 6
  40   printf(" %d \n", ptr2[1]);
  41   free(ptr2);
  42
  43   a[1] = 111;
  44   int *p = &a[0];
  45   // CHECK: 111
  46   printf("%d %p %p\n", p[1], p, &p); // 111 hst_p1 hst_p2
  47 #pragma omp target data map(to : p[1 : 3]) map(p)
  48 #pragma omp target data use_device_addr(p)
  49   {
  50 #pragma omp target has_device_addr(p)
  51     {
  52       // CHECK: 111
  53       printf("%d %p %p\n", p[1], p, &p); // 111 dev_p1 dev_p2
  54       p[1] = 222;
  55       // CHECK: 222
  56       printf("%d %p %p\n", p[1], p, &p); // 222 dev_p1 dev_p2
  57     }
  58   }
  59   // CHECK: 111
  60   printf("%d %p %p\n", p[1], p, &p); // 111 hst_p1 hst_p2
  61   return 0;
  62 }