offload/test/mapping/target_has_device_addr.c

   1 // RUN: %libomptarget-compile-generic -fopenmp-version=51
   2 // RUN: %libomptarget-run-generic 2>&1 \
   3 // RUN: | %fcheck-generic
   4
   5 // UNSUPPORTED: amdgcn-amd-amdhsa
   6
   7 #include <omp.h>
   8 #include <stdio.h>
   9 #include <stdlib.h>
  10
  11 #define N 1024
  12 #define FROM 64
  13 #define LENGTH 128
  14
  15 void foo() {
  16   const int device_id = omp_get_default_device();
  17   float *A;
  18   A = (float *)omp_target_alloc((FROM + LENGTH) * sizeof(float), device_id);
  19
  20   float *A_dev = NULL;
  21 #pragma omp target has_device_addr(A[FROM : LENGTH]) map(A_dev)
  22   { A_dev = A; }
  23   // CHECK: Success
  24   if (A_dev == NULL || A_dev != A)
  25     fprintf(stderr, "Failure %p %p \n", A_dev, A);
  26   else
  27     fprintf(stderr, "Success\n");
  28 }
  29
  30 void bar() {
  31   short x[10];
  32   short *xp = &x[0];
  33
  34   x[1] = 111;
  35 #pragma omp target data map(tofrom : xp[0 : 2]) use_device_addr(xp[0 : 2])
  36 #pragma omp target has_device_addr(xp[0 : 2])
  37   {
  38     xp[1] = 222;
  39     // CHECK: 222
  40     printf("%d %p\n", xp[1], &xp[1]);
  41   }
  42   // CHECK: 222
  43   printf("%d %p\n", xp[1], &xp[1]);
  44 }
  45
  46 void moo() {
  47   short *b = malloc(sizeof(short));
  48   b = b - 1;
  49
  50   b[1] = 111;
  51 #pragma omp target data map(tofrom : b[1]) use_device_addr(b[1])
  52 #pragma omp target has_device_addr(b[1])
  53   {
  54     b[1] = 222;
  55     // CHECK: 222
  56     printf("%hd %p %p %p\n", b[1], b, &b[1], &b);
  57   }
  58   // CHECK: 222
  59   printf("%hd %p %p %p\n", b[1], b, &b[1], &b);
  60 }
  61
  62 void zoo() {
  63   short x[10];
  64   short *(xp[10]);
  65   xp[1] = &x[0];
  66   short **xpp = &xp[0];
  67
  68   x[1] = 111;
  69 #pragma omp target data map(tofrom : xpp[1][1]) use_device_addr(xpp[1][1])
  70 #pragma omp target has_device_addr(xpp[1][1])
  71   {
  72     xpp[1][1] = 222;
  73     // CHECK: 222
  74     printf("%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
  75   }
  76   // CHECK: 222
  77   printf("%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
  78 }
  79 void xoo() {
  80   short a[10], b[10];
  81   a[1] = 111;
  82   b[1] = 111;
  83 #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
  84 #pragma omp target has_device_addr(a) has_device_addr(b[0])
  85   {
  86     a[1] = 222;
  87     b[1] = 222;
  88     // CHECK: 222 222
  89     printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b);
  90   }
  91   // CHECK:111
  92   printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b); // 111 111 p1d p2d p3d
  93 }
  94 int main() {
  95   foo();
  96   bar();
  97   moo();
  98   zoo();
  99   xoo();
 100   return 0;
 101 }