libgomp/testsuite/libgomp.c/target-20.c

   1 /* { dg-require-effective-target offload_device_nonshared_as } */
   2
   3 #include <stdlib.h>
   4 #include <assert.h>
   5
   6 #define N 40
   7
   8 int sum;
   9 int var1 = 1;
  10 int var2 = 2;
  11
  12 #pragma omp declare target
  13 int D[N];
  14 #pragma omp end declare target
  15
  16 void enter_data (int *X)
  17 {
  18   #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
  19 }
  20
  21 void exit_data_0 (int *D)
  22 {
  23   #pragma omp target exit data map(delete: D[:N])
  24 }
  25
  26 void exit_data_1 ()
  27 {
  28   #pragma omp target exit data map(from: var1)
  29 }
  30
  31 void exit_data_2 (int *X)
  32 {
  33   #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
  34 }
  35
  36 void exit_data_3 (int *p)
  37 {
  38   #pragma omp target exit data map(from: p[:0])
  39 }
  40
  41 void test_nested ()
  42 {
  43   int X = 0, Y = 0, Z = 0;
  44
  45   #pragma omp target data map(from: X, Y, Z)
  46     {
  47       #pragma omp target data map(from: X, Y, Z)
  48         {
  49           #pragma omp target map(from: X, Y, Z)
  50             X = Y = Z = 1337;
  51           assert (X == 0);
  52           assert (Y == 0);
  53           assert (Z == 0);
  54
  55           #pragma omp target exit data map(from: X) map(release: Y)
  56           assert (X == 0);
  57           assert (Y == 0);
  58
  59           #pragma omp target exit data map(release: Y) map(delete: Z)
  60           assert (Y == 0);
  61           assert (Z == 0);
  62         }
  63       assert (X == 1337);
  64       assert (Y == 0);
  65       assert (Z == 0);
  66
  67       #pragma omp target map(from: X)
  68         X = 2448;
  69       assert (X == 2448);
  70       assert (Y == 0);
  71       assert (Z == 0);
  72
  73       X = 4896;
  74     }
  75   assert (X == 4896);
  76   assert (Y == 0);
  77   assert (Z == 0);
  78 }
  79
  80 int main ()
  81 {
  82   int *X = malloc (N * sizeof (int));
  83   int *Y = malloc (N * sizeof (int));
  84   X[10] = 10;
  85   Y[20] = 20;
  86   enter_data (X);
  87
  88   exit_data_0 (D); /* This should have no effect on D.  */
  89
  90   #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) map(always from: sum)
  91     {
  92       var1 += X[10];
  93       var2 += Y[20];
  94       sum = var1 + var2;
  95       D[sum]++;
  96     }
  97
  98   assert (var1 == 1);
  99   assert (var2 == 2);
 100   assert (sum == 33);
 101
 102   exit_data_1 ();
 103   assert (var1 == 11);
 104   assert (var2 == 2);
 105
 106   /* Increase refcount of already mapped X[0:N].  */
 107   #pragma omp target enter data map(alloc: X[16:1])
 108
 109   exit_data_2 (X);
 110   assert (var2 == 22);
 111
 112   exit_data_3 (X + 5); /* Unmap X[0:N].  */
 113
 114   free (X);
 115   free (Y);
 116
 117   test_nested ();
 118
 119   return 0;
 120 }