libgomp/testsuite/libgomp.c-c++-common/icv-6.c

   1 /* { dg-do run } */
   2 /* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } */
   3 /* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } */
   4 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } */
   5 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } */
   6 /* { dg-set-target-env-var OMP_SCHEDULE_ALL "guided,4" } */
   7 /* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */
   8 /* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "45" } */
   9 /* { dg-set-target-env-var OMP_NUM_THREADS_ALL "46,3,2" } */
  10 /* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "47" } */
  11 /* { dg-set-target-env-var OMP_PROC_BIND_ALL "spread" } */
  12 /* { dg-set-target-env-var OMP_WAIT_POLICY_ALL "active" } */
  13
  14 /* This tests the hierarchical usage of ICVs on the device, i.e. if
  15    OMP_NUM_TEAMS_DEV_<device_num> is not configured, then the value of
  16    OMP_NUM_TEAMS_DEV should be used.  And if OMP_NUM_TEAMS (without suffix) is
  17    not defined, then OMP_NUM_TEAMS_ALL should be used for the host.  */
  18
  19 #include <omp.h>
  20 #include <stdlib.h>
  21 #include <string.h>
  22
  23 int
  24 main ()
  25 {
  26   enum omp_sched_t kind;
  27   int chunk_size;
  28   omp_get_schedule(&kind, &chunk_size);
  29
  30   if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 3)
  31       || (!getenv ("OMP_DYNAMIC") && !omp_get_dynamic ())
  32       || (!getenv ("OMP_SCHEDULE") && (kind != 3 || chunk_size != 4))
  33       || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 2)
  34       || (!getenv ("OMP_THREAD_LIMIT") && omp_get_thread_limit () != 45)
  35       || (!getenv ("OMP_NUM_THREADS") && omp_get_max_threads () != 46)
  36       || (!getenv ("OMP_PROC_BIND") && omp_get_proc_bind () != omp_proc_bind_spread)
  37       || (!getenv ("OMP_MAX_ACTIVE_LEVELS") && omp_get_max_active_levels () != 47))
  38     abort ();
  39
  40   int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
  41   for (int i = 0; i < num_devices; i++)
  42     {
  43       char name[sizeof ("OMP_NUM_TEAMS_DEV_1")];
  44       strcpy (name, "OMP_NUM_TEAMS_DEV_1");
  45       name[sizeof ("OMP_NUM_TEAMS_DEV_1") - 2] = '0' + i;
  46       if (getenv (name))
  47         continue;
  48
  49       #pragma omp target device (i)
  50       if (omp_get_max_teams () != 4
  51           || omp_get_teams_thread_limit () != 3)
  52         abort ();
  53       #pragma omp target device (i)
  54       #pragma omp teams
  55       {
  56         if (omp_get_num_teams () > 4
  57             || omp_get_team_num () >= 4)
  58           abort ();
  59         #pragma omp parallel
  60         if (omp_get_thread_limit () > 3
  61             || omp_get_thread_num () >= 3)
  62           abort ();
  63       }
  64
  65       #pragma omp target device (i)
  66       {
  67         omp_set_num_teams (3 + i);
  68         omp_set_teams_thread_limit (2 + i);
  69         if (omp_get_max_teams () != 3 + i
  70             || omp_get_teams_thread_limit () != 2 + i)
  71           abort ();
  72       }
  73
  74      /* omp_set_num_teams above set the value of nteams-var ICV on device 'i',
  75          which has scope 'device' and should be avaible in subsequent target
  76          regions.  */
  77       #pragma omp target device (i)
  78       if (omp_get_max_teams () != 3 + i
  79           || omp_get_teams_thread_limit () != 2 + i)
  80         abort ();
  81
  82       #pragma omp target device (i)
  83       #pragma omp teams
  84       {
  85         if (omp_get_num_teams () > 3 + i
  86             || omp_get_team_num () >= 3 + i)
  87           abort ();
  88         #pragma omp parallel
  89         if (omp_get_thread_limit () > 2 + i
  90             || omp_get_thread_num () >= 2 + i)
  91           abort ();
  92       }
  93     }
  94
  95   return 0;
  96 }