gcc/libgomp/testsuite/libgomp.c-c++-common/target-memcpy-async-2.c

/* Test for omp_target_memcpy_async considering dependence objects.  */

#include <omp.h>
#include <stdlib.h>

int
main ()
{
  int d = omp_get_default_device ();
  int id = omp_get_initial_device ();
  int a[128], b[64], c[32], e[16], q[128], i;
  void *p;

  if (d < 0 || d >= omp_get_num_devices ())
    d = id;

  p = omp_target_alloc (130 * sizeof (int), d);
  if (p == NULL)
    return 0;

  for (i = 0; i < 128; ++i)
    a[i] = i + 1;
  for (i = 0; i < 64; ++i)
    b[i] = i + 2;
  for (i = 0; i < 32; i++)
    c[i] = 0;
  for (i = 0; i < 16; i++)
    e[i] = i + 4;

  omp_depend_t obj[2];

  #pragma omp parallel num_threads(5)
  #pragma omp single
  {
    #pragma omp task depend(out: p)
    omp_target_memcpy (p, a, 128 * sizeof (int), 0, 0, d, id);

    #pragma omp task depend(inout: p)
    omp_target_memcpy (p, b, 64 * sizeof (int), 0, 0, d, id);

    #pragma omp task depend(out: c)
    for (i = 0; i < 32; i++)
      c[i] = i + 3;

    #pragma omp depobj(obj[0]) depend(inout: p)
    #pragma omp depobj(obj[1]) depend(in: c)
    omp_target_memcpy_async (p, c, 32 * sizeof (int), 0, 0, d, id, 2, obj);

    #pragma omp task depend(in: p)
    omp_target_memcpy (p, e, 16 * sizeof (int), 0, 0, d, id);
  }

  #pragma omp taskwait

  for (i = 0; i < 128; ++i)
    q[i] = 0;
  omp_target_memcpy (q, p, 128 * sizeof(int), 0, 0, id, d);
  for (i = 0; i < 16; ++i)
    if (q[i] != i + 4)
      abort ();
  for (i = 16; i < 32; ++i)
    if (q[i] != i + 3)
      abort ();
  for (i = 32; i < 64; ++i)
    if (q[i] != i + 2)
      abort ();
  for (i = 64; i < 128; ++i)
    if (q[i] != i + 1)
      abort ();

  omp_target_free (p, d);

  return 0;
}