[OpenMP] Add bounds to num_teams clause (OpenMP 5.1)

[Nawrin]

Summary

This patch adds lower-bound and upper-bound to num_teams clause according to OpenMP 5.1 specification. The number of initial teams created is implementation defined, but it will be greater than or equal to lower bound and less than or equal to the upper bound. If num_teams clause is not present, the number of initial teams created is implementation defined, but it will be greater than or equal to 1.

[jdoerfert]

openmp/runtime/src/kmp.h

3367

Given the version name, should we go with _51?

openmp/runtime/src/kmp_runtime.cpp

7545

IIRC, this is not correct anymore in 5.1. We cannot cap the number of teams. We need to implement a loop if the number of physical teams is smaller than the number requested teams to basically emulate as many as have been requested.

[Nawrin]

openmp/runtime/src/kmp_runtime.cpp

7545

According to the spec, when num_teams clause is not present, the number of teams created should be implementation defined, but it will be greater or equal to 1.

5.1 spec - "If the num_teams clause is not specified and the value of the nteams-var ICV is greater than zero, the number of teams created is less or equal to the value of the nteams-var ICV. Otherwise, the number of teams created is implementation defined, but it will be greater than or equal to 1."

[Nawrin]

openmp/runtime/src/kmp.h

3367

Sounds good. I'll change it to _51.

[jdoerfert]

openmp/runtime/src/kmp_runtime.cpp

7545

I'm not arguing against that. What I'm trying to say is that we cannot cap the number of teams. So we cannot do nteams = min(nteams, __kmp_teams_max_nth).

OpenMP 5.1 says, you get a number between the lower and upper bound. So if the lower bound is not smaller or equal than __kmp_teams_max_nth, we cannot cap at __kmp_teams_max_nth. This is explicitly different from num_threads on a parallel. For example, to implement N teams with a single thread, e.g., if __kmp_teams_max_nth was 1, you would do something like this:

for (unsigned team_no = 0; team_no < num_teams_requested; ++team_no) {
  set_team_no_for_thread(team_no);
  execute_team_body();
}

[AndreyChurbanov]

openmp/runtime/src/kmp_runtime.cpp

7545

I'm not arguing against that. What I'm trying to say is that we cannot cap the number of teams. So we cannot do nteams = min(nteams, __kmp_teams_max_nth).

Why? Once the specification allows to choose any number of teams, why would we complicate things without the need?

OpenMP 5.1 says, you get a number between the lower and upper bound. So if the lower bound is not smaller or equal than __kmp_teams_max_nth, we cannot cap at __kmp_teams_max_nth.

Right. But this is not the case in this code branch. Here we have neither lower nor upper bound.

We don't cap the number of teams when we have lower and/or upper bound. But again, it is not the case here.

[jdoerfert]

openmp/runtime/src/kmp_runtime.cpp

7545

Right. But this is not the case in this code branch. Here we have neither lower nor upper bound.

You are both right, in this branch there is no problem. I was under the impression we cap it in this conditional and I choose a bad location for my comment. If there is no capping anywhere, there is no problem. Apologies for the noise.

[Nawrin]

Update function name from _5 to _51

[AndreyChurbanov]

LGTM

[hoy]

openmp/runtime/test/teams/kmp_num_teams.c

69

Hello, I'm wondering if nteams could have an uninitialized use here. It is defined conditionally at line 41. I was seeing nteams had a random value. This happened in a very specific testing environment and sorry I cannot provide you a repro. Thanks.

[jdoerfert]

openmp/runtime/test/teams/kmp_num_teams.c

69

When the encountering thread finished executing the target construct it should have also executed line 41. That said, there might be a capturing issue here. Do we need to capture nteams as shared explicitly to avoid a firstprivate copy in the target teams region?

[AndreyChurbanov]

openmp/runtime/test/teams/kmp_num_teams.c

69

Once target construct does not have shared clause, map(tofrom:..) probably needed for a, nteams and nthreads.

[hoy]

openmp/runtime/test/teams/kmp_num_teams.c

69

