libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c

   1 /* PR middle-end/110270 */
   2
   3 /* Same as target-implicit-map-3.c but uses the following requiement
   4    and for not mapping the stack variables 'A' and 'B' (not mapped
   5    but accessible -> USM makes this tested feature even more important.)  */
   6
   7 #pragma omp requires unified_shared_memory
   8
   9 /* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2
  10    semantic, i.e. keeping the pointer value even if not mapped;
  11    before OpenMP 5.0/5.1 required that it is NULL. */
  12
  13 #include <stdio.h>
  14 #include <stdlib.h>
  15 #include <stdint.h>
  16 #include <omp.h>
  17
  18 void
  19 test_device (int dev)
  20 {
  21   int *p1 = (int*) 0x12345;
  22   int *p1a = (int*) 0x67890;
  23   int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev);
  24   int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev);
  25   intptr_t ip = (intptr_t) p2;
  26   intptr_t ipa = (intptr_t) p2a;
  27
  28   int A[3] = {1,2,3};
  29   int B[5] = {4,5,6,7,8};
  30   int *p3 = &A[0];
  31   int *p3a = &B[0];
  32
  33   const omp_alloctrait_t traits[]
  34     = { { omp_atk_alignment, 128 },
  35         { omp_atk_pool_size, 1024 }};
  36   omp_allocator_handle_t a = omp_init_allocator (omp_default_mem_space, 2, traits);
  37
  38   int *p4 = (int*) malloc (sizeof (int) * 5);
  39   int *p4a = (int*) omp_alloc (sizeof (int) * 10, a);
  40   intptr_t ip4 = (intptr_t) p4;
  41   intptr_t ip4a = (intptr_t) p4a;
  42
  43   for (int i = 0; i < 5; i++)
  44     p4[i] = -31*i;
  45
  46   for (int i = 0; i < 10; i++)
  47     p4a[i] = -43*i;
  48
  49   /* Note: 'A' is not mapped but USM accessible.  */
  50   #pragma omp target device(dev) /* defaultmap(default:pointer) */
  51   {
  52     /* The pointees aren't mapped. */
  53     /* OpenMP 5.2 -> same value as before the target region. */
  54     if ((intptr_t) p1 != 0x12345) abort ();
  55     if ((intptr_t) p2 != ip) abort ();
  56     for (int i = 0; i < 5; i++)
  57       p2[i] = 13*i;
  58
  59     for (int i = 0; i < 10; i++)
  60       ((int *)ipa)[i] = 7*i;
  61
  62     /* OpenMP: Points to 'A'. */
  63     if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3)
  64       abort ();
  65     p3[0] = -11; p3[1] = -22; p3[2] = -33;
  66
  67     /* USM accesible allocated host memory. */
  68     if ((intptr_t) p4 != ip4)
  69       abort ();
  70     for (int i = 0; i < 5; i++)
  71       if (p4[i] != -31*i)
  72         abort ();
  73     for (int i = 0; i < 10; i++)
  74       if (((int *)ip4a)[i] != -43*i)
  75         abort ();
  76     for (int i = 0; i < 5; i++)
  77       p4[i] = 9*i;
  78     for (int i = 0; i < 10; i++)
  79       ((int *)ip4a)[i] = 18*i;
  80   }
  81
  82   if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33)
  83     abort ();
  84
  85   for (int i = 0; i < 5; i++)
  86     if (p4[i] != 9*i)
  87       abort ();
  88   for (int i = 0; i < 10; i++)
  89     if (p4a[i] != 18*i)
  90       abort ();
  91   for (int i = 0; i < 5; i++)
  92     p4[i] = -77*i;
  93   for (int i = 0; i < 10; i++)
  94     p4a[i] = -65*i;
  95
  96   // With defaultmap:
  97
  98   /* Note: 'B' is not mapped but USM accessible.  */
  99   #pragma omp target device(dev) defaultmap(default:pointer)
 100   {
 101     /* The pointees aren't mapped. */
 102     /* OpenMP 5.2 -> same value as before the target region. */
 103     if ((intptr_t) p1a != 0x67890) abort ();
 104     if ((intptr_t) p2a != ipa) abort ();
 105
 106     for (int i = 0; i < 5; i++)
 107       ((int *)ip)[i] = 13*i;
 108
 109     for (int i = 0; i < 10; i++)
 110       p2a[i] = 7*i;
 111
 112     /* USM accesible allocated host memory. */
 113     if ((intptr_t) p4a != ip4a) abort ();
 114
 115     /* OpenMP: Points to 'B'. */
 116     if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8)
 117       abort ();
 118     p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88;
 119
 120     /* USM accesible allocated host memory. */
 121     if ((intptr_t) p4a != ip4a)
 122       abort ();
 123     for (int i = 0; i < 5; i++)
 124       if (((int *)ip4)[i] != -77*i)
 125         abort ();
 126     for (int i = 0; i < 10; i++)
 127       if (p4a[i] != -65*i)
 128         abort ();
 129     for (int i = 0; i < 5; i++)
 130       p4[i] = 36*i;
 131     for (int i = 0; i < 10; i++)
 132       ((int *)ip4a)[i] = 4*i;
 133   }
 134
 135   if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88)
 136     abort ();
 137
 138   for (int i = 0; i < 5; i++)
 139     if (p4[i] != 36*i)
 140       abort ();
 141   for (int i = 0; i < 10; i++)
 142     if (p4a[i] != 4*i)
 143       abort ();
 144
 145   omp_target_free (p2, dev);
 146   omp_target_free (p2a, dev);
 147   free (p4);
 148   omp_free (p4a, a);
 149   omp_destroy_allocator (a);
 150 }
 151
 152 int
 153 main()
 154 {
 155   int ntgts = omp_get_num_devices();
 156   for (int i = 0; i <= ntgts; i++)
 157     test_device (i);
 158   return 0;
 159 }