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  }