1  /* Verify that OpenACC 'attach'/'detach' doesn't interfere with reference
       2     counting.  */
       3  
       4  #include <assert.h>
       5  #include <stdlib.h>
       6  #include <openacc.h>
       7  
       8  /* Need to shared this (and, in particular, implicit '&data_work' in
       9     'attach'/'detach' clauses) between 'test' and 'test_'.  */
      10  static unsigned char *data_work;
      11  
      12  static void test_(unsigned variant,
      13  		  unsigned char *data,
      14  		  void *data_d)
      15  {
      16    assert(acc_is_present(&data_work, sizeof data_work));
      17    assert(data_work == data);
      18  
      19    acc_update_self(&data_work, sizeof data_work);
      20    assert(data_work == data);
      21  
      22    if (variant & 1)
      23      {
      24  #pragma acc enter data attach(data_work)
      25      }
      26    else
      27      acc_attach((void **) &data_work);
      28    acc_update_self(&data_work, sizeof data_work);
      29    assert(data_work == data_d);
      30  
      31    if (variant & 4)
      32      {
      33        if (variant & 2)
      34  	{ // attach some more
      35  	  data_work = data;
      36  	  acc_attach((void **) &data_work);
      37  #pragma acc enter data attach(data_work)
      38  	  acc_attach((void **) &data_work);
      39  #pragma acc enter data attach(data_work)
      40  #pragma acc enter data attach(data_work)
      41  #pragma acc enter data attach(data_work)
      42  	  acc_attach((void **) &data_work);
      43  	  acc_attach((void **) &data_work);
      44  #pragma acc enter data attach(data_work)
      45  	}
      46        else
      47  	{}
      48      }
      49    else
      50      { // detach
      51        data_work = data;
      52        if (variant & 2)
      53  	{
      54  #pragma acc exit data detach(data_work)
      55  	}
      56        else
      57  	acc_detach((void **) &data_work);
      58        acc_update_self(&data_work, sizeof data_work);
      59        assert(data_work == data);
      60  
      61        // now not attached anymore
      62  
      63  #if 0
      64        if (TODO)
      65  	{
      66  	  acc_detach(&data_work); //TODO PR95203 "libgomp: attach count underflow"
      67  	  acc_update_self(&data_work, sizeof data_work);
      68  	  assert(data_work == data);
      69  	}
      70  #endif
      71      }
      72  
      73    assert(acc_is_present(&data_work, sizeof data_work));
      74  }
      75  
      76  static void test(unsigned variant)
      77  {
      78    const int size = sizeof (void *);
      79    unsigned char *data = (unsigned char *) malloc(size);
      80    assert(data);
      81    void *data_d = acc_create(data, size);
      82    assert(data_d);
      83    assert(acc_is_present(data, size));
      84  
      85    data_work = data;
      86  
      87    if (variant & 8)
      88      {
      89  #pragma acc data copyin(data_work)
      90        test_(variant, data, data_d);
      91      }
      92    else
      93      {
      94        acc_copyin(&data_work, sizeof data_work);
      95        test_(variant, data, data_d);
      96        acc_delete(&data_work, sizeof data_work);
      97      }
      98  #if ACC_MEM_SHARED
      99    assert(acc_is_present(&data_work, sizeof data_work));
     100  #else
     101    assert(!acc_is_present(&data_work, sizeof data_work));
     102  #endif
     103    data_work = NULL;
     104  
     105    assert(acc_is_present(data, size));
     106    acc_delete(data, size);
     107    data_d = NULL;
     108  #if ACC_MEM_SHARED
     109    assert(acc_is_present(data, size));
     110  #else
     111    assert(!acc_is_present(data, size));
     112  #endif
     113    free(data);
     114    data = NULL;
     115  }
     116  
     117  int main()
     118  {
     119    for (size_t i = 0; i < 16; ++i)
     120      test(i);
     121  
     122    return 0;
     123  }