libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c

   1 /* PR libgomp/109837 */
   2
   3 #include <assert.h>
   4 #include <omp.h>
   5 #include <stdio.h>
   6 #include <stdlib.h>
   7
   8 #pragma omp requires unified_address
   9
  10 #define N 15
  11
  12 void
  13 test_device (int dev)
  14 {
  15   struct st {
  16     int *ptr;
  17     int n;
  18   };
  19   struct st s;
  20
  21   s.n = 10;
  22   s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev);
  23   int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev);
  24   assert (s.ptr != NULL);
  25   assert (ptr1 != NULL);
  26
  27   int q[4] = {1,2,3,4};
  28   int *qptr;
  29   #pragma omp target enter data map(q) device(device_num: dev)
  30   #pragma omp target data use_device_addr(q) device(device_num: dev)
  31     qptr = q;
  32
  33   #pragma omp target map(to:s) device(device_num: dev)
  34   for (int i = 0; i < s.n; i++)
  35     s.ptr[i] = 23*i;
  36
  37   int *ptr2 = &s.ptr[3];
  38
  39   #pragma omp target firstprivate(qptr) map(tofrom:ptr2) device(device_num: dev)
  40   for (int i = 0; i < 4; i++)
  41     *(qptr++) = ptr2[i];
  42
  43   #pragma omp target exit data map(q) device(device_num: dev)
  44   for (int i = 0; i < 4; i++)
  45     q[i] = 23 * (i+3);
  46
  47   #pragma omp target map(to: ptr1) device(device_num: dev)
  48   for (int i = 0; i < N; i++)
  49     ptr1[i] = 11*i;
  50
  51   int *ptr3 = (int *) malloc (sizeof (int)*N);
  52   assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0,
  53                                  omp_get_initial_device(), dev));
  54   for (int i = 0; i < N; i++)
  55     assert (ptr3[i] == 11*i);
  56
  57   free (ptr3);
  58   omp_target_free (ptr1, dev);
  59   omp_target_free (s.ptr, dev);
  60 }
  61
  62 int
  63 main()
  64 {
  65   int ntgts = omp_get_num_devices();
  66   if (ntgts)
  67     fprintf (stderr, "Offloading devices exist\n");  /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */
  68   else
  69     fprintf (stderr, "Only host fallback\n");        /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */
  70
  71   for (int i = 0; i <= ntgts; i++)
  72     test_device (i);
  73   return 0;
  74 }