Example #1
0
/* version 1: use omp parallel, i.e. each host thread responsible for one dev */
void axpy_mdev_v1(REAL* x, REAL* y, int n, REAL a) {
  int ndev = omp_get_num_devices(); /* standard omp call, see ticket 167 */
  #pragma omp parallel num_threads(ndev)
  {
        int i;
        /* chunking it for each device */
        int devid = omp_get_thread_num();
        int remain = n % ndev;
        int esize = n / ndev;
        int partsize, starti, endi;
        if (devid < remain) { /* each of the first remain dev has one more element */
                partsize = esize+1;
                starti = partsize*devid;
        } else {
                partsize = esize;
                starti = esize*devid+remain;
        }
        endi=starti + partsize;

#pragma omp target device (devid) map(inout: y[starti:endi]) map(in: x[starti:endi],a,partsize)
#pragma omp parallel for shared(x, y, partsize, a) private(i)
        for (i = 0; i < partsize; ++i)
          y[i] += a * x[i];
  }
}
Example #2
0
int main ()
{
  int i;
  int offload[N];
  int num = omp_get_num_devices();

  #pragma omp parallel for
    for (i = 0; i < N; i++)
      #pragma omp target device(i) map(from: offload[i:1])
	offload[i] = omp_is_initial_device ();

  for (i = 0; i < num; i++)
    if (offload[i])
      abort ();

  for (i = num; i < N; i++)
    if (!offload[i])
      abort ();

  return 0;
}
Example #3
0
int
main ()
{
  int d = omp_get_default_device ();
  int id = omp_get_initial_device ();

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

  int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
  int *b = a;
  int shared_mem = 0;
  #pragma omp target map (alloc: shared_mem)
  shared_mem = 1;
  if (omp_target_is_present (b, d) != shared_mem)
    abort ();
  #pragma omp target enter data map (to: a)
  if (omp_target_is_present (b, d) == 0)
    abort ();
  #pragma omp target enter data map (alloc: b[:0])
  if (omp_target_is_present (b, d) == 0)
    abort ();
  #pragma omp target exit data map (release: b[:0])
  if (omp_target_is_present (b, d) == 0)
    abort ();
  #pragma omp target exit data map (release: b[:0])
  if (omp_target_is_present (b, d) != shared_mem)
    abort ();
  #pragma omp target enter data map (to: a)
  if (omp_target_is_present (b, d) == 0)
    abort ();
  #pragma omp target enter data map (always, to: b[:0])
  if (omp_target_is_present (b, d) == 0)
    abort ();
  #pragma omp target exit data map (delete: b[:0])
  if (omp_target_is_present (b, d) != shared_mem)
    abort ();
  #pragma omp target exit data map (from: b[:0])
  return 0;
}
Example #4
0
int main()
{
  int i = 1;
  int j = 2;
  int k = 3;

  int num_devices = omp_get_num_devices();
  printf(" There are %d devices.\n ",num_devices);
  if(num_devices == 0) printf(" ... if there was a device...\n");
  #pragma omp target map(to:j) map(tofrom:k) 
  {
     printf(" Inside target ifs is %d.  It should be 42\n",  ifs);
     printf(" Inside target i   is %d.  It should be undefined \n",  i);
     printf(" Inside target j   is %d.  It should be 2\n",  j);
     printf(" Inside target k   is %d.  It should be 3\n",   k);
     ifs++; i++; j++; k++;
  }
  printf(" Outside target ifs is %d.  It should be 42\n",  ifs);
  printf(" Outside target i   is %d.  It should be 1\n",  i);
  printf(" Outside target j   is %d.  It should be 2\n",  j);
  printf(" Outside target k   is %d.  It should be 4\n",   k);

}
Example #5
0
int32_t
omp_get_num_devices_ (void)
{
  return omp_get_num_devices ();
}
Example #6
0
int
main ()
{
  int d = omp_get_default_device ();
  int id = omp_get_initial_device ();
  int err;
  int q[128], i;
  void *p;

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

  for (i = 0; i < 128; i++)
    q[i] = i;

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

  if (omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, NULL,
			      d, id) < 3
      || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
				 NULL, id, d) < 3
      || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
				 NULL, id, id) < 3)
    abort ();

  if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) == 0)
    {
      size_t volume[3] = { 128, 0, 0 };
      size_t dst_offsets[3] = { 0, 0, 0 };
      size_t src_offsets[3] = { 1, 0, 0 };
      size_t dst_dimensions[3] = { 128, 0, 0 };
      size_t src_dimensions[3] = { 128, 0, 0 };

      if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) != 0)
	abort ();

      if (omp_target_is_present (q, d) != 1
	  || omp_target_is_present (&q[32], d) != 1
	  || omp_target_is_present (&q[127], d) != 1)
	abort ();

      if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
			     d, id) != 0)
	abort ();

      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
      {
	int j;
	err = 0;
	for (j = 0; j < 128; j++)
	  if (q[j] != j)
	    err = 1;
	  else
	    q[j] += 4;
      }

      if (err)
	abort ();

      if (omp_target_memcpy_rect (q, p, sizeof (int), 1, volume,
				  dst_offsets, src_offsets, dst_dimensions,
				  src_dimensions, id, d) != 0)
	abort ();

      for (i = 0; i < 128; i++)
	if (q[i] != i + 4)
	  abort ();

      volume[2] = 2;
      volume[1] = 3;
      volume[0] = 6;
      dst_offsets[2] = 1;
      dst_offsets[1] = 0;
      dst_offsets[0] = 0;
      src_offsets[2] = 1;
      src_offsets[1] = 0;
      src_offsets[0] = 3;
      dst_dimensions[2] = 2;
      dst_dimensions[1] = 3;
      dst_dimensions[0] = 6;
      src_dimensions[2] = 3;
      src_dimensions[1] = 4;
      src_dimensions[0] = 6;
      if (omp_target_memcpy_rect (p, q, sizeof (int), 3, volume,
				  dst_offsets, src_offsets, dst_dimensions,
				  src_dimensions, d, id) != 0)
	abort ();

      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
      {
	int j, k, l;
	err = 0;
	for (j = 0; j < 6; j++)
	  for (k = 0; k < 3; k++)
	    for (l = 0; l < 2; l++)
	      if (q[j * 6 + k * 2 + l] != 3 * 12 + 4 + 1 + l + k * 3 + j * 12)
		err = 1;
      }

      if (err)
	abort ();

      if (omp_target_memcpy (p, p, 10 * sizeof (int), 51 * sizeof (int),
			     111 * sizeof (int), d, d) != 0)
	abort ();

      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
	{
	  int j;
	  err = 0;
	  for (j = 0; j < 10; j++)
	    if (q[50 + j] != q[110 + j])
	      err = 1;
	}

      if (err)
	abort ();

      if (omp_target_disassociate_ptr (q, d) != 0)
	abort ();
    }

  omp_target_free (p, d);
  return 0;
}
Example #7
0
__attribute__((noinline, noclone)) void
foo (S<C, I, L, UCR, CAR, SH, IPR> s)
{
  int d = omp_get_default_device ();
  int id = omp_get_initial_device ();
  int sep = 1;

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

  int err;
  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
  {
    err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
    err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
    err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
    err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
    s.a = 35; s.b[0] = 36; s.b[1] = 37;
    s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
    s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
    s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
    sep = 0;
  }
  if (err) abort ();
  err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
  err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
  err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
  err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
  if (err) abort ();
  s.a = 50; s.b[0] = 49; s.b[1] = 48;
  s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
  s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
  s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
  if (sep
      && (omp_target_is_present (&s.a, d)
	  || omp_target_is_present (s.b, d)
	  || omp_target_is_present (&s.c[1], d)
	  || omp_target_is_present (s.d, d)
	  || omp_target_is_present (&s.d[-2], d)
	  || omp_target_is_present (&s.e, d)
	  || omp_target_is_present (s.f, d)
	  || omp_target_is_present (&s.g[1], d)
	  || omp_target_is_present (&s.h, d)
	  || omp_target_is_present (&s.h[2], d)))
    abort ();
  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  {
    if (!omp_target_is_present (&s.a, d)
	|| !omp_target_is_present (s.b, d)
	|| !omp_target_is_present (&s.c[1], d)
	|| !omp_target_is_present (s.d, d)
	|| !omp_target_is_present (&s.d[-2], d)
	|| !omp_target_is_present (&s.e, d)
	|| !omp_target_is_present (s.f, d)
	|| !omp_target_is_present (&s.g[1], d)
	|| !omp_target_is_present (&s.h, d)
	|| !omp_target_is_present (&s.h[2], d))
      abort ();
    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
    {
      err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
      err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
      err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
      err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
      s.a = 17; s.b[0] = 18; s.b[1] = 19;
      s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
      s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
      s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
    }
    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  }
  if (sep
      && (omp_target_is_present (&s.a, d)
	  || omp_target_is_present (s.b, d)
	  || omp_target_is_present (&s.c[1], d)
	  || omp_target_is_present (s.d, d)
	  || omp_target_is_present (&s.d[-2], d)
	  || omp_target_is_present (&s.e, d)
	  || omp_target_is_present (s.f, d)
	  || omp_target_is_present (&s.g[1], d)
	  || omp_target_is_present (&s.h, d)
	  || omp_target_is_present (&s.h[2], d)))
    abort ();
  if (err) abort ();
  err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
  err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
  err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
  err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
  if (err) abort ();
  s.a = 33; s.b[0] = 34; s.b[1] = 35;
  s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
  s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
  s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  if (!omp_target_is_present (&s.a, d)
      || !omp_target_is_present (s.b, d)
      || !omp_target_is_present (&s.c[1], d)
      || !omp_target_is_present (s.d, d)
      || !omp_target_is_present (&s.d[-2], d)
      || !omp_target_is_present (&s.e, d)
      || !omp_target_is_present (s.f, d)
      || !omp_target_is_present (&s.g[1], d)
      || !omp_target_is_present (&s.h, d)
      || !omp_target_is_present (&s.h[2], d))
    abort ();
  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
  {
    err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
    err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
    err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
    err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
    s.a = 49; s.b[0] = 48; s.b[1] = 47;
    s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
    s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
    s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
  }
  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  if (!omp_target_is_present (&s.a, d)
      || !omp_target_is_present (s.b, d)
      || !omp_target_is_present (&s.c[1], d)
      || !omp_target_is_present (s.d, d)
      || !omp_target_is_present (&s.d[-2], d)
      || !omp_target_is_present (&s.e, d)
      || !omp_target_is_present (s.f, d)
      || !omp_target_is_present (&s.g[1], d)
      || !omp_target_is_present (&s.h, d)
      || !omp_target_is_present (&s.h[2], d))
    abort ();
  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
  if (sep
      && (omp_target_is_present (&s.a, d)
	  || omp_target_is_present (s.b, d)
	  || omp_target_is_present (&s.c[1], d)
	  || omp_target_is_present (s.d, d)
	  || omp_target_is_present (&s.d[-2], d)
	  || omp_target_is_present (&s.e, d)
	  || omp_target_is_present (s.f, d)
	  || omp_target_is_present (&s.g[1], d)
	  || omp_target_is_present (&s.h, d)
	  || omp_target_is_present (&s.h[2], d)))
    abort ();
  if (err) abort ();
  err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
  err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
  err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
  err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
  if (err) abort ();
}
Example #8
0
int main(void){
  #if CHECK
    check_offloading();
  #endif

  /*
   * Default device
   */
  printf("Is%s initial device\n", omp_is_initial_device() ? "" : " not");
  printf("Initial device: %d\n", omp_get_initial_device());
  omp_set_default_device(1);
  printf("Default device before task: %d\n", omp_get_default_device());
  #pragma omp task
  {
    printf("Default device inside task: %d\n", omp_get_default_device());
    omp_set_default_device(2);
    printf("Default device inside task after resetting: %d\n",
        omp_get_default_device());
  }
  #pragma omp taskwait
  printf("Default device outside task: %d\n", omp_get_default_device());

  // default device can set to whatever, if target fails, it goes to the host
  const int default_device = 0;
  omp_set_default_device(default_device);

  // default device for omp target call MUST be >= 0 and <omp_get_num_devices() or
  // the initial device. So when there are no devices, it must be the initial device
  int default_device_omp_target_call = default_device;
  if (omp_get_num_devices() == 0) {
    default_device_omp_target_call = omp_get_initial_device();
  } 
  #if DEBUG
    printf("test on machine with %d devices\n", omp_get_num_devices());
  #endif
  /*
   * Target alloc & target memcpy
   */
  double A[N], B[N], C[N], D[N], E[N];
  double *pA, *pB, *pC, *pD, *pE;
  // map ptrs
  pA = &A[0];
  pB = &B[0];
  pC = &C[0];
  pD = &D[0];
  pE = &E[0];

  INIT();

  pA = pA - 10;
  pC = pC - 20;
  pD = pD - 30;
  void *device_A = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
  void *device_C = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
  void *device_D = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
  double *dpA = (double *) device_A - 100;
  double *dpC = (double *) device_C - 200;
  double *dpD = (double *) device_D - 300;
  printf("omp_target_alloc %s\n", device_A && device_C && device_D ?
      "succeeded" : "failed");

  omp_target_memcpy(dpC, pC, N*sizeof(double), 200*sizeof(double),
      20*sizeof(double), default_device_omp_target_call, omp_get_initial_device());
  omp_target_memcpy(dpD, pD, N*sizeof(double), 300*sizeof(double),
      30*sizeof(double), default_device_omp_target_call, omp_get_initial_device());

  #pragma omp target is_device_ptr(dpA, dpC, dpD) device(default_device)
  {
    #pragma omp parallel for schedule(static,1)
    for (int i = 0; i < 992; i++)
      dpA[i+100] = dpC[i+200] + dpD[i+300] + 1;
  }

  omp_target_memcpy(pA, dpA, N*sizeof(double), 10*sizeof(double),
      100*sizeof(double), omp_get_initial_device(), default_device_omp_target_call);

  int fail = 0;
  VERIFY(0, N, A[i], (double)(i+2));
  if (fail) {
    printf ("Test omp_target_memcpy: Failed\n");
  } else {
    printf ("Test omp_target_memcpy: Succeeded\n");
  }

  /*
   * target_is_present and target_associate/disassociate_ptr
   */
  INIT();
  if (offloading_disabled()) {
    // If offloading is disabled just recreate the messages so that this can
    // also be tested with no device.
    printf("C is not present, associating it...\n");
    printf("omp_target_associate_ptr C %s\n", 1 ? "succeeded" : "failed");
  } else if (!omp_target_is_present(C, default_device_omp_target_call)) {
    printf("C is not present, associating it...\n");
    int rc = omp_target_associate_ptr(C, dpC, N*sizeof(double),
        200*sizeof(double), default_device_omp_target_call);
    printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");
  }
  if (offloading_disabled()) {
    // If offloading is disabled just recreate the messages so that this can
    // also be tested with no device.
    printf("D is not present, associating it...\n");
    printf("omp_target_associate_ptr D %s\n", 1 ? "succeeded" : "failed");
  } else if (!omp_target_is_present(D, default_device_omp_target_call)) {
    printf("D is not present, associating it...\n");
    int rc = omp_target_associate_ptr(D, dpD, N*sizeof(double),
        300*sizeof(double), default_device_omp_target_call);
    printf("omp_target_associate_ptr D %s\n", !rc ? "succeeded" : "failed");
  }
  #pragma omp target data map(from: C, D) device(default_device)
  {
    printf("Inside target data: A is%s present\n",
        (omp_target_is_present(A, default_device_omp_target_call) && !offloading_disabled()) ? "" : " not");
    printf("Inside target data: C is%s present\n",
        omp_target_is_present(C, default_device_omp_target_call) ? "" : " not");
    printf("Inside target data: D is%s present\n",
        omp_target_is_present(D, default_device_omp_target_call) ? "" : " not");

    // C and D are mapped "from", so there is no copy from host to device.
    // If the association was successful, their corresponding device arrays
    // are already populated from previous omp_target_memcpy with the correct
    // values and the following target for-loop must yield the correct results.
    #pragma omp target map(from: A) device(default_device)
    {
      #pragma omp parallel for schedule(static,1)
      for (int i = 0; i < 992; i++)
        A[i] = C[i] + D[i] + 1;
    }
  }


  if (offloading_disabled()) {
    printf("C is present, disassociating it...\n");
    printf("omp_target_disassociate_ptr C %s\n", 1 ? "succeeded" : "failed");
  } else if (omp_target_is_present(C, default_device_omp_target_call)) {
    printf("C is present, disassociating it...\n");
    int rc = omp_target_disassociate_ptr(C, default_device_omp_target_call);
    printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");
  }
  if (offloading_disabled()) {
    printf("D is present, disassociating it...\n");
    printf("omp_target_disassociate_ptr D %s\n", 1 ? "succeeded" : "failed");
  } else if (omp_target_is_present(D, default_device_omp_target_call)) {
    printf("D is present, disassociating it...\n");
    int rc = omp_target_disassociate_ptr(D, default_device_omp_target_call);
    printf("omp_target_disassociate_ptr D %s\n", !rc ? "succeeded" : "failed");
  }

  fail = 0;
  VERIFY(0, N, A[i], (double)(i+2));
  if (fail) {
    printf ("Test omp_target_associate_ptr: Failed\n");
  } else {
    printf ("Test omp_target_associate_ptr: Succeeded\n");
  }

  omp_target_free(device_A, default_device_omp_target_call);
  omp_target_free(device_C, default_device_omp_target_call);
  omp_target_free(device_D, default_device_omp_target_call);

  return 0;
}
Example #9
0
void
foo (int f)
{
  int d = f ? omp_get_num_devices () : omp_get_default_device ();
  int h = 5;
  #pragma omp target device (d)
  if (omp_get_level () != 0)
    abort ();
  #pragma omp target if (v > 1)
  if (omp_get_level () != 0 || !omp_is_initial_device ())
    abort ();
  #pragma omp target device (d) if (v > 1)
  if (omp_get_level () != 0 || !omp_is_initial_device ())
    abort ();
  #pragma omp target if (v <= 1)
  if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
    abort ();
  #pragma omp target device (d) if (v <= 1)
  if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
    abort ();
  #pragma omp target if (0)
  if (omp_get_level () != 0 || !omp_is_initial_device ())
    abort ();
  #pragma omp target device (d) if (0)
  if (omp_get_level () != 0 || !omp_is_initial_device ())
    abort ();
  #pragma omp target if (1)
  if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
    abort ();
  #pragma omp target device (d) if (1)
  if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
    abort ();
  #pragma omp target data device (d) map (to: h)
  {
    #pragma omp target device (d)
    if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
      abort ();
    #pragma omp target update device (d) from (h)
  }
  #pragma omp target data if (v > 1) map (to: h)
  {
    #pragma omp target if (v > 1)
    if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
      abort ();
    #pragma omp target update if (v > 1) from (h)
  }
  #pragma omp target data device (d) if (v > 1) map (to: h)
  {
    #pragma omp target device (d) if (v > 1)
    if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
      abort ();
    #pragma omp target update device (d) if (v > 1) from (h)
  }
  #pragma omp target data if (v <= 1) map (to: h)
  {
    #pragma omp target if (v <= 1)
    if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 8)
      abort ();
    #pragma omp target update if (v <= 1) from (h)
  }
  #pragma omp target data device (d) if (v <= 1) map (to: h)
  {
    #pragma omp target device (d) if (v <= 1)
    if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
      abort ();
    #pragma omp target update device (d) if (v <= 1) from (h)
  }
  #pragma omp target data if (0) map (to: h)
  {
    #pragma omp target if (0)
    if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
      abort ();
    #pragma omp target update if (0) from (h)
  }
  #pragma omp target data device (d) if (0) map (to: h)
  {
    #pragma omp target device (d) if (0)
    if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
      abort ();
    #pragma omp target update device (d) if (0) from (h)
  }
  #pragma omp target data if (1) map (to: h)
  {
    #pragma omp target if (1)
    if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 12)
      abort ();
    #pragma omp target update if (1) from (h)
  }
  #pragma omp target data device (d) if (1) map (to: h)
  {
    #pragma omp target device (d) if (1)
    if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
      abort ();
    #pragma omp target update device (d) if (1) from (h)
  }
  if (h != 14)
    abort ();
}