1  /* Test "subset" subarray mappings.  */
       2  
       3  /* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
       4  
       5  #include <openacc.h>
       6  #include <acc_prof.h>
       7  #include <stdbool.h>
       8  #include <stdint.h>
       9  #include <stdlib.h>
      10  #include <assert.h>
      11  
      12  
      13  static bool cb_ev_alloc_expected;
      14  static size_t cb_ev_alloc_bytes;
      15  static const void *cb_ev_alloc_device_ptr;
      16  static void
      17  cb_ev_alloc (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
      18  {
      19    assert (cb_ev_alloc_expected);
      20    cb_ev_alloc_expected = false;
      21  
      22    cb_ev_alloc_bytes = event_info->data_event.bytes;
      23    cb_ev_alloc_device_ptr = event_info->data_event.device_ptr;
      24  }
      25  
      26  static bool cb_ev_free_expected;
      27  static const void *cb_ev_free_device_ptr;
      28  static void
      29  cb_ev_free (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
      30  {
      31    assert (cb_ev_free_expected);
      32    cb_ev_free_expected = false;
      33  
      34    cb_ev_free_device_ptr = event_info->data_event.device_ptr;
      35  }
      36  
      37  
      38  /* Match the alignment processing that
      39     'libgomp/target.c:gomp_map_vars_internal' is doing; simplified, not
      40     considering special alignment requirements of certain data types.  */
      41  
      42  static size_t
      43  aligned_size (size_t tgt_size)
      44  {
      45    size_t tgt_align = sizeof (void *);
      46    return tgt_size + tgt_align - 1;
      47  }
      48  
      49  static const void *
      50  aligned_address (const void *tgt_start)
      51  {
      52    size_t tgt_align = sizeof (void *);
      53    return (void *) (((uintptr_t) tgt_start + tgt_align - 1) & ~(tgt_align - 1));
      54  }
      55  
      56  
      57  #define SIZE 1024
      58  
      59  
      60  int
      61  main ()
      62  {
      63    cb_ev_alloc_expected = false;
      64    cb_ev_free_expected = false;
      65    acc_prof_register (acc_ev_alloc, cb_ev_alloc, acc_reg);
      66    acc_prof_register (acc_ev_free, cb_ev_free, acc_reg);
      67  
      68    char *block1 = (char *) malloc (SIZE);
      69    char *block2 = (char *) malloc (SIZE);
      70    char *block3 = (char *) malloc (SIZE);
      71    cb_ev_alloc_expected = true;
      72  #pragma acc data create (block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
      73    {
      74      void *s_block1_d = acc_deviceptr (&block1[1]);
      75      void *s_block2_d = acc_deviceptr (&block2[20]);
      76      void *s_block3_d = acc_deviceptr (&block3[300]);
      77      assert (!cb_ev_alloc_expected);
      78      /* 'block1', 'block2', 'block3' get mapped in one device memory object, in
      79         reverse order.  */
      80      assert (cb_ev_alloc_bytes == aligned_size (3 * SIZE));
      81      assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 2 * SIZE + 1) == s_block1_d);
      82      assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 1 * SIZE + 20) == s_block2_d);
      83      assert ((void *) ((uintptr_t) aligned_address (cb_ev_alloc_device_ptr) + 0 * SIZE + 300) == s_block3_d);
      84  
      85      void *s_block1_p_d = acc_pcopyin (&block1[1], SIZE - 3);
      86      void *s_block2_p_d = acc_pcopyin (&block2[20], SIZE - 33);
      87      void *s_block3_p_d = acc_pcopyin (&block3[300], SIZE - 333);
      88      assert (s_block1_p_d == s_block1_d);
      89      assert (s_block2_p_d == s_block2_d);
      90      assert (s_block3_p_d == s_block3_d);
      91  
      92      acc_delete (block1, SIZE);
      93      acc_delete (block2, SIZE);
      94      acc_delete (block3, SIZE);
      95      assert (acc_is_present (block1, SIZE));
      96      assert (acc_is_present (block2, SIZE));
      97      assert (acc_is_present (block3, SIZE));
      98  
      99      cb_ev_free_expected = true;
     100    }
     101    assert (!cb_ev_free_expected);
     102    assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
     103    assert (!acc_is_present (block1, SIZE));
     104    assert (!acc_is_present (block2, SIZE));
     105    assert (!acc_is_present (block3, SIZE));
     106  
     107    free (block1);
     108    free (block2);
     109    free (block3);
     110  
     111    acc_prof_unregister (acc_ev_alloc, cb_ev_alloc, acc_reg);
     112    acc_prof_unregister (acc_ev_free, cb_ev_free, acc_reg);
     113  
     114    return 0;
     115  }