https://gcc.gnu.org/g:9bfd50a42d53b02860fe0f58c6be1d58c791e486
commit r17-1746-g9bfd50a42d53b02860fe0f58c6be1d58c791e486 Author: Tobias Burnus <[email protected]> Date: Mon Jun 22 13:35:41 2026 +0200 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. Diff: --- 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 | 393 ++++++++++++++++++++- 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, 1112 insertions(+), 12 deletions(-) diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index e4dcf063a5b9..d0307aeba88b 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 02b62d54ba60..eef0b418b9b4 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 42570eb445e2..793d4bb78388 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 4af9aaef18a5..af6bf6c57107 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 59d1de2886e9..a28b86ae666e 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 fc6ad65074b3..6c331451983a 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 88841ac5817b..60ac7201597a 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 f2aabdbc8068..9510b148ed8c 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 7e8bb7ebec3b..dbde88c4d5f5 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 9fcfbf96963f..05de63e4020c 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,376 @@ 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 dim)} +@end multitable + +@item @emph{Fortran}: +@multitable @columnfractions .20 .80 +@item @emph{Interface}: @tab @code{integer function omp_get_supported_threads_dim(device_num, dim)} +@item @tab @code{integer :: device_num, dim} +@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 +2445,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 cfb104dccbdd..605ea9bb5a6b 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 fc58b90f2d1f..740b5a886410 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 d5e39e6c770d..cce424eaa350 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 754745227e05..fe7939ddcce2 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 cfeb648aa4a2..81426168da95 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 447457cb448f..59f2ce07f0ed 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 d07adc2956e5..47fdd7458ea8 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 1dad41278e3e..bc80eb4f2ca3 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 000000000000..edd3154628ab --- /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 000000000000..11ef17e29a7f --- /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 000000000000..ed11d01107de --- /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
