libgomp/testsuite/libgomp.c/target-12.c

   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 }