libgomp/testsuite/libgomp.c/target-map-zero-sized.c

/* { dg-do run } */
/* { dg-additional-options "-O0" }  */

/* Issue showed up in the real world when large data was distributed
   over multiple MPI progresses - such that for one process n == 0
   happend at run time.

   Before  map(var[:0])  and  map(var[:n])  with n > 0 was handled,
   this patch now also handles  map(var[:n]) with n == 0.

   Failed before with "libgomp: pointer target not mapped for attach".  */

/* Here, the base address is shifted - which should have no effect,
   but must work as well.  */
void
with_offset ()
{
  struct S {
     int *ptr1, *ptr2;
  };
  struct S s1, s2;
  int *a, *b, *c, *d;
  s1.ptr1 = (int *) 0L;
  s1.ptr2 = (int *) 0xdeedbeef;
  s2.ptr1 = (int *) 0L;
  s2.ptr2 = (int *) 0xdeedbeef;
  a = (int *) 0L;
  b = (int *) 0xdeedbeef;
  c = (int *) 0L;
  d = (int *) 0xdeedbeef;

  int n1, n2, n3, n4;
  n1 = n2 = n3 = n4 = 0;

  #pragma omp target enter data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])

  #pragma omp target map(s2.ptr1[4:n1], s2.ptr2[2:n2], c[6:n3], d[9:n4])
  {
    if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
	|| c != (void *) 0L || d != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
  {
    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target
  {
    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target exit data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
}

int
main ()
{
  struct S {
     int *ptr1, *ptr2;
  };
  struct S s1, s2;
  int *a, *b, *c, *d;
  s1.ptr1 = (int *) 0L;
  s1.ptr2 = (int *) 0xdeedbeef;
  s2.ptr1 = (int *) 0L;
  s2.ptr2 = (int *) 0xdeedbeef;
  a = (int *) 0L;
  b = (int *) 0xdeedbeef;
  c = (int *) 0L;
  d = (int *) 0xdeedbeef;

  int n1, n2, n3, n4;
  n1 = n2 = n3 = n4 = 0;

  #pragma omp target enter data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])

  #pragma omp target map(s2.ptr1[:n1], s2.ptr2[:n2], c[:n3], d[:n4])
  {
    if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
	|| c != (void *) 0L || d != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
  {
    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target
  {
    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
      __builtin_abort ();
  }

  #pragma omp target exit data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])

  with_offset ();
}