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