libgomp/testsuite/libgomp.c/target-teams-1.c

   1 /* { dg-do run } */
   2
   3 #include <omp.h>
   4 #include <stdlib.h>
   5
   6 int v = 6;
   7
   8 void
   9 bar (long *x, long *y)
  10 {
  11   *x += 2;
  12   *y += 3;
  13 }
  14
  15 int
  16 baz (void)
  17 {
  18   return 5;
  19 }
  20
  21 #pragma omp declare target to (bar, baz, v)
  22
  23 __attribute__((noinline, noclone)) void
  24 foo (int a, int b, long c, long d)
  25 {
  26   int err;
  27   if (omp_get_num_teams () != 1)
  28     abort ();
  29   /* The OpenMP 4.5 spec says that these expressions are evaluated before
  30      target region on combined target teams, so those cases are always
  31      fine.  */
  32   #pragma omp target map(from: err)
  33   err = omp_get_num_teams () != 1;
  34   if (err)
  35     abort ();
  36   #pragma omp target map(from: err)
  37   #pragma omp teams
  38   {
  39     err = omp_get_num_teams () < 1;
  40     #pragma omp parallel if(0)
  41     err |= omp_get_thread_limit () < 1;
  42   }
  43   if (err)
  44     abort ();
  45   #pragma omp target teams map(from: err)
  46   {
  47     err = omp_get_num_teams () < 1;
  48     #pragma omp parallel if(0)
  49     err |= omp_get_thread_limit () < 1;
  50   }
  51   if (err)
  52     abort ();
  53   #pragma omp target map(from: err)
  54   #pragma omp teams num_teams (4)
  55   {
  56     err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
  57     #pragma omp parallel if(0)
  58     err |= omp_get_thread_limit () < 1;
  59   }
  60   if (err)
  61     abort ();
  62   #pragma omp target teams num_teams (4) map(from: err)
  63   {
  64     err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
  65     #pragma omp parallel if(0)
  66     err |= omp_get_thread_limit () < 1;
  67   }
  68   if (err)
  69     abort ();
  70   #pragma omp target map(from: err)
  71   #pragma omp teams thread_limit (7)
  72   {
  73     err = omp_get_num_teams () < 1;
  74     #pragma omp parallel if(0)
  75     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7;
  76   }
  77   if (err)
  78     abort ();
  79   #pragma omp target teams thread_limit (7) map(from: err)
  80   {
  81     err = omp_get_num_teams () < 1;
  82     #pragma omp parallel if(0)
  83     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7;
  84   }
  85   if (err)
  86     abort ();
  87   #pragma omp target map(from: err)
  88   #pragma omp teams num_teams (4) thread_limit (8)
  89   {
  90     {
  91       err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
  92     }
  93     #pragma omp parallel if(0)
  94     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 8;
  95   }
  96   if (err)
  97     abort ();
  98   #pragma omp target teams num_teams (4) thread_limit (8) map(from: err)
  99   {
 100     err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
 101     #pragma omp parallel if(0)
 102     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 8;
 103   }
 104   if (err)
 105     abort ();
 106   #pragma omp target map(from: err)
 107   #pragma omp teams num_teams (a) thread_limit (b)
 108   {
 109     err = omp_get_num_teams () < 1 || omp_get_num_teams () > a;
 110     #pragma omp parallel if(0)
 111     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > b;
 112   }
 113   if (err)
 114     abort ();
 115   #pragma omp target teams num_teams (a) thread_limit (b) map(from: err)
 116   {
 117     err = omp_get_num_teams () < 1 || omp_get_num_teams () > a;
 118     #pragma omp parallel if(0)
 119     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > b;
 120   }
 121   if (err)
 122     abort ();
 123   #pragma omp target map(from: err)
 124   #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
 125   {
 126     err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
 127     #pragma omp parallel if(0)
 128     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
 129   }
 130   if (err)
 131     abort ();
 132   #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err)
 133   {
 134     err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
 135     #pragma omp parallel if(0)
 136     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
 137   }
 138   if (err)
 139     abort ();
 140   #pragma omp target map (always, to: c, d) map(from: err)
 141   #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
 142   {
 143     err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
 144     #pragma omp parallel if(0)
 145     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
 146   }
 147   if (err)
 148     abort ();
 149   #pragma omp target data map (to: c, d)
 150   {
 151     #pragma omp target defaultmap (tofrom: scalar)
 152     bar (&c, &d);
 153     /* This is one of the cases which can't be generally optimized,
 154        the c and d are (or could be) already mapped and whether
 155        their device and original values match is unclear.  */
 156     #pragma omp target map (to: c, d) map(from: err)
 157     #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
 158     {
 159       err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
 160       #pragma omp parallel if(0)
 161       err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
 162     }
 163     if (err)
 164       abort ();
 165   }
 166   /* This can't be optimized, there are function calls inside of
 167      target involved.  */
 168   #pragma omp target map(from: err)
 169   #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
 170   {
 171     err = omp_get_num_teams () < 1 || omp_get_num_teams () > baz () + 1;
 172     #pragma omp parallel if(0)
 173     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > baz () - 1;
 174   }
 175   if (err)
 176     abort ();
 177   #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err)
 178   {
 179     err = omp_get_num_teams () < 1 || omp_get_num_teams () > baz () + 1;
 180     #pragma omp parallel if(0)
 181     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > baz () - 1;
 182   }
 183   if (err)
 184     abort ();
 185   /* This one can't be optimized, as v might have different value between
 186      host and target.  */
 187   #pragma omp target map(from: err)
 188   #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
 189   {
 190     err = omp_get_num_teams () < 1 || omp_get_num_teams () > v + 1;
 191     #pragma omp parallel if(0)
 192     err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > v - 1;
 193   }
 194   if (err)
 195     abort ();
 196 }
 197
 198 int
 199 main ()
 200 {
 201   foo (3, 5, 7, 9);
 202   return 0;
 203 }