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   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
  39   if (err)
  40     abort ();
  41   #pragma omp target teams map(from: err)
  42   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
  43   if (err)
  44     abort ();
  45   #pragma omp target map(from: err)
  46   #pragma omp teams num_teams (4)
  47   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  48         || omp_get_num_teams () > 4;
  49   if (err)
  50     abort ();
  51   #pragma omp target teams num_teams (4) map(from: err)
  52   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  53         || omp_get_num_teams () > 4;
  54   if (err)
  55     abort ();
  56   #pragma omp target map(from: err)
  57   #pragma omp teams thread_limit (7)
  58   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  59         || omp_get_thread_limit () > 7;
  60   if (err)
  61     abort ();
  62   #pragma omp target teams thread_limit (7) map(from: err)
  63   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  64         || omp_get_thread_limit () > 7;
  65   if (err)
  66     abort ();
  67   #pragma omp target map(from: err)
  68   #pragma omp teams num_teams (4) thread_limit (8)
  69   {
  70     {
  71       err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  72             || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
  73     }
  74   }
  75   if (err)
  76     abort ();
  77   #pragma omp target teams num_teams (4) thread_limit (8) map(from: err)
  78   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  79         || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
  80   if (err)
  81     abort ();
  82   #pragma omp target map(from: err)
  83   #pragma omp teams num_teams (a) thread_limit (b)
  84   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  85         || omp_get_num_teams () > a || omp_get_thread_limit () > b;
  86   if (err)
  87     abort ();
  88   #pragma omp target teams num_teams (a) thread_limit (b) map(from: err)
  89   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  90         || omp_get_num_teams () > a || omp_get_thread_limit () > b;
  91   if (err)
  92     abort ();
  93   #pragma omp target map(from: err)
  94   #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
  95   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
  96         || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
  97   if (err)
  98     abort ();
  99   #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err)
 100   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 101         || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
 102   if (err)
 103     abort ();
 104   #pragma omp target map (always, to: c, d) map(from: err)
 105   #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
 106   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 107         || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
 108   if (err)
 109     abort ();
 110   #pragma omp target data map (to: c, d)
 111   {
 112     #pragma omp target defaultmap (tofrom: scalar)
 113     bar (&c, &d);
 114     /* This is one of the cases which can't be generally optimized,
 115        the c and d are (or could be) already mapped and whether
 116        their device and original values match is unclear.  */
 117     #pragma omp target map (to: c, d) map(from: err)
 118     #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
 119     err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 120           || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
 121     if (err)
 122       abort ();
 123   }
 124   /* This can't be optimized, there are function calls inside of
 125      target involved.  */
 126   #pragma omp target map(from: err)
 127   #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
 128   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 129         || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
 130   if (err)
 131     abort ();
 132   #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err)
 133   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 134         || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
 135   if (err)
 136     abort ();
 137   /* This one can't be optimized, as v might have different value between
 138      host and target.  */
 139   #pragma omp target map(from: err)
 140   #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
 141   err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
 142         || omp_get_num_teams () > v + 1 || omp_get_thread_limit () > v - 1;
 143   if (err)
 144     abort ();
 145 }
 146
 147 int
 148 main ()
 149 {
 150   foo (3, 5, 7, 9);
 151   return 0;
 152 }