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  }