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

   1 /* PR middle-end/110270 */
   2
   3 /* OpenMP 5.2's 'defaultmap(default : pointer) for C/C++ pointers retains the
   4    pointer value instead of setting it to NULL if the pointer cannot be found.
   5    Contrary to requires-unified-addr-1.c which is valid OpenMP 5.0/5.1/5.2,
   6    this testcase is only valid since OpenMP 5.2.  */
   7
   8 /* This is kind of a follow-up to the requires-unified-addr-1.c testcase
   9    and  PR libgomp/109837 */
  10
  11
  12 #include <assert.h>
  13 #include <omp.h>
  14 #include <stdio.h>
  15 #include <stdlib.h>
  16
  17 #pragma omp requires unified_address
  18
  19 #define N 15
  20
  21 void
  22 test_device (int dev)
  23 {
  24   struct st {
  25     int *ptr;
  26     int n;
  27   };
  28   struct st s;
  29
  30   s.n = 10;
  31   s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev);
  32   int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev);
  33   assert (s.ptr != NULL);
  34   assert (ptr1 != NULL);
  35
  36   int q[4] = {1,2,3,4};
  37   int *qptr;
  38   #pragma omp target enter data map(q) device(device_num: dev)
  39   #pragma omp target data use_device_addr(q) device(device_num: dev)
  40     qptr = q;
  41
  42   #pragma omp target map(to:s) device(device_num: dev)
  43   for (int i = 0; i < s.n; i++)
  44     s.ptr[i] = 23*i;
  45
  46   int *ptr2 = &s.ptr[3];
  47
  48   /* s.ptr is not mapped (but omp_target_alloc'ed) thus ptr2 shall retain its value.  */
  49   #pragma omp target device(device_num: dev)  /* implied: defaultmap(default : pointer) */
  50   for (int i = 0; i < 4; i++)
  51     *(qptr++) = ptr2[i];
  52
  53   #pragma omp target exit data map(q) device(device_num: dev)
  54   for (int i = 0; i < 4; i++)
  55     q[i] = 23 * (i+3);
  56
  57   /* ptr1 retains the value as it is not mapped (but it is omp_target_alloc'ed). */
  58   #pragma omp target defaultmap(default : pointer) device(device_num: dev)
  59   for (int i = 0; i < N; i++)
  60     ptr1[i] = 11*i;
  61
  62   int *ptr3 = (int *) malloc (sizeof (int)*N);
  63   assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0,
  64                                  omp_get_initial_device(), dev));
  65   for (int i = 0; i < N; i++)
  66     assert (ptr3[i] == 11*i);
  67
  68   free (ptr3);
  69   omp_target_free (ptr1, dev);
  70   omp_target_free (s.ptr, dev);
  71 }
  72
  73 int
  74 main()
  75 {
  76   int ntgts = omp_get_num_devices();
  77   if (ntgts)
  78     fprintf (stderr, "Offloading devices exist\n");  /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */
  79   else
  80     fprintf (stderr, "Only host fallback\n");        /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */
  81
  82   for (int i = 0; i <= ntgts; i++)
  83     test_device (i);
  84   return 0;
  85 }