offload/test/api/omp_target_memcpy_rect_async2.c

   1 // RUN: %libomptarget-compile-and-run-generic
   2
   3 #include <omp.h>
   4 #include <stdlib.h>
   5
   6 #define NUM_DIMS 3
   7
   8 int main() {
   9   int d = omp_get_default_device();
  10   int id = omp_get_initial_device();
  11   int a[128], b[64], c[128], e[16], q[128], i;
  12   void *p;
  13
  14   if (d < 0 || d >= omp_get_num_devices())
  15     d = id;
  16
  17   p = omp_target_alloc(130 * sizeof(int), d);
  18   if (p == NULL)
  19     return 0;
  20
  21   for (i = 0; i < 128; i++)
  22     q[i] = 0;
  23   if (omp_target_memcpy(p, q, 128 * sizeof(int), 0, 0, d, id) != 0)
  24     abort();
  25
  26   size_t volume[NUM_DIMS] = {2, 2, 3};
  27   size_t dst_offsets[NUM_DIMS] = {0, 0, 0};
  28   size_t src_offsets[NUM_DIMS] = {0, 0, 0};
  29   size_t dst_dimensions[NUM_DIMS] = {3, 4, 5};
  30   size_t src_dimensions[NUM_DIMS] = {2, 3, 4};
  31
  32   for (i = 0; i < 128; i++)
  33     a[i] = 42;
  34   for (i = 0; i < 64; i++)
  35     b[i] = 24;
  36   for (i = 0; i < 128; i++)
  37     c[i] = 0;
  38   for (i = 0; i < 16; i++)
  39     e[i] = 77;
  40
  41   omp_depend_t obj[2];
  42
  43 #pragma omp parallel num_threads(5)
  44 #pragma omp single
  45   {
  46 #pragma omp task depend(out : p)
  47     omp_target_memcpy(p, a, 128 * sizeof(int), 0, 0, d, id);
  48
  49 #pragma omp task depend(inout : p)
  50     omp_target_memcpy(p, b, 64 * sizeof(int), 0, 0, d, id);
  51
  52 #pragma omp task depend(out : c)
  53     for (i = 0; i < 128; i++)
  54       c[i] = i + 1;
  55
  56 #pragma omp depobj(obj[0]) depend(inout : p)
  57 #pragma omp depobj(obj[1]) depend(in : c)
  58
  59     /*  This produces: 1 2 3 - - 5 6 7 - - at positions 0..9 and
  60         13 14 15 - - 17 18 19 - - at positions 20..29.  */
  61     omp_target_memcpy_rect_async(p, c, sizeof(int), NUM_DIMS, volume,
  62                                  dst_offsets, src_offsets, dst_dimensions,
  63                                  src_dimensions, d, id, 2, obj);
  64
  65 #pragma omp task depend(in : p)
  66     omp_target_memcpy(p, e, 16 * sizeof(int), 0, 0, d, id);
  67   }
  68
  69 #pragma omp taskwait
  70
  71   if (omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d) != 0)
  72     abort();
  73
  74   for (i = 0; i < 16; ++i)
  75     if (q[i] != 77)
  76       abort();
  77   if (q[20] != 13 || q[21] != 14 || q[22] != 15 || q[25] != 17 || q[26] != 18 ||
  78       q[27] != 19)
  79     abort();
  80   for (i = 28; i < 64; ++i)
  81     if (q[i] != 24)
  82       abort();
  83   for (i = 64; i < 128; ++i)
  84     if (q[i] != 42)
  85       abort();
  86
  87   omp_target_free(p, d);
  88   return 0;
  89 }