1  #include <omp.h>
       2  #include <stdlib.h>
       3  
       4  int
       5  main ()
       6  {
       7    int d = omp_get_default_device ();
       8    int id = omp_get_initial_device ();
       9    int err;
      10    int q[128], i;
      11    void *p;
      12  
      13    if (d < 0 || d >= omp_get_num_devices ())
      14      d = id;
      15  
      16    for (i = 0; i < 128; i++)
      17      q[i] = i;
      18  
      19    p = omp_target_alloc (130 * sizeof (int), d);
      20    if (p == NULL)
      21      return 0;
      22  
      23    if (omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, NULL,
      24  			      d, id) < 3
      25        || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
      26  				 NULL, id, d) < 3
      27        || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
      28  				 NULL, id, id) < 3)
      29      abort ();
      30  
      31    if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) == 0)
      32      {
      33        size_t volume[3] = { 128, 0, 0 };
      34        size_t dst_offsets[3] = { 0, 0, 0 };
      35        size_t src_offsets[3] = { 1, 0, 0 };
      36        size_t dst_dimensions[3] = { 128, 0, 0 };
      37        size_t src_dimensions[3] = { 128, 0, 0 };
      38  
      39        if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) != 0)
      40  	abort ();
      41  
      42        if (omp_target_is_present (q, d) != 1
      43  	  || omp_target_is_present (&q[32], d) != 1
      44  	  || omp_target_is_present (&q[127], d) != 1)
      45  	abort ();
      46  
      47        if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
      48  			     d, id) != 0)
      49  	abort ();
      50  
      51        #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
      52        {
      53  	int j;
      54  	err = 0;
      55  	for (j = 0; j < 128; j++)
      56  	  if (q[j] != j)
      57  	    err = 1;
      58  	  else
      59  	    q[j] += 4;
      60        }
      61  
      62        if (err)
      63  	abort ();
      64  
      65        if (omp_target_memcpy_rect (q, p, sizeof (int), 1, volume,
      66  				  dst_offsets, src_offsets, dst_dimensions,
      67  				  src_dimensions, id, d) != 0)
      68  	abort ();
      69  
      70        for (i = 0; i < 128; i++)
      71  	if (q[i] != i + 4)
      72  	  abort ();
      73  
      74        volume[2] = 2;
      75        volume[1] = 3;
      76        volume[0] = 6;
      77        dst_offsets[2] = 1;
      78        dst_offsets[1] = 0;
      79        dst_offsets[0] = 0;
      80        src_offsets[2] = 1;
      81        src_offsets[1] = 0;
      82        src_offsets[0] = 3;
      83        dst_dimensions[2] = 2;
      84        dst_dimensions[1] = 3;
      85        dst_dimensions[0] = 6;
      86        src_dimensions[2] = 3;
      87        src_dimensions[1] = 4;
      88        src_dimensions[0] = 6;
      89        if (omp_target_memcpy_rect (p, q, sizeof (int), 3, volume,
      90  				  dst_offsets, src_offsets, dst_dimensions,
      91  				  src_dimensions, d, id) != 0)
      92  	abort ();
      93  
      94        #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
      95        {
      96  	int j, k, l;
      97  	err = 0;
      98  	for (j = 0; j < 6; j++)
      99  	  for (k = 0; k < 3; k++)
     100  	    for (l = 0; l < 2; l++)
     101  	      if (q[j * 6 + k * 2 + l] != 3 * 12 + 4 + 1 + l + k * 3 + j * 12)
     102  		err = 1;
     103        }
     104  
     105        if (err)
     106  	abort ();
     107  
     108        if (omp_target_memcpy (p, p, 10 * sizeof (int), 51 * sizeof (int),
     109  			     111 * sizeof (int), d, d) != 0)
     110  	abort ();
     111  
     112        #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
     113  	{
     114  	  int j;
     115  	  err = 0;
     116  	  for (j = 0; j < 10; j++)
     117  	    if (q[50 + j] != q[110 + j])
     118  	      err = 1;
     119  	}
     120  
     121        if (err)
     122  	abort ();
     123  
     124        if (omp_target_disassociate_ptr (q, d) != 0)
     125  	abort ();
     126      }
     127  
     128    omp_target_free (p, d);
     129    return 0;
     130  }