Spending too much time with OpenMP, I got dragged into this little
private weekend project. OpenMP is slowly adding features that permit
writing code as fast as with low-level GPU languages like CUDA and its
cousin HIP. This consists of four parts as explained
in Kevin Sala's talk "Accelerating with OpenMP - Latest and Upcoming
Features" at https://link.openmp.org/sc25talks
The attached patch is about the section "Multidimensional Grid
Programming" in the talk. The actual feature requires quite some
work; however, adding the new API routines is rather simple.
Hence:
This patch adds OpenMP 6.1's multidimensional support routines,
or to be precise the routines added to the OpenMP 6.1 draft, to
be released as TR 15. (The final OpenMP 6.1 spec is expected to be released in
early November in time for SC26.)
Without multi-dim support in GCC those routines are rather boring:
omp_get_num_threads_dim yields for dim == 0 the same result as the
existing unidimensional omp_get_num_threads - and for dim > 0
it always returns 1.
Thus, the routines are working & compliant but don't provide any
benefit until actual multi-dim support is added to GCC.
However, there is an exception:omp_get_supported_teams_dim andomp_get_supported_threads_dim actually
return (for dim == 0) an upper bound bound for the supported number of
teams and threads, which can be useful: On the host, that's just
INT_MAX. (No algorithmic limit but at least num_threads(INT_MAX) is
likely to fail with out of memory.) However, for the GPUs devices, it
shows actually useful limits based on GCC's used teams/thread limits.
(Some are due to implementation choices others actual some hardware
restrictions.) Those are upper bounds, i.e. expect that the actual
values are lower. In particular, specifying values in multiple
dimensions (once support) whose product exceeds the dim==0 value will
surely fail but also mixing threads + teams or ... will lower it. On my
Laptop, I get for #include <omp.h> int main() { for (int i = -1; i <
omp_get_num_devices (); i++) __builtin_printf("%d: %d / %d\n", i,
omp_get_supported_teams_dim(i, 0), omp_get_supported_threads_dim(i,0));
} the following output for the initial device (host, '-1') and an Nvidia
GPU - namely, for teams and threads the upper bound is: -1: 2147483647 /
2147483647 0: 40960 / 32 And I think that can be useful information! * *
* Any comments, remarks, suggestions before I commit it? Tobias
libgomp: Add stubs for OpenMP's multidim routines
This commit adds OpenMP 6.1's multidimensional support routines. 'stub' is
slightly misleading as they are fully functional; however, as GCC currently
only supports a single active spatial dimension for both teams and target,
the non-'omp_supported' functions act as their unidimensional counter part
for dimension 0. And the supported-active-dims functions just return one.
However, the functions omp_get_supported_teams_dim and
omp_get_supported_threads_dim actually return (for dim == 0) an upper bound
for the number of teams and threads, respectively, that can be useful to
know for unidimensional use as well.
gcc/ChangeLog:
* omp-general.cc (omp_runtime_api_procname): Update for added
functions and older additions.
* omp-low.cc (scan_omp_1_stmt): Update for added _dim function.
libgomp/ChangeLog:
* config/gcn/teams.c (omp_get_num_teams_dim,
omp_get_team_num_dim): New functions.
* config/nvptx/teams.c (omp_get_num_teams_dim,
omp_get_team_num_dim): Likewise.
* fortran.c (omp_get_supported_active_team_dims_,
omp_get_supported_active_league_dims_): Likewise.
* icv.c (omp_get_thread_limit_dim): Likewise.
* libgomp-plugin.h (GOMP_OFFLOAD_get_numa_node): Remove spurious
const of 'const int' function return type.
(GOMP_OFFLOAD_supported_teams_dim,
GOMP_OFFLOAD_supported_threads_dim): Declare.
* libgomp.h (struct gomp_device_descr): Add supported_teams_dim_func
and supported_threads_dim_func members.
* libgomp.map (OMP_6.1): Add new OpenMP routines.
* libgomp.texi (Multidimensional Support Routines): Add new section.
(omp_get_device_distances): Fix chapter of spec reference.
* omp.h.in (omp_get_num_threads_dim, omp_get_thread_num_dim,
omp_get_thread_limit_dim, omp_get_supported_active_team_dims,
omp_get_supported_active_league_dims, omp_get_supported_teams_dim,
omp_get_supported_threads_dim, omp_get_num_teams_dim,
omp_get_team_num_dim, omp_get_device_distances): Declare.
* omp_lib.f90.in: Likewise.
* omp_lib.h.in: Likewise.
* parallel.c (omp_get_num_threads_dim,
omp_get_thread_num_dim): New functions.
* plugin/plugin-gcn.c (limit_worker_threads): Add comment.
(GOMP_OFFLOAD_supported_teams_dim,
GOMP_OFFLOAD_supported_threads_dim): New functions.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_teams_dim,
GOMP_OFFLOAD_supported_threads_dim): New functions.
(nvptx_adjust_launch_bounds): Add comment.
* target.c (omp_get_supported_active_team_dims,
omp_get_supported_active_league_dims,
omp_get_supported_teams_dim,
omp_get_supported_threads_dim): New functions.
(gomp_load_plugin_for_device): dlsym added functions.
* teams.c (omp_get_num_teams_dim,
omp_get_team_num_dim): New functions.
* testsuite/libgomp.c/multi-dim-routines.c: New test.
* testsuite/libgomp.fortran/multi-dim-routines-2.F90: New test.
* testsuite/libgomp.fortran/multi-dim-routines.F90: New test.
gcc/omp-general.cc | 16 +-
gcc/omp-low.cc | 10 +-
libgomp/config/gcn/teams.c | 18 +
libgomp/config/nvptx/teams.c | 18 +
libgomp/fortran.c | 14 +
libgomp/icv.c | 9 +
libgomp/libgomp-plugin.h | 4 +-
libgomp/libgomp.h | 2 +
libgomp/libgomp.map | 11 +
libgomp/libgomp.texi | 394 ++++++++++++++++++++-
libgomp/omp.h.in | 11 +-
libgomp/omp_lib.f90.in | 84 +++++
libgomp/omp_lib.h.in | 85 +++++
libgomp/parallel.c | 18 +
libgomp/plugin/plugin-gcn.c | 26 ++
libgomp/plugin/plugin-nvptx.c | 47 ++-
libgomp/target.c | 54 +++
libgomp/teams.c | 18 +
libgomp/testsuite/libgomp.c/multi-dim-routines.c | 135 +++++++
.../libgomp.fortran/multi-dim-routines-2.F90 | 2 +
.../libgomp.fortran/multi-dim-routines.F90 | 149 ++++++++
21 files changed, 1113 insertions(+), 12 deletions(-)
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index e4dcf063a5b..d0307aeba88 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -3959,16 +3959,25 @@ omp_runtime_api_procname (const char *name)
static const char *omp_runtime_apis[] =
{
/* This array has 3 sections. First omp_* calls that don't
- have any suffixes. */
+ have any suffixes; this implies bind(C) either in the
+ specification or implementation choice. */
"aligned_alloc",
"aligned_calloc",
"alloc",
"calloc",
"free",
+ "get_device_distances",
"get_interop_int",
"get_interop_ptr",
"get_mapped_ptr",
"get_num_interop_properties",
+ "get_num_teams_dim",
+ "get_num_threads_dim",
+ "get_supported_teams_dim",
+ "get_supported_threads_dim",
+ "get_team_num_dim",
+ "get_thread_limit_dim",
+ "get_thread_num_dim",
"realloc",
"target_alloc",
"target_associate_ptr",
@@ -4016,7 +4025,9 @@ omp_runtime_api_procname (const char *name)
"get_partition_num_places",
"get_place_num",
"get_proc_bind",
+ "get_supported_active_league_dims",
"get_supported_active_levels",
+ "get_supported_active_team_dims",
"get_team_num",
"get_teams_thread_limit",
"get_thread_limit",
@@ -4059,7 +4070,8 @@ omp_runtime_api_procname (const char *name)
"set_num_teams",
"set_num_threads",
"set_schedule",
- "set_teams_thread_limit"
+ "set_teams_thread_limit",
+ "target_memset"
};
int mode = 0;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 02b62d54ba6..eef0b418b9b 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4167,10 +4167,18 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
!= strlen ("omp_get_num_teams"))
|| strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
"omp_get_num_teams") != 0)
+ && ((IDENTIFIER_LENGTH (DECL_NAME (fndecl))
+ != strlen ("omp_get_num_teams_dim"))
+ || strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
+ "omp_get_num_teams_dim") != 0)
&& ((IDENTIFIER_LENGTH (DECL_NAME (fndecl))
!= strlen ("omp_get_team_num"))
|| strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
- "omp_get_team_num") != 0))
+ "omp_get_team_num") != 0)
+ && ((IDENTIFIER_LENGTH (DECL_NAME (fndecl))
+ != strlen ("omp_get_team_num_dim"))
+ || strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
+ "omp_get_team_num_dim") != 0))
{
remove = true;
error_at (gimple_location (stmt),
diff --git a/libgomp/config/gcn/teams.c b/libgomp/config/gcn/teams.c
index 42570eb445e..793d4bb7838 100644
--- a/libgomp/config/gcn/teams.c
+++ b/libgomp/config/gcn/teams.c
@@ -44,6 +44,14 @@ omp_get_num_teams (void)
return gomp_num_teams_var + 1;
}
+int
+omp_get_num_teams_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_num_teams ();
+ return 1;
+}
+
int
omp_get_team_num (void)
{
@@ -51,5 +59,15 @@ omp_get_team_num (void)
return *gomp_team_num;
}
+int
+omp_get_team_num_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_team_num ();
+ return 0;
+}
+
ialias (omp_get_num_teams)
+ialias (omp_get_num_teams_dim)
ialias (omp_get_team_num)
+ialias (omp_get_team_num_dim)
diff --git a/libgomp/config/nvptx/teams.c b/libgomp/config/nvptx/teams.c
index 4af9aaef18a..af6bf6c5710 100644
--- a/libgomp/config/nvptx/teams.c
+++ b/libgomp/config/nvptx/teams.c
@@ -47,11 +47,29 @@ omp_get_num_teams (void)
return gomp_num_teams_var + 1;
}
+int
+omp_get_num_teams_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_num_teams ();
+ return 1;
+}
+
int
omp_get_team_num (void)
{
return __gomp_team_num;
}
+int
+omp_get_team_num_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_team_num ();
+ return 0;
+}
+
ialias (omp_get_num_teams)
+ialias (omp_get_num_teams_dim)
ialias (omp_get_team_num)
+ialias (omp_get_team_num_dim)
diff --git a/libgomp/fortran.c b/libgomp/fortran.c
index 59d1de2886e..a28b86ae666 100644
--- a/libgomp/fortran.c
+++ b/libgomp/fortran.c
@@ -104,6 +104,8 @@ ialias_redirect (omp_get_interop_name)
ialias_redirect (omp_get_interop_type_desc)
ialias_redirect (omp_get_interop_rc_desc)
ialias_redirect (omp_control_tool)
+ialias_redirect (omp_get_supported_active_team_dims)
+ialias_redirect (omp_get_supported_active_league_dims)
#endif
#ifndef LIBGOMP_GNU_SYMBOL_VERSIONING
@@ -861,6 +863,18 @@ omp_get_uid_from_device_8_ (const char **res, size_t *res_len,
#ifndef LIBGOMP_OFFLOADED_ONLY
+int
+omp_get_supported_active_team_dims_ (int *device_num, int *native_support)
+{
+ return omp_get_supported_active_team_dims (*device_num, *native_support);
+}
+
+int
+omp_get_supported_active_league_dims_ (int *device_num, int *native_support)
+{
+ return omp_get_supported_active_league_dims (*device_num, *native_support);
+}
+
void
omp_display_env_ (const int32_t *verbose)
{
diff --git a/libgomp/icv.c b/libgomp/icv.c
index fc6ad65074b..6c331451983 100644
--- a/libgomp/icv.c
+++ b/libgomp/icv.c
@@ -118,6 +118,14 @@ omp_get_thread_limit (void)
return icv->thread_limit_var > INT_MAX ? INT_MAX : icv->thread_limit_var;
}
+int
+omp_get_thread_limit_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_thread_limit ();
+ return 1;
+}
+
void
omp_set_max_active_levels (int max_levels)
{
@@ -239,6 +247,7 @@ ialias (omp_set_schedule)
ialias (omp_get_schedule)
ialias (omp_get_max_threads)
ialias (omp_get_thread_limit)
+ialias (omp_get_thread_limit_dim)
ialias (omp_set_max_active_levels)
ialias (omp_get_max_active_levels)
ialias (omp_get_supported_active_levels)
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 88841ac5817..60ac7201597 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -159,7 +159,9 @@ extern void GOMP_PLUGIN_target_rev (uint64_t, uint64_t, uint64_t, uint64_t,
/* Prototypes for functions implemented by libgomp plugins. */
extern const char *GOMP_OFFLOAD_get_name (void);
extern const char *GOMP_OFFLOAD_get_uid (int);
-extern const int GOMP_OFFLOAD_get_numa_node (int);
+extern int GOMP_OFFLOAD_get_numa_node (int);
+extern int GOMP_OFFLOAD_supported_teams_dim (int, int);
+extern int GOMP_OFFLOAD_supported_threads_dim (int, int);
extern unsigned int GOMP_OFFLOAD_get_caps (void);
extern int GOMP_OFFLOAD_get_type (void);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index f2aabdbc806..9510b148ed8 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1449,6 +1449,8 @@ struct gomp_device_descr
__typeof (GOMP_OFFLOAD_get_name) *get_name_func;
__typeof (GOMP_OFFLOAD_get_uid) *get_uid_func;
__typeof (GOMP_OFFLOAD_get_numa_node) *get_numa_node_func;
+ __typeof (GOMP_OFFLOAD_supported_teams_dim) *supported_teams_dim_func;
+ __typeof (GOMP_OFFLOAD_supported_threads_dim) *supported_threads_dim_func;
__typeof (GOMP_OFFLOAD_get_caps) *get_caps_func;
__typeof (GOMP_OFFLOAD_get_type) *get_type_func;
__typeof (GOMP_OFFLOAD_get_num_devices) *get_num_devices_func;
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 7e8bb7ebec3..dbde88c4d5f 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -250,6 +250,17 @@ OMP_6.0 {
OMP_6.1 {
global:
omp_get_device_distances;
+ omp_get_num_teams_dim;
+ omp_get_num_threads_dim;
+ omp_get_supported_active_league_dims;
+ omp_get_supported_active_league_dims_;
+ omp_get_supported_active_team_dims;
+ omp_get_supported_active_team_dims_;
+ omp_get_supported_teams_dim;
+ omp_get_supported_threads_dim;
+ omp_get_team_num_dim;
+ omp_get_thread_limit_dim;
+ omp_get_thread_num_dim;
} OMP_6.0;
GOMP_1.0 {
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 9fcfbf96963..9482e8753a2 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -701,6 +701,7 @@ specification in version 5.2.
* Thread Team Routines::
* Thread Affinity Routines::
* Teams Region Routines::
+* Multidimensional Support Routines::
* Tasking Routines::
* Resource Relinquishing Routines::
* Device Information Routines::
@@ -784,7 +785,7 @@ the program @code{omp_get_num_threads} returns 1.
The default team size may be initialized at startup by the
@env{OMP_NUM_THREADS} environment variable. At runtime, the size
-of the current team may be set either by the @code{NUM_THREADS}
+of the current team may be set either by the @code{num_threads}
clause or by @code{omp_set_num_threads}. If none of the above were
used to define a specific value and @env{OMP_DYNAMIC} is disabled,
one thread per CPU online is used.
@@ -800,7 +801,8 @@ one thread per CPU online is used.
@end multitable
@item @emph{See also}:
-@ref{omp_get_max_threads}, @ref{omp_set_num_threads}, @ref{OMP_NUM_THREADS}
+@ref{omp_get_max_threads}, @ref{omp_set_num_threads},
+@ref{omp_get_num_threads_dim}, @ref{OMP_NUM_THREADS}
@item @emph{Reference}:
@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.2.
@@ -855,7 +857,8 @@ value of the primary thread of a team is always 0.
@end multitable
@item @emph{See also}:
-@ref{omp_get_num_threads}, @ref{omp_get_ancestor_thread_num}
+@ref{omp_get_thread_num_dim}, @ref{omp_get_num_threads},
+@ref{omp_get_ancestor_thread_num}
@item @emph{Reference}:
@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.4.
@@ -1421,6 +1424,9 @@ Returns the number of teams in the current team region.
@item @emph{Interface}: @tab @code{integer function omp_get_num_teams()}
@end multitable
+@item @emph{See also}:
+@ref{omp_get_num_teams_dim}, @ref{OMP_NUM_TEAMS}, @ref{omp_get_team_num}
+
@item @emph{Reference}:
@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.32.
@end table
@@ -1443,6 +1449,9 @@ Returns the team number of the calling thread.
@item @emph{Interface}: @tab @code{integer function omp_get_team_num()}
@end multitable
+@item @emph{See also}:
+@ref{omp_get_num_teams_dim}, @ref{OMP_NUM_TEAMS}, @ref{omp_get_num_teams}
+
@item @emph{Reference}:
@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.33.
@end table
@@ -1536,7 +1545,9 @@ for each team created by the teams construct which does not specify a
@subsection @code{omp_get_thread_limit} -- Maximum number of threads
@table @asis
@item @emph{Description}:
-Return the maximum number of threads of the program.
+Return the maximal number of threads that are available in the current
+contention group to execute tasks, i.e. the unidimenional value of
+the @var{thread-limit-var} ICV.
@item @emph{C/C++}:
@multitable @columnfractions .20 .80
@@ -1549,7 +1560,7 @@ Return the maximum number of threads of the program.
@end multitable
@item @emph{See also}:
-@ref{omp_get_max_threads}, @ref{OMP_THREAD_LIMIT}
+@ref{omp_get_thread_limit_dim}, @ref{omp_get_max_threads}, @ref{OMP_THREAD_LIMIT}
@item @emph{Reference}:
@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.14.
@@ -1557,6 +1568,377 @@ Return the maximum number of threads of the program.
+@node Multidimensional Support Routines
+@section Multidimensional Support Routines
+
+Routines that retrieve information about the execution environment with
+multiple spatial dimensions. They have C linkage and do not throw exceptions.
+
+The dimensional space is unidimensional, if the number of active spaces
+is one. To create a multidimensional space, the @code{dim} modifier must
+be specified with a dimension greater than one in the @code{num_teams},
+@code{num_thread_limits}, or @code{num_threads} clause.
+
+The relation between the unidimensional and the spatial values is as
+follows: For the size, the unidimensional size is the product of the
+spatial sizes. For the thread or team number, the unidimensional number
+is the respective index of dimension 0 plus the index of dimension 1
+times the size of dimension 0 plus the index of dimension 2 times the
+size of dimension 1 times the size of dimension 0 etc. If the active
+spatial dimensions is one, the spatial dimensional result for dimension
+zero is the same as returned by the respective unidimensional routine.
+
+Note that there is no means to obtain the number of active dimensions in
+the general case as the size of an inactive spatial dimension is one,
+which is also a permitted size for an active spatial dimension.
+
+@menu
+* omp_get_num_threads_dim:: Size of a dimension of the active team
+* omp_get_thread_num_dim:: Dimensional ID of the current thread
+* omp_get_num_teams_dim:: Number of teams along a dimension
+* omp_get_team_num_dim:: Get team number in a dimension
+* omp_get_thread_limit_dim:: Maximum number of threads in a dimension
+* omp_get_supported_active_team_dims:: Number of supported team dimensions
+* omp_get_supported_active_league_dims:: Number of supported thread dimensions
+* omp_get_supported_threads_dim:: Get upper limit of threads in a dimension
+* omp_get_supported_teams_dim:: Get upper limit of teams in a dimension
+@end menu
+
+
+@node omp_get_num_threads_dim
+@subsection @code{omp_get_num_threads_dim} -- Size of a dimension of the active team
+@table @asis
+Returns the size in dimension @var{DIM} of threads in the current team and,
+therefore, of the parallel region that binds to that team, i.e. the value
+of the @var{team-size-var} ICV in dimension @var{dim}. For inactive spatial
+dimensions, the routine returns 1. The dimension argument @var{DIM} must be
+non-negative.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_num_threads_dim(int dim)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_num_threads_dim(dim)}
+@item @tab @code{integer :: dim}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_num_threads}, @ref{omp_get_thread_num_dim},
+@ref{omp_get_thread_limit_dim}, @ref{omp_get_max_threads},
+@ref{omp_set_num_threads}, @ref{OMP_NUM_THREADS},
+@ref{omp_get_supported_active_league_dims}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.1
+@end table
+
+
+
+@node omp_get_thread_num_dim
+@subsection @code{omp_get_thread_num_dim} -- Dimensional ID of the current thread
+@table @asis
+@item @emph{Description}:
+Returns the value of dimension @var{dim} of the multidimensional thread
+identification number for the invoking thread in the current team, i.e. the
+value of the @var{thread-num-var} ICV in dimension @var{dim}. For inactive
+spatial dimensions, the routine returns 0. The dimension argument @var{DIM}
+must be non-negative.
+
+The returned value is between zero and one less than the value returned by
+@code{omp_get_num_threads} for the dimension @var{dim}.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_thread_num(int dim)}
+@item @tab @code{integer :: dim}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_thread_num(dim)}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_thread_num}, @ref{omp_get_num_threads_dim}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.2
+@end table
+
+
+
+@node omp_get_num_teams_dim
+@subsection @code{omp_get_num_teams_dim} -- Number of teams along a dimension
+@table @asis
+@item @emph{description}:
+Returns the size in dimension @var{DIM} of the number of initial teams, i.e.
+the value of the @var{league-size-var} ICV in dimension @var{dim}. For inactive
+spatial dimensions, the routine returns 1. The dimension argument @var{DIM}
+must be non-negative.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{c/c++}:
+@multitable @columnfractions .20 .80
+@item @emph{prototype}: @tab @code{int omp_get_num_teams(int dim)}
+@end multitable
+
+@item @emph{fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{interface}: @tab @code{integer function omp_get_num_teams(dim)}
+@item @tab @code{integer :: dim}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_num_teams}, @ref{OMP_NUM_TEAMS}, @ref{omp_get_team_num_dim}
+@ref{omp_get_supported_active_team_dims}
+
+
+@item @emph{reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.3
+@end table
+
+
+
+@node omp_get_team_num_dim
+@subsection @code{omp_get_team_num_dim} -- Get team number in a dimension
+@table @asis
+@item @emph{description}:
+Returns the value of dimension @var{dim} of the multidimensional team
+identification number in the current team, i.e. the value of the
+@var{team-num-var} ICV in dimension @var{dim}. For inactive spatial
+dimensions, the routine returns 0. The dimension argument @var{DIM}
+must be non-negative.
+
+The returned value is between zero and one less than the value returned by
+@code{omp_get_num_teams_dim} for the dimension @var{dim}.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{c/c++}:
+@multitable @columnfractions .20 .80
+@item @emph{prototype}: @tab @code{int omp_get_team_num_dim(int dim);}
+@end multitable
+
+@item @emph{fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{interface}: @tab @code{integer function omp_get_team_num_dim(dim)}
+@item @tab @code{integer :: dim}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_team_num_dim}, @ref{OMP_NUM_TEAMS}, @ref{omp_get_num_teams_dim}
+
+@item @emph{reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.4
+@end table
+
+
+
+@node omp_get_thread_limit_dim
+@subsection @code{omp_get_thread_limit_dim} -- Maximum number of threads in a dimension
+@table @asis
+@item @emph{Description}:
+Return the maximal number of threads in dimension @var{dim} that are available
+in the current contention group in that dimension, i.e. the value of the
+@var{thread-limit-var} ICV in dimension @var{dim}. The routine returns one if
+the dimension @var{dim} exceeds the number of dimensions for which values have
+been specified in the @var{thread-limit-var} ICV. The dimension argument
+@var{DIM} must be non-negative.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_thread_limit_dim(int dim)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_thread_limit_dim(dim)}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_thread_limit}, @ref{omp_get_thread_limit_dim},
+@ref{omp_get_max_threads}, @ref{OMP_THREAD_LIMIT}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.5
+@end table
+
+
+
+@node omp_get_supported_active_team_dims
+@subsection @code{omp_get_supported_active_team_dims} -- Number of supported team dimensions
+@table @asis
+@item @emph{Description}:
+The routine returns the number of active spatial dimensions that the device with
+@var{device_num} supports for teams. If @var{native_support} is set to true, it
+returns the maximal number active spatial dimensions that are natively supported
+by the device.
+
+The effect when the routine is invoked from within a @code{target} region is unspecified.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_supported_active_team_dims(int device_num, int native_support)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_supported_active_team_dims(device_num, native_support)}
+@item @tab @code{integer :: device_num}
+@item @tab @code{logical :: native_support}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_supported_teams_dim}, @ref{omp_get_num_teams_dim}, @ref{OMP_NUM_TEAMS}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.6
+@end table
+
+
+
+@node omp_get_supported_active_league_dims
+@subsection @code{omp_get_supported_active_league_dims} -- Number of supported thread dimensions
+@table @asis
+@item @emph{Description}:
+The routine returns the number of active spatial dimensions that the device with
+@var{device_num} supports for league spaces, i.e. the number of threads available
+to each team in the league. If @var{native_support} is set to true, it returns
+the maximal number active spatial dimensions that are natively supported by the
+device.
+
+The effect when the routine is invoked from within a @code{target} region is unspecified.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_supported_active_league_dims(int device_num, int native_support)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_supported_active_league_dims(device_num, native_support)}
+@item @tab @code{integer :: device_num}
+@item @tab @code{logical :: native_support}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_supported_threads_dim}, @ref{omp_get_thread_limit_dim},
+@ref{omp_get_num_teams_dim}, @ref{OMP_THREAD_LIMIT}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.7
+@end table
+
+
+
+@node omp_get_supported_threads_dim
+@subsection @code{omp_get_supported_threads_dim} -- Get upper limit of threads in a dimension
+@table @asis
+@item @emph{Description}:
+The routine returns an upper bound for the number of threads supported in
+dimension @var{dim} on device @var{device_num}. If @var{dim} specifies a number
+that exceeds the number of supported active spatial dimensions for threads, the
+routine returns one.
+
+Note: The actually supported number of threads by dimension might be lower,
+depending on used number of threads in other dimensions, the number of
+teams, and other parameters like the register use of the code to be executed.
+
+The effect when the routine is invoked from within a @code{target} region is unspecified.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_supported_threads_dim(int device_num, int native_support)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_supported_threads_dim(device_num, native_support)}
+@item @tab @code{integer :: device_num}
+@item @tab @code{logical :: native_support}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_supported_active_league_dims}, @ref{omp_get_thread_limit_dim},
+@ref{omp_get_max_threads}, @ref{omp_get_num_teams_dim}, @ref{OMP_THREAD_LIMIT}
+
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.8
+@end table
+
+
+
+@node omp_get_supported_teams_dim
+@subsection @code{omp_get_supported_teams_dim} -- Get upper limit of teams in a dimension
+@table @asis
+@item @emph{Description}:
+The routine returns an upper bound for the number of teams supported in
+dimension @var{dim} on device @var{device_num}. If @var{dim} specifies a number
+that exceeds the number of supported active spatial dimensions for teams, the
+routine returns one.
+
+Note: The actually supported number of teams by dimension might be lower,
+depending on used number of teams in other dimensions and other parameters.
+
+The effect when the routine is invoked from within a @code{target} region is unspecified.
+
+Fortran note: In GCC, this routine does not support passing default-kind integer
+arguments with @code{-fdefault-integer-8}.
+
+@item @emph{C/C++}:
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_supported_teams_dim(int device_num, int dim)}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_supported_teams_dim(device_num, dim)}
+@item @tab @code{integer :: device_num, dim}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_supported_active_team_dims}, @ref{omp_get_num_teams_dim}, @ref{OMP_NUM_TEAMS}
+
+@item @emph{Reference}:
+@c TODO: Update once OpenMP 6.1 is released
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 29.9
+@end table
+
+
+
@node Tasking Routines
@section Tasking Routines
@@ -2064,7 +2446,7 @@ arguments with @code{-fdefault-integer-8}.
@item @emph{Reference}:
@c TODO: Update once OpenMP 6.1 is released
-@uref{https://www.openmp.org, OpenMP specification Technical Report 15}, Section 3.16
+@uref{https://www.openmp.org, OpenMP Technical Report 15}, Section 31.16
@end table
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index cfb104dccbd..605ea9bb5a6 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -269,8 +269,10 @@ extern "C" {
extern void omp_set_num_threads (int) __GOMP_NOTHROW;
extern int omp_get_num_threads (void) __GOMP_NOTHROW;
+extern int omp_get_num_threads_dim (int) __GOMP_NOTHROW;
extern int omp_get_max_threads (void) __GOMP_NOTHROW;
extern int omp_get_thread_num (void) __GOMP_NOTHROW;
+extern int omp_get_thread_num_dim (int) __GOMP_NOTHROW;
extern int omp_get_num_procs (void) __GOMP_NOTHROW;
extern int omp_in_parallel (void) __GOMP_NOTHROW;
@@ -303,8 +305,13 @@ extern double omp_get_wtick (void) __GOMP_NOTHROW;
extern void omp_set_schedule (omp_sched_t, int) __GOMP_NOTHROW;
extern void omp_get_schedule (omp_sched_t *, int *) __GOMP_NOTHROW;
extern int omp_get_thread_limit (void) __GOMP_NOTHROW;
+extern int omp_get_thread_limit_dim (int) __GOMP_NOTHROW;
extern void omp_set_max_active_levels (int) __GOMP_NOTHROW;
extern int omp_get_max_active_levels (void) __GOMP_NOTHROW;
+extern int omp_get_supported_active_team_dims (int, int) __GOMP_NOTHROW;
+extern int omp_get_supported_active_league_dims (int, int) __GOMP_NOTHROW;
+extern int omp_get_supported_teams_dim (int, int) __GOMP_NOTHROW;
+extern int omp_get_supported_threads_dim (int, int) __GOMP_NOTHROW;
extern int omp_get_supported_active_levels (void) __GOMP_NOTHROW;
extern int omp_get_level (void) __GOMP_NOTHROW;
extern int omp_get_ancestor_thread_num (int) __GOMP_NOTHROW;
@@ -328,7 +335,9 @@ extern int omp_get_default_device (void) __GOMP_NOTHROW;
extern int omp_get_num_devices (void) __GOMP_NOTHROW;
extern int omp_get_device_num (void) __GOMP_NOTHROW;
extern int omp_get_num_teams (void) __GOMP_NOTHROW;
+extern int omp_get_num_teams_dim (int) __GOMP_NOTHROW;
extern int omp_get_team_num (void) __GOMP_NOTHROW;
+extern int omp_get_team_num_dim (int) __GOMP_NOTHROW;
extern int omp_is_initial_device (void) __GOMP_NOTHROW;
extern int omp_get_initial_device (void) __GOMP_NOTHROW;
@@ -449,7 +458,7 @@ extern const char *omp_get_interop_rc_desc (const omp_interop_t,
omp_interop_rc_t) __GOMP_NOTHROW;
extern int omp_get_device_from_uid (const char *) __GOMP_NOTHROW;
-extern void omp_get_device_distances(int, const int *, int *) __GOMP_NOTHROW;
+extern void omp_get_device_distances (int, const int *, int *) __GOMP_NOTHROW;
extern const char *omp_get_uid_from_device (int) __GOMP_NOTHROW;
extern omp_control_tool_result_t omp_control_tool (omp_control_tool_t, int,
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index fc58b90f2d1..740b5a88641 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -388,12 +388,30 @@
end function omp_get_num_threads
end interface
+ interface
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_num_threads_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_num_threads_dim
+ integer (c_int), value :: dim
+ end function omp_get_num_threads_dim
+ end interface
+
interface
function omp_get_thread_num ()
integer (4) :: omp_get_thread_num
end function omp_get_thread_num
end interface
+ interface
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_thread_num_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_thread_num_dim
+ integer (c_int), value :: dim
+ end function omp_get_thread_num_dim
+ end interface
+
interface
function omp_test_nest_lock (nvar)
use omp_lib_kinds
@@ -446,6 +464,15 @@
end function omp_get_thread_limit
end interface
+ interface
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_thread_limit_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_thread_limit_dim
+ integer (c_int), value :: dim
+ end function omp_get_thread_limit_dim
+ end interface
+
interface omp_set_max_active_levels
subroutine omp_set_max_active_levels (max_levels)
integer (4), intent (in) :: max_levels
@@ -605,12 +632,30 @@
end function omp_get_num_teams
end interface
+ interface
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_num_teams_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_num_teams_dim
+ integer (c_int), value :: dim
+ end function omp_get_num_teams_dim
+ end interface
+
interface
function omp_get_team_num ()
integer (4) :: omp_get_team_num
end function omp_get_team_num
end interface
+ interface
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_team_num_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_team_num_dim
+ integer (c_int), value :: dim
+ end function omp_get_team_num_dim
+ end interface
+
interface
function omp_is_initial_device ()
logical (4) :: omp_is_initial_device
@@ -1056,6 +1101,45 @@
end function omp_get_device_from_uid
end interface
+ interface omp_get_supported_active_team_dims
+ ! Implementation choice: no default-kind-8 integer
+ function omp_get_supported_active_team_dims (device_num, &
+ native_support)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (4) :: omp_get_supported_active_team_dims
+ integer (4), intent(in) :: device_num
+ logical (4), intent(in) :: native_support
+ end function omp_get_supported_active_team_dims
+ end interface
+
+ interface omp_get_supported_active_league_dims
+ ! Implementation choice: no default-kind-8 integer
+ function omp_get_supported_active_league_dims ( &
+ device_num, native_support)
+ integer (4) :: omp_get_supported_active_league_dims
+ integer (4), intent(in) :: device_num
+ logical (4), intent(in) :: native_support
+ end function omp_get_supported_active_league_dims
+ end interface
+
+ interface omp_get_supported_teams_dim
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_supported_teams_dim (device_num, dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_supported_teams_dim
+ integer (c_int), value :: device_num, dim
+ end function omp_get_supported_teams_dim
+ end interface
+
+ interface omp_get_supported_threads_dim
+ ! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_supported_threads_dim (device_num, dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_supported_threads_dim
+ integer (c_int), value :: device_num, dim
+ end function omp_get_supported_threads_dim
+ end interface
+
interface omp_get_uid_from_device
! Deviation from OpenMP 6.0: VALUE added.
character(:) function omp_get_uid_from_device (device_num)
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index d5e39e6c770..cce424eaa35 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -337,6 +337,51 @@
integer(4) omp_get_default_device, omp_get_num_devices
integer(4) omp_get_num_teams, omp_get_team_num
+ interface
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_num_threads_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_num_threads_dim
+ integer (c_int), value :: dim
+ end function omp_get_num_threads_dim
+ end interface
+
+ interface
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_thread_num_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_thread_num_dim
+ integer (c_int), value :: dim
+ end function omp_get_thread_num_dim
+ end interface
+
+ interface
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_thread_limit_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_thread_limit_dim
+ integer (c_int), value :: dim
+ end function omp_get_thread_limit_dim
+ end interface
+
+ interface
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_num_teams_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_num_teams_dim
+ integer (c_int), value :: dim
+ end function omp_get_num_teams_dim
+ end interface
+
+ interface
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_team_num_dim (dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_team_num_dim
+ integer (c_int), value :: dim
+ end function omp_get_team_num_dim
+ end interface
+
external omp_is_initial_device
logical(4) omp_is_initial_device
external omp_get_initial_device
@@ -687,6 +732,46 @@
end function omp_get_uid_from_device_8
end interface omp_get_uid_from_device
+ interface omp_get_supported_active_team_dims
+! Implementation choice: no default-kind-8 integer
+ function omp_get_supported_active_team_dims (device_num, &
+ & native_support)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (4) :: omp_get_supported_active_team_dims
+ integer (4), intent(in) :: device_num
+ logical (4), intent(in) :: native_support
+ end function omp_get_supported_active_team_dims
+ end interface
+
+ interface omp_get_supported_active_league_dims
+! Implementation choice: no default-kind-8 integer
+ function omp_get_supported_active_league_dims ( &
+ & device_num, native_support)
+ integer (4) :: omp_get_supported_active_league_dims
+ integer (4), intent(in) :: device_num
+ logical (4), intent(in) :: native_support
+ end function omp_get_supported_active_league_dims
+ end interface
+
+ interface omp_get_supported_teams_dim
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_supported_teams_dim (device_num, dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_supported_teams_dim
+ integer (c_int), value :: device_num, dim
+ end function omp_get_supported_teams_dim
+ end interface
+
+ interface omp_get_supported_threads_dim
+! Implementation choice: bind(C) and no default-kind-8 integer
+ function omp_get_supported_threads_dim (device_num, dim) bind(C)
+ use, intrinsic :: iso_c_binding, only : c_int
+ integer (c_int) :: omp_get_supported_threads_dim
+ integer (c_int), value :: device_num, dim
+ end function omp_get_supported_threads_dim
+ end interface
+
+
interface
! This function only supports integer(4) for simplicity.
! Contrary to OpenMP 6.1, it also uses bind(C) and value
diff --git a/libgomp/parallel.c b/libgomp/parallel.c
index 754745227e0..fe7939ddcce 100644
--- a/libgomp/parallel.c
+++ b/libgomp/parallel.c
@@ -280,12 +280,28 @@ omp_get_num_threads (void)
return team ? team->nthreads : 1;
}
+int
+omp_get_num_threads_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_num_threads ();
+ return 1;
+}
+
int
omp_get_thread_num (void)
{
return gomp_thread ()->ts.team_id;
}
+int
+omp_get_thread_num_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_thread_num ();
+ return 0;
+}
+
/* This wasn't right for OpenMP 2.5. Active region used to be non-zero
when the IF clause doesn't evaluate to false, starting with OpenMP 3.0
it is non-zero with more than one thread in the team. */
@@ -334,7 +350,9 @@ omp_get_active_level (void)
}
ialias (omp_get_num_threads)
+ialias (omp_get_num_threads_dim)
ialias (omp_get_thread_num)
+ialias (omp_get_thread_num_dim)
ialias (omp_in_parallel)
ialias (omp_get_level)
ialias (omp_get_ancestor_thread_num)
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index cfeb648aa4a..81426168da9 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -1264,6 +1264,7 @@ limit_worker_threads (int threads)
/* FIXME Do something more intelligent here.
GCN can always run 4 threads within a Compute Unit, but
more than that depends on register usage. */
+ /* Keep in sync with GOMP_OFFLOAD_supported_threads_dim. */
if (threads > 16)
threads = 16;
return threads;
@@ -3773,6 +3774,31 @@ GOMP_OFFLOAD_get_numa_node (int ord)
return numa_node;
}
+/* Number of teams supported (by dimension) as reported by OpenMP.
+ For dim < 0 (invalid) and for dim > supported dims, return 1. */
+
+int
+GOMP_OFFLOAD_supported_teams_dim (int ord, int dim)
+{
+ if (dim > 0 /* max supported dims */ || dim < 0)
+ return 1;
+ struct agent_info *agent = get_agent_info (ord);
+ return limit_teams (__INT_MAX__, agent);
+}
+
+/* Number of threads supported (by dimension) as reported by OpenMP.
+ For dim < 0 (invalid) and for dim > supported dims, return 1. */
+
+int
+GOMP_OFFLOAD_supported_threads_dim (int ord, int dim)
+{
+ (void) ord;
+ if (dim > 0 /* max supported dims */ || dim < 0)
+ return 1;
+
+ return limit_worker_threads (__INT_MAX__);
+}
+
/* Return the specific capabilities the HSA accelerator have. */
unsigned int
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 447457cb448..59f2ce07f0e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1383,6 +1383,49 @@ GOMP_OFFLOAD_get_numa_node (int ord)
return numa_node;
}
+/* Number of teams supported (by dimension) as reported by OpenMP.
+ For dim < 0 (invalid) and for dim > supported dims, return 1. */
+
+int
+GOMP_OFFLOAD_supported_teams_dim (int ord, int dim)
+{
+ if (dim > 0 /* max supported dims */ || dim < 0)
+ return 1;
+
+ struct ptx_device *ptx_dev = ptx_devices[ord];
+
+ /* Keep in sync with nvptx_adjust_launch_bounds; assume 1 for the
+ following as upper bound. */
+ int num_threads = 1, regs_per_thread = 1;
+
+ int regs_per_block = regs_per_thread * 32 * num_threads;
+
+ int max_blocks = ptx_dev->regs_per_sm / regs_per_block * ptx_dev->num_sms;
+ /* This is an estimate of how many blocks the device can host simultaneously.
+ Actual limit, which may be lower, can be queried with "occupancy control"
+ driver interface (since CUDA 6.0). */
+ return max_blocks;
+}
+
+/* Number of threads supported (by dimension) as reported by OpenMP.
+ For dim < 0 (invalid) and for dim > supported dims, return 1. */
+
+int
+GOMP_OFFLOAD_supported_threads_dim (int ord, int dim)
+{
+ if (dim > 0 /* max supported dims */ || dim < 0)
+ return 1;
+
+ /* Keep in sync with nvptx_adjust_launch_bounds. */
+ struct ptx_device *ptx_dev = ptx_devices[ord];
+
+ int max_warps_block = ptx_dev->max_threads_per_block / 32;
+ /* Maximum 32 warps per block is an implementation limit in NVPTX backend
+ and libgcc, which matches documented limit of all GPUs as of 2015. */
+
+ return max_warps_block;
+}
+
unsigned int
GOMP_OFFLOAD_get_caps (void)
{
@@ -2644,7 +2687,9 @@ GOMP_OFFLOAD_openacc_get_property (int n, enum goacc_property prop)
/* Adjust launch dimensions: pick good values for number of blocks and warps
and ensure that number of warps does not exceed CUDA limits as well as GCC's
- own limits. */
+ own limits.
+ Keep in sync with GOMP_OFFLOAD_supported_teams_dims and
+ GOMP_OFFLOAD_supported_threads_dim. */
static void
nvptx_adjust_launch_bounds (struct targ_fn_descriptor *fn,
diff --git a/libgomp/target.c b/libgomp/target.c
index d07adc2956e..47fdd7458ea 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -6032,6 +6032,54 @@ GOMP_interop (int device_num, int n_init, struct interop_obj_t ***init,
}
}
+/* Query device-related data. */
+
+int
+omp_get_supported_active_team_dims (int device_num, int native_support)
+{
+ (void) device_num;
+ (void) native_support;
+ return 1;
+}
+
+int
+omp_get_supported_active_league_dims (int device_num, int native_support)
+{
+ (void) device_num;
+ (void) native_support;
+ return 1;
+}
+
+int
+omp_get_supported_teams_dim (int device_num, int dim)
+{
+ if (device_num == omp_default_device)
+ device_num = gomp_get_default_device ();
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
+ return dim == 0 ? INT_MAX : 1;
+
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
+ if (devicep == NULL || !devicep->supported_teams_dim_func)
+ return -1;
+ return devicep->supported_teams_dim_func (devicep->target_id, dim);
+}
+
+int
+omp_get_supported_threads_dim (int device_num, int dim)
+{
+ if (device_num == omp_default_device)
+ device_num = gomp_get_default_device ();
+ if (device_num == omp_initial_device
+ || device_num == gomp_get_num_devices ())
+ return dim == 0 ? INT_MAX : 1;
+
+ struct gomp_device_descr *devicep = resolve_device (device_num, false);
+ if (devicep == NULL || !devicep->supported_threads_dim_func)
+ return -1;
+ return devicep->supported_threads_dim_func (devicep->target_id, dim);
+}
+
static const char *
gomp_get_uid_for_device (struct gomp_device_descr *devicep, int device_num)
{
@@ -6136,6 +6184,10 @@ omp_get_device_distances (int ndevs, const int *devs, int *distances)
}
}
+ialias (omp_get_supported_active_team_dims)
+ialias (omp_get_supported_active_league_dims)
+ialias (omp_get_supported_threads_dim)
+ialias (omp_get_supported_teams_dim)
ialias (omp_get_uid_from_device)
ialias (omp_get_device_from_uid)
ialias (omp_get_device_distances)
@@ -6184,6 +6236,8 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
DLSYM (get_name);
DLSYM_OPT (get_uid, get_uid);
DLSYM_OPT (get_numa_node, get_numa_node);
+ DLSYM_OPT (supported_threads_dim, supported_threads_dim);
+ DLSYM_OPT (supported_teams_dim, supported_teams_dim);
DLSYM (get_caps);
DLSYM (get_type);
DLSYM (get_num_devices);
diff --git a/libgomp/teams.c b/libgomp/teams.c
index 1dad41278e3..bc80eb4f2ca 100644
--- a/libgomp/teams.c
+++ b/libgomp/teams.c
@@ -65,6 +65,14 @@ omp_get_num_teams (void)
return thr->num_teams + 1;
}
+int
+omp_get_num_teams_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_num_teams ();
+ return 1;
+}
+
int
omp_get_team_num (void)
{
@@ -72,5 +80,15 @@ omp_get_team_num (void)
return thr->team_num;
}
+int
+omp_get_team_num_dim (int dim)
+{
+ if (dim == 0)
+ return omp_get_team_num ();
+ return 0;
+}
+
ialias (omp_get_num_teams)
+ialias (omp_get_num_teams_dim)
ialias (omp_get_team_num)
+ialias (omp_get_team_num_dim)
diff --git a/libgomp/testsuite/libgomp.c/multi-dim-routines.c b/libgomp/testsuite/libgomp.c/multi-dim-routines.c
new file mode 100644
index 00000000000..edd3154628a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/multi-dim-routines.c
@@ -0,0 +1,135 @@
+#include <omp.h>
+
+int
+main ()
+{
+ /* Supported active teams/league dimension - the assumed result for either,
+ once implemented, is:
+ - Host: native == 1, but supporting (emulating) any dimension
+ - Non-host: native == 3 and not supporting higher dimension,
+ i.e. expected that no one implements this. */
+
+ /* Always true: All devices support at least one dimension,
+ with native_support set to true or to false. */
+
+ if (omp_get_supported_active_team_dims (omp_initial_device, true) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_team_dims (omp_initial_device, false) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (omp_initial_device, true) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (omp_initial_device, false) < 1)
+ __builtin_abort ();
+ for (int dev = 0; dev < omp_get_num_devices (); dev++)
+ {
+ if (omp_get_supported_active_team_dims (dev, true) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_team_dims (dev, false) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (dev, true) < 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (dev, false) < 1)
+ __builtin_abort ();
+ }
+
+ /* For each supported dimension, expect that it supports more than
+ one team and thread. - And for each unsupported dimension, expect
+ that exactly one thread and team is reported as supported. */
+
+ for (int dev = 0; dev < omp_get_num_devices (); dev++)
+ {
+ int max_dims = omp_get_supported_active_team_dims (dev, false);
+ int test_size = max_dims == __INT_MAX__ ? 64 : max_dims;
+ for (int dim = 1; dim <= test_size; dim++)
+ if (omp_get_supported_teams_dim (dev, dim) < 1)
+ __builtin_abort ();
+ int test_end = __INT_MAX__ - 12 > test_size ? test_size + 12 : __INT_MAX__;
+ if (max_dims < __INT_MAX__)
+ for (int dim = test_size + 1; dim < test_end; dim++)
+ if (omp_get_supported_teams_dim (dev, dim) != 1)
+ __builtin_abort ();
+
+ max_dims = omp_get_supported_active_league_dims (dev, false);
+ test_size = max_dims == __INT_MAX__ ? 64 : max_dims;
+ for (int dim = 1; dim <= test_size; dim++)
+ if (omp_get_supported_threads_dim (dev, dim) < 1)
+ __builtin_abort ();
+ test_end = __INT_MAX__ - 12 > test_size ? test_size + 12 : __INT_MAX__;
+ if (max_dims < __INT_MAX__)
+ for (int dim = test_size + 1; dim < test_end; dim++)
+ if (omp_get_supported_threads_dim (dev, dim) != 1)
+ __builtin_abort ();
+ if (max_dims == __INT_MAX__)
+ for (int dim = __INT_MAX__ - 12; dim < __INT_MAX__; dim++)
+ if (omp_get_supported_threads_dim (dev, dim) != 1)
+ __builtin_abort ();
+ }
+
+ /* Implementation specific to GCC:
+ Currently, exactly one dimension is supported on the host and
+ on the device, both natively and emulated. */
+
+ if (omp_get_supported_active_team_dims (omp_initial_device, true) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_team_dims (omp_initial_device, false) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (omp_initial_device, true) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (omp_initial_device, false) != 1)
+ __builtin_abort ();
+ for (int dev = 0; dev < omp_get_num_devices (); dev++)
+ {
+ if (omp_get_supported_active_team_dims (dev, true) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_team_dims (dev, false) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (dev, true) != 1)
+ __builtin_abort ();
+ if (omp_get_supported_active_league_dims (dev, false) != 1)
+ __builtin_abort ();
+ }
+
+ /* TODO: Once implemented, add a check for 'dim(N)', N > 1 below. */
+
+
+ /* The following should always pass as only a single spacial dimension
+ is used. */
+
+ for (int dev = omp_initial_device; dev < omp_get_num_devices (); dev++)
+ {
+ #pragma omp target teams device(dev) num_teams(10) thread_limit(8)
+ {
+ if (omp_get_num_teams () != 10
+ || omp_get_num_teams_dim (0) != 10)
+ __builtin_abort ();
+ if (omp_get_team_num () != omp_get_team_num_dim (0)
+ || omp_get_team_num_dim (0) < 0
+ || omp_get_team_num_dim (0) >= 10)
+ __builtin_abort ();
+ for (int dim = 1; dim <= 12; dim++)
+ if (omp_get_num_teams_dim (dim) != 1
+ || omp_get_team_num_dim (dim) != 0)
+ __builtin_abort ();
+ #pragma omp parallel num_threads(8)
+ // FIXME: Assumes 'num_threads(strict:', add once modifier is supported
+ {
+ if (omp_get_thread_limit_dim (0) != 8)
+ __builtin_abort ();
+ for (int dim = 1; dim <= 12; dim++)
+ if (omp_get_thread_limit_dim (dim) != 1)
+ __builtin_abort ();
+ if (omp_get_num_threads () != 8
+ || omp_get_num_threads_dim (0) != 8)
+ __builtin_abort ();
+ if (omp_get_team_num () != omp_get_team_num_dim (0)
+ || omp_get_thread_num_dim (0) < 0
+ || omp_get_thread_num_dim (0) >= 8)
+ __builtin_abort ();
+ for (int dim = 1; dim <= 12; dim++)
+ if (omp_get_num_threads_dim (dim) != 1
+ || omp_get_thread_num_dim (dim) != 0)
+ __builtin_abort ();
+ }
+ }
+ }
+}
diff --git a/libgomp/testsuite/libgomp.fortran/multi-dim-routines-2.F90 b/libgomp/testsuite/libgomp.fortran/multi-dim-routines-2.F90
new file mode 100644
index 00000000000..11ef17e29a7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/multi-dim-routines-2.F90
@@ -0,0 +1,2 @@
+#define USE_OMP_HEADER 1
+#include "multi-dim-routines.F90"
diff --git a/libgomp/testsuite/libgomp.fortran/multi-dim-routines.F90 b/libgomp/testsuite/libgomp.fortran/multi-dim-routines.F90
new file mode 100644
index 00000000000..ed11d01107d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/multi-dim-routines.F90
@@ -0,0 +1,149 @@
+program main
+#ifndef USE_OMP_HEADER
+ use omp_lib
+#endif
+ implicit none (type, external)
+
+#ifdef USE_OMP_HEADER
+ include "omp_lib.h"
+#endif
+
+ integer :: dev, dim
+
+ ! Supported active teams/league dimension - the assumed result for either,
+ ! once implemented, is:
+ ! - Host: native == 1, but supporting (emulating) any dimension
+ ! - Non-host: native == 3 and not supporting higher dimension,
+ ! i.e. expected that no one implements this.
+
+ ! Always true: All devices support at least one dimension,
+ ! with native_support set to true or to false.
+
+ if (omp_get_supported_active_team_dims (omp_initial_device, .true.) < 1) &
+ stop 1
+ if (omp_get_supported_active_team_dims (omp_initial_device, .false.) < 1) &
+ stop 2
+ if (omp_get_supported_active_league_dims (omp_initial_device, .true.) < 1) &
+ stop 3
+ if (omp_get_supported_active_league_dims (omp_initial_device, .false.) < 1) &
+ stop 4
+ do dev = 0, omp_get_num_devices () - 1
+ if (omp_get_supported_active_team_dims (dev, .true.) < 1) &
+ stop 5
+ if (omp_get_supported_active_team_dims (dev, .false.) < 1) &
+ stop 6
+ if (omp_get_supported_active_league_dims (dev, .true.) < 1) &
+ stop 7
+ if (omp_get_supported_active_league_dims (dev, .false.) < 1) &
+ stop 8
+ end do
+
+ ! For each supported dimension, expect that it supports more than
+ ! one team and thread. - And for each unsupported dimension, expect
+ ! that exactly one thread and team is reported as supported.
+
+ do dev = 0, omp_get_num_devices () - 1
+ block
+ integer :: max_dims, test_size, test_end
+ max_dims = omp_get_supported_active_team_dims (dev, .false.)
+ test_size = (max_dims == huge(max_dims) ? 64 : max_dims)
+ do dim = 1, test_size
+ if (omp_get_supported_teams_dim (dev, dim) < 1) &
+ stop 9
+ end do
+ test_end = (huge (test_size) - 12 > test_size ? test_size + 12 : huge(test_size))
+ if (max_dims < huge(max_dims)) then
+ do dim = test_size + 1, test_end - 1
+ if (omp_get_supported_teams_dim (dev, dim) /= 1) &
+ stop 10
+ end do
+ end if
+
+ max_dims = omp_get_supported_active_league_dims (dev, .false.)
+ test_size = (max_dims == huge(max_dims) ? 64 : max_dims)
+ do dim = 1, test_size
+ if (omp_get_supported_threads_dim (dev, dim) < 1) &
+ stop 11
+ end do
+ test_end = (huge(test_size) - 12 > test_size ? test_size + 12 : huge(test_size))
+ if (max_dims < huge(max_dims)) then
+ do dim = test_size + 1, test_end - 1
+ if (omp_get_supported_threads_dim (dev, dim) /= 1) &
+ stop 12
+ end do
+ else
+ do dim = huge (max_dims) - 12, huge (max_dims) - 1
+ if (omp_get_supported_threads_dim (dev, dim) /= 1) &
+ stop 13
+ end do
+ end if
+ end block
+ end do
+
+ ! Implementation specific to GCC:
+ ! Currently, exactly one dimension is supported on the host and
+ ! on the device, both natively and emulated.
+
+ if (omp_get_supported_active_team_dims (omp_initial_device, .true.) /= 1) &
+ stop 14
+ if (omp_get_supported_active_team_dims (omp_initial_device, .false.) /= 1) &
+ stop 15
+ if (omp_get_supported_active_league_dims (omp_initial_device, .true.) /= 1) &
+ stop 16
+ if (omp_get_supported_active_league_dims (omp_initial_device, .false.) /= 1) &
+ stop 17
+ do dev = 0, omp_get_num_devices () - 1
+ if (omp_get_supported_active_team_dims (dev, .true.) /= 1) &
+ stop 18
+ if (omp_get_supported_active_team_dims (dev, .false.) /= 1) &
+ stop 19
+ if (omp_get_supported_active_league_dims (dev, .true.) /= 1) &
+ stop 20
+ if (omp_get_supported_active_league_dims (dev, .false.) /= 1) &
+ stop 21
+ end do
+
+ ! TODO: Once implemented, add a check for 'dim(N)', N > 1 below.
+
+
+ ! The following should always pass as only a single spacial dimension
+ ! is used.
+
+ do dev = omp_initial_device, omp_get_num_devices () - 1
+ !$omp target teams device(dev) num_teams(10) thread_limit(8)
+ if (omp_get_num_teams () /= 10 &
+ .or. omp_get_num_teams_dim (0) /= 10) &
+ stop 22
+ if (omp_get_team_num () /= omp_get_team_num_dim (0) &
+ .or. omp_get_team_num_dim (0) < 0 &
+ .or. omp_get_team_num_dim (0) >= 10) &
+ stop 23
+ do dim = 1, 12
+ if (omp_get_num_teams_dim (dim) /= 1 &
+ .or. omp_get_team_num_dim (dim) /= 0) &
+ stop 24
+ end do
+ !$omp parallel num_threads(8)
+ ! FIXME: Assumes 'num_threads(strict:', add once modifier is supported
+ if (omp_get_thread_limit_dim (0) /= 8) &
+ stop 25
+ do dim = 1, 12
+ if (omp_get_thread_limit_dim (dim) /= 1) &
+ stop 26
+ end do
+ if (omp_get_num_threads () /= 8 &
+ .or. omp_get_num_threads_dim (0) /= 8) &
+ stop 27
+ if (omp_get_team_num () /= omp_get_team_num_dim (0) &
+ .or. omp_get_thread_num_dim (0) < 0 &
+ .or. omp_get_thread_num_dim (0) >= 8) &
+ stop 28
+ do dim = 1, 12
+ if (omp_get_num_threads_dim (dim) /= 1 &
+ .or. omp_get_thread_num_dim (dim) /= 0) &
+ stop 29
+ end do
+ !$omp end parallel
+ !$omp end target teams
+ end do
+end program