OK, I'm seeing nteams has different addresses in the master and team threads. Maybe shared is needed.

[AndreyChurbanov]

openmp/runtime/test/teams/kmp_num_teams.c

69

I tried to fix the test in commit 091e8da.
Strictly the a needs to be in in_reduction and reduction, but I hope atomic to work here. Otherwise map(tofrom) should be changed to in_reduction(+:a) reduction(+:a).

[AndreyChurbanov]

openmp/runtime/test/teams/kmp_num_teams.c

69

Sorry, I was in a hurry. My previous fix should not work with real offload, because the test was intended to check the teams construct on a host. With real offload the test cannot pass because num_teams and thread_limit won't be propagated to the device. So the proper fix would be to remove the target construct, and execute the teams on host. I've committed another patch 5631842d1810 that should definitely fix the problem. I apologize to miss the test problem during review (I actually don't have a machine with an accelerator at hand).

[hoy]

openmp/runtime/test/teams/kmp_num_teams.c

69

Thanks for the quick fix!

openmp/runtime/src/dllexports

%ifndef stub
         __kmpc_dist_dispatch_init_4         251
         __kmpc_dist_dispatch_init_4u        252
         __kmpc_dist_dispatch_init_8         253
         __kmpc_dist_dispatch_init_8u        254
         __kmpc_team_static_init_4           255
         __kmpc_team_static_init_4u          256
         __kmpc_team_static_init_8           257
         __kmpc_team_static_init_8u          258
+        __kmpc_push_num_teams_51            284
 %endif
 
 %ifndef stub
         __kmpc_proxy_task_completed         259
         __kmpc_proxy_task_completed_ooo     260
         __kmpc_doacross_init                261
         __kmpc_doacross_wait                262
         __kmpc_doacross_post                263

openmp/runtime/src/i18n/en_US.txt

 HierSchedInvalid             "Hierarchy ignored: unsupported level: %1$s."
 AffFormatDefault             "OMP: pid %1$s tid %2$s thread %3$s bound to OS proc set {%4$s}"
 APIDeprecated                "%1$s routine deprecated, please use %2$s instead."
 GompFeatureNotSupported      "libgomp compatibility layer does not support OpenMP feature: %1$s"
 AffHWSubsetManyDies          "KMP_HW_SUBSET ignored: too many Dies requested."
 AffUseGlobCpuidL             "%1$s: Affinity capable, using global cpuid leaf %2$d info"
 AffNotCapableUseLocCpuidL    "%1$s: Affinity not capable, using local cpuid leaf %2$d info"
 AffNotUsingHwloc             "%1$s: Affinity not capable, using hwloc."
+FailedToCreateTeam           "Failed to create teams between lower bound (%1$d) and upper bound (%2$d)."
 
 # --------------------------------------------------------------------------------------------------
 -*- HINTS -*-
 # --------------------------------------------------------------------------------------------------
 
 # Hints. Hint may be printed after a message. Usually it is longer explanation text or suggestion.
 # To maintain hint numbers (they are visible to customers), add new hints to the end.
 

 CheckEnvVar                  "Check %1$s environment variable, its value is \"%2$s\"."
 OBSOLETE                     "You may want to use an %1$s library that supports %2$s interface with version %3$s."
 OBSOLETE                     "You may want to use an %1$s library with version %2$s."
 BadExeFormat                 "System error #193 is \"Bad format of EXE or DLL file\". "
                              "Usually it means the file is found, but it is corrupted or "
                              "a file for another architecture. "
                              "Check whether \"%1$s\" is a file for %2$s architecture."
 SystemLimitOnThreads         "System-related limit on the number of threads."
-
+SetNewBound                  "Try setting new bounds (preferably less than or equal to %1$d) for num_teams clause."
 
 
 # --------------------------------------------------------------------------------------------------
 # end of file #
 # --------------------------------------------------------------------------------------------------

openmp/runtime/src/kmp.h

 #define KMP_INTERNAL_CALLOC(n, sz) calloc((n), (sz))
 
 extern void __kmp_push_num_threads(ident_t *loc, int gtid, int num_threads);
 
 extern void __kmp_push_proc_bind(ident_t *loc, int gtid,
                                  kmp_proc_bind_t proc_bind);
 extern void __kmp_push_num_teams(ident_t *loc, int gtid, int num_teams,
                                  int num_threads);
+extern void __kmp_push_num_teams_51(ident_t *loc, int gtid, int num_teams_lb,
+                                    int num_teams_ub, int num_threads);
 
 extern void __kmp_yield();
 
 extern void __kmpc_dispatch_init_4(ident_t *loc, kmp_int32 gtid,
                                    enum sched_type schedule, kmp_int32 lb,
                                    kmp_int32 ub, kmp_int32 st, kmp_int32 chunk);
 extern void __kmpc_dispatch_init_4u(ident_t *loc, kmp_int32 gtid,
                                     enum sched_type schedule, kmp_uint32 lb,

 KMP_EXPORT void __kmpc_push_num_threads(ident_t *loc, kmp_int32 global_tid,
                                         kmp_int32 num_threads);
 
 KMP_EXPORT void __kmpc_push_proc_bind(ident_t *loc, kmp_int32 global_tid,
                                       int proc_bind);
 KMP_EXPORT void __kmpc_push_num_teams(ident_t *loc, kmp_int32 global_tid,
                                       kmp_int32 num_teams,
                                       kmp_int32 num_threads);
+/* Function for OpenMP 5.1 num_teams clause */
+KMP_EXPORT void __kmpc_push_num_teams_51(ident_t *loc, kmp_int32 global_tid,
+                                         kmp_int32 num_teams_lb,
+                                         kmp_int32 num_teams_ub,
+                                         kmp_int32 num_threads);
 KMP_EXPORT void __kmpc_fork_teams(ident_t *loc, kmp_int32 argc,
                                   kmpc_micro microtask, ...);
 struct kmp_dim { // loop bounds info casted to kmp_int64
   kmp_int64 lo; // lower
   kmp_int64 up; // upper
   kmp_int64 st; // stride
 };
 KMP_EXPORT void __kmpc_doacross_init(ident_t *loc, kmp_int32 gtid,

openmp/runtime/src/kmp_csupport.cpp

KA_TRACE(20,
            ("__kmpc_push_num_teams: enter T#%d num_teams=%d num_threads=%d\n",
             global_tid, num_teams, num_threads));
   __kmp_assert_valid_gtid(global_tid);
   __kmp_push_num_teams(loc, global_tid, num_teams, num_threads);
 }
 
 /*!
 @ingroup PARALLEL
-@param loc  source location information
+@param loc source location information
+@param global_tid global thread number
+@param num_teams_lo lower bound on number of teams requested for the teams
+construct
+@param num_teams_up upper bound on number of teams requested for the teams
+construct
+@param num_threads number of threads per team requested for the teams construct
+
+Set the number of teams to be used by the teams construct. The number of initial
+teams cretaed will be greater than or equal to the lower bound and less than or
+equal to the upper bound.
+This call is only required if the teams construct has a `num_teams` clause
+or a `thread_limit` clause (or both).
+*/
+void __kmpc_push_num_teams_51(ident_t *loc, kmp_int32 global_tid,
+                              kmp_int32 num_teams_lb, kmp_int32 num_teams_ub,
+                              kmp_int32 num_threads) {
+  KA_TRACE(20, ("__kmpc_push_num_teams_51: enter T#%d num_teams_lb=%d"
+                " num_teams_ub=%d num_threads=%d\n",
+                global_tid, num_teams_lb, num_teams_ub, num_threads));
+  __kmp_assert_valid_gtid(global_tid);
+  __kmp_push_num_teams_51(loc, global_tid, num_teams_lb, num_teams_ub,
+                          num_threads);
+}
+
+/*!
+@ingroup PARALLEL
+@param loc  source location information
 @param argc  total number of arguments in the ellipsis
 @param microtask  pointer to callback routine consisting of outlined teams
 construct
 @param ...  pointers to shared variables that aren't global
 
 Do the actual fork and call the microtask in the relevant number of threads.
 */
 void __kmpc_fork_teams(ident_t *loc, kmp_int32 argc, kmpc_micro microtask,

openmp/runtime/src/kmp_runtime.cpp

 
 void __kmp_push_num_threads(ident_t *id, int gtid, int num_threads) {
   kmp_info_t *thr = __kmp_threads[gtid];
 
   if (num_threads > 0)
     thr->th.th_set_nproc = num_threads;
 }
 
-/* this sets the requested number of teams for the teams region and/or
-   the number of threads for the next parallel region encountered  */
-void __kmp_push_num_teams(ident_t *id, int gtid, int num_teams,
+static void __kmp_push_thread_limit(kmp_info_t *thr, int num_teams,
                                     int num_threads) {
-  kmp_info_t *thr = __kmp_threads[gtid];
-  KMP_DEBUG_ASSERT(num_teams >= 0);
-  KMP_DEBUG_ASSERT(num_threads >= 0);
-
-  if (num_teams == 0) {
-    if (__kmp_nteams > 0) {
-      num_teams = __kmp_nteams;
-    } else {
-      num_teams = 1; // default number of teams is 1.
-    }
-  }
-  if (num_teams > __kmp_teams_max_nth) { // if too many teams requested?
-    if (!__kmp_reserve_warn) {
-      __kmp_reserve_warn = 1;
-      __kmp_msg(kmp_ms_warning,
-                KMP_MSG(CantFormThrTeam, num_teams, __kmp_teams_max_nth),
-                KMP_HNT(Unset_ALL_THREADS), __kmp_msg_null);
-    }
-    num_teams = __kmp_teams_max_nth;
-  }
-  // Set number of teams (number of threads in the outer "parallel" of the
-  // teams)
-  thr->th.th_set_nproc = thr->th.th_teams_size.nteams = num_teams;
-
+  KMP_DEBUG_ASSERT(thr);
   // Remember the number of threads for inner parallel regions
   if (!TCR_4(__kmp_init_middle))
     __kmp_middle_initialize(); // get internal globals calculated
   KMP_DEBUG_ASSERT(__kmp_avail_proc);
   KMP_DEBUG_ASSERT(__kmp_dflt_team_nth);
+
   if (num_threads == 0) {
     if (__kmp_teams_thread_limit > 0) {
       num_threads = __kmp_teams_thread_limit;
     } else {
       num_threads = __kmp_avail_proc / num_teams;
     }
     // adjust num_threads w/o warning as it is not user setting
     // num_threads = min(num_threads, nthreads-var, thread-limit-var)
     // no thread_limit clause specified -  do not change thread-limit-var ICV
     if (num_threads > __kmp_dflt_team_nth) {
       num_threads = __kmp_dflt_team_nth; // honor nthreads-var ICV
     }
     if (num_threads > thr->th.th_current_task->td_icvs.thread_limit) {
       num_threads = thr->th.th_current_task->td_icvs.thread_limit;
     } // prevent team size to exceed thread-limit-var
     if (num_teams * num_threads > __kmp_teams_max_nth) {
       num_threads = __kmp_teams_max_nth / num_teams;
     }
+    if (num_threads == 0) {
+      num_threads = 1;
+    }
   } else {
     // This thread will be the master of the league masters
     // Store new thread limit; old limit is saved in th_cg_roots list
     thr->th.th_current_task->td_icvs.thread_limit = num_threads;
     // num_threads = min(num_threads, nthreads-var)
     if (num_threads > __kmp_dflt_team_nth) {
       num_threads = __kmp_dflt_team_nth; // honor nthreads-var ICV
     }
     if (num_teams * num_threads > __kmp_teams_max_nth) {
       int new_threads = __kmp_teams_max_nth / num_teams;
+      if (new_threads == 0) {
+        new_threads = 1;
+      }
+      if (new_threads != num_threads) {
         if (!__kmp_reserve_warn) { // user asked for too many threads
           __kmp_reserve_warn = 1; // conflicts with KMP_TEAMS_THREAD_LIMIT
           __kmp_msg(kmp_ms_warning,
                     KMP_MSG(CantFormThrTeam, num_threads, new_threads),
                     KMP_HNT(Unset_ALL_THREADS), __kmp_msg_null);
         }
+      }
       num_threads = new_threads;
     }
   }
   thr->th.th_teams_size.nth = num_threads;
 }
 
+/* this sets the requested number of teams for the teams region and/or
+   the number of threads for the next parallel region encountered  */
+void __kmp_push_num_teams(ident_t *id, int gtid, int num_teams,
+                          int num_threads) {
+  kmp_info_t *thr = __kmp_threads[gtid];
+  KMP_DEBUG_ASSERT(num_teams >= 0);
+  KMP_DEBUG_ASSERT(num_threads >= 0);
+
+  if (num_teams == 0) {
+    if (__kmp_nteams > 0) {
+      num_teams = __kmp_nteams;
+    } else {
+      num_teams = 1; // default number of teams is 1.
+    }
+  }
+  if (num_teams > __kmp_teams_max_nth) { // if too many teams requested?
+    if (!__kmp_reserve_warn) {
+      __kmp_reserve_warn = 1;
+      __kmp_msg(kmp_ms_warning,
+                KMP_MSG(CantFormThrTeam, num_teams, __kmp_teams_max_nth),
+                KMP_HNT(Unset_ALL_THREADS), __kmp_msg_null);
+    }
+    num_teams = __kmp_teams_max_nth;
+  }
+  // Set number of teams (number of threads in the outer "parallel" of the
+  // teams)
+  thr->th.th_set_nproc = thr->th.th_teams_size.nteams = num_teams;
+
+  __kmp_push_thread_limit(thr, num_teams, num_threads);
+}
+
+/* This sets the requested number of teams for the teams region and/or
+   the number of threads for the next parallel region encountered  */
+void __kmp_push_num_teams_51(ident_t *id, int gtid, int num_teams_lb,
+                             int num_teams_ub, int num_threads) {
+  kmp_info_t *thr = __kmp_threads[gtid];
+  KMP_DEBUG_ASSERT(num_teams_lb >= 0 && num_teams_ub >= 0);
+  KMP_DEBUG_ASSERT(num_teams_ub >= num_teams_lb);
+  KMP_DEBUG_ASSERT(num_threads >= 0);
+
+  if (num_teams_lb > num_teams_ub) {
+    __kmp_fatal(KMP_MSG(FailedToCreateTeam, num_teams_lb, num_teams_ub),
+                KMP_HNT(SetNewBound, __kmp_teams_max_nth), __kmp_msg_null);
+  }
+
+  int num_teams = 1; // defalt number of teams is 1.
+
+  if (num_teams_lb == 0 && num_teams_ub > 0)
+    num_teams_lb = num_teams_ub;
+
+  if (num_teams_lb == 0 && num_teams_ub == 0) { // no num_teams clause
+    num_teams = (__kmp_nteams > 0) ? __kmp_nteams : num_teams;
+    if (num_teams > __kmp_teams_max_nth) {
+      if (!__kmp_reserve_warn) {
+        __kmp_reserve_warn = 1;
+        __kmp_msg(kmp_ms_warning,
+                  KMP_MSG(CantFormThrTeam, num_teams, __kmp_teams_max_nth),
+                  KMP_HNT(Unset_ALL_THREADS), __kmp_msg_null);
+      }
+      num_teams = __kmp_teams_max_nth;
+    }
+  } else if (num_teams_lb == num_teams_ub) { // requires exact number of teams
+    num_teams = num_teams_ub;
+  } else { // num_teams_lb <= num_teams <= num_teams_ub
+    if (num_threads == 0) {
+      if (num_teams_ub > __kmp_teams_max_nth) {
+        num_teams = num_teams_lb;
+      } else {
+        num_teams = num_teams_ub;
+      }
+    } else {
+      num_teams = (num_threads > __kmp_teams_max_nth)
+                      ? num_teams
+                      : __kmp_teams_max_nth / num_threads;
+      if (num_teams < num_teams_lb) {
+        num_teams = num_teams_lb;
+      } else if (num_teams > num_teams_ub) {
+        num_teams = num_teams_ub;
+      }
+    }
+  }
+  // Set number of teams (number of threads in the outer "parallel" of the
+  // teams)
+  thr->th.th_set_nproc = thr->th.th_teams_size.nteams = num_teams;
+
+  __kmp_push_thread_limit(thr, num_teams, num_threads);
+}
+
 // Set the proc_bind var to use in the following parallel region.
 void __kmp_push_proc_bind(ident_t *id, int gtid, kmp_proc_bind_t proc_bind) {
   kmp_info_t *thr = __kmp_threads[gtid];
   thr->th.th_set_proc_bind = proc_bind;
 }
 
 /* Launch the worker threads into the microtask. */
 

openmp/runtime/test/teams/kmp_num_teams.c

+// RUN: %libomp-compile-and-run
+// UNSUPPORTED: gcc
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+#define NT 8
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+typedef int kmp_int32;
+typedef struct ident {
+  kmp_int32 reserved_1;
+  kmp_int32 flags;
+  kmp_int32 reserved_2;
+  kmp_int32 reserved_3;
+  char const *psource;
+} ident_t;
+extern int __kmpc_global_thread_num(ident_t *);
+extern void __kmpc_push_num_teams_51(ident_t *, kmp_int32, kmp_int32, kmp_int32,
+                                     kmp_int32);
+#ifdef __cplusplus
+}
+#endif
+
+void check_num_teams(int num_teams_lb, int num_teams_ub, int thread_limit) {
+  int nteams, nthreads;
+  int a = 0;
+
+  int gtid = __kmpc_global_thread_num(NULL);
+  __kmpc_push_num_teams_51(NULL, gtid, num_teams_lb, num_teams_ub,
+                           thread_limit);
+
+#pragma omp target teams
+  {
+    int priv_nteams;
+    int team_num = omp_get_team_num();
+    if (team_num == 0)
+      nteams = omp_get_num_teams();
+    priv_nteams = omp_get_num_teams();
+#pragma omp parallel
+    {
+      int priv_nthreads;
+      int thread_num = omp_get_thread_num();
+      int teams_ub, teams_lb, thr_limit;
+      if (team_num == 0 && thread_num == 0)
+        nthreads = omp_get_num_threads();
+      priv_nthreads = omp_get_num_threads();
+
+      teams_ub = (num_teams_ub ? num_teams_ub : priv_nteams);
+      teams_lb = (num_teams_lb ? num_teams_lb : teams_ub);
+      thr_limit = (thread_limit ? thread_limit : priv_nthreads);
+
+      if (priv_nteams < teams_lb || priv_nteams > teams_ub) {
+        fprintf(stderr, "error: invalid number of teams=%d\n", priv_nteams);
+        exit(1);
+      }
+      if (priv_nthreads > thr_limit) {
+        fprintf(stderr, "error: invalid number of threads=%d\n", priv_nthreads);
+        exit(1);
+      }
+#pragma omp atomic
+      a++;
+    }
+  }
+  if (a != nteams * nthreads) {
+    fprintf(stderr, "error: a (%d) != nteams * nthreads (%d)\n", a,
+            nteams * nthreads);
+    exit(1);
+  } else {
+    printf("#teams %d, #threads %d: Hello!\n", nteams, nthreads);
+  }
+}
+
+int main(int argc, char *argv[]) {
+  omp_set_num_threads(NT);
+
+  check_num_teams(1, 8, 2);
+  check_num_teams(2, 2, 2);
+  check_num_teams(2, 2, 0);
+  check_num_teams(8, 16, 2);
+  check_num_teams(9, 16, 0);
+  check_num_teams(9, 16, 2);
+  check_num_teams(2, 3, 0);
+  check_num_teams(0, 0, 2);
+  check_num_teams(0, 4, 0);
+  check_num_teams(0, 2, 2);
+
+  printf("Test Passed\n");
+  return 0;
+}