1  /* Verify that 'acc_unmap_data' unmaps even in presence of structured and
       2     dynamic reference counts, but the device memory remains allocated.  */
       3  
       4  /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
       5  
       6  #include <assert.h>
       7  #include <stdlib.h>
       8  #include <string.h>
       9  #include <openacc.h>
      10  
      11  int
      12  main ()
      13  {
      14    const int N = 180;
      15    const int N_i = 537;
      16    const int C = 37;
      17  
      18    unsigned char *h = (unsigned char *) malloc (N);
      19    assert (h);
      20    unsigned char *d = (unsigned char *) acc_malloc (N);
      21    assert (d);
      22  
      23    for (int i = 0; i < N_i; ++i)
      24      {
      25        acc_map_data (h, d, N);
      26        assert (acc_is_present (h, N));
      27  #pragma acc parallel present(h[0:N])
      28        {
      29  	if (i == 0)
      30  	  memset (h, C, N);
      31        }
      32  
      33        unsigned char *d_ = (unsigned char *) acc_create (h + 3, N - 77);
      34        assert (d_ == d + 3);
      35  
      36  #pragma acc data create(h[6:N - 44])
      37        {
      38  	d_ = (unsigned char *) acc_create (h, N);
      39  	assert (d_ == d);
      40  
      41  #pragma acc enter data create(h[0:N])
      42  
      43  	assert (acc_is_present (h, N));
      44  	acc_unmap_data (h);
      45  	assert (!acc_is_present (h, N));
      46        }
      47  
      48        /* We can however still access the device memory.  */
      49  #pragma acc parallel loop deviceptr(d)
      50        for (int j = 0; j < N; ++j)
      51  	d[j] += i * j;
      52      }
      53  
      54    acc_memcpy_from_device(h, d, N);
      55    for (int j = 0; j < N; ++j)
      56      assert (h[j] == ((C + N_i * (N_i - 1) / 2 * j) % 256));
      57  
      58    acc_free (d);
      59  
      60    return 0;
      61  }