libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c

   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 }