https://gcc.gnu.org/g:31ba6ead0d5083010698d9fc0d45c31a4ae1f61d
commit r17-1804-g31ba6ead0d5083010698d9fc0d45c31a4ae1f61d Author: Paul-Antoine Arras <[email protected]> Date: Wed Jun 17 16:23:00 2026 +0200 openmp: Add GOMP_loop_static_worksharing and GOMP_distribute_static_worksharing For static-schedule worksharing loops (`omp for schedule(static)') the compiler emits two separate calls -- omp_get_thread_num() and omp_get_num_threads() -- to obtain the thread id and count. Similarly, for `omp distribute' constructs, it emits omp_get_team_num() and omp_get_num_teams(). Because these are public OpenMP API routines, OMPT cannot distinguish them from user-level queries. Introduce two new libgomp entry points: - GOMP_loop_static_worksharing(): returns both the thread id and the thread count packed into a single complex value. - GOMP_distribute_static_worksharing(): same, but for team id and team count. Using a single return value rather than two output-pointer arguments preserves optimisations in later passes. gcc/ChangeLog: * builtin-types.def (BT_COMPLEX_INT, BT_FN_COMPLEX_INT): New types. * omp-builtins.def (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING): New builtin. (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING): Likewise. * omp-expand.cc (expand_omp_for_static_nochunk): Replace separate calls to public API functions with a single call to GOMP_loop_static_worksharing or GOMP_distribute_static_worksharing, then unpack id and count. (expand_omp_for_static_chunk): Likewise. gcc/fortran/ChangeLog: * types.def (BT_COMPLEX_INT): New type. (BT_FN_COMPLEX_INT): New function type. libgomp/ChangeLog: * config/gcn/teams.c (GOMP_distribute_static_worksharing): New function. * config/nvptx/teams.c (GOMP_distribute_static_worksharing): Likewise. * libgomp.map (GOMP_6.0.2): Export GOMP_loop_static_worksharing and GOMP_distribute_static_worksharing. * libgomp_g.h (GOMP_loop_static_worksharing): New declaration. (GOMP_distribute_static_worksharing): Likewise. * parallel.c (GOMP_loop_static_worksharing): New function. * teams.c (GOMP_distribute_static_worksharing): Likewise. gcc/testsuite/ChangeLog: * c-c++-common/gomp/for-8.c: New test. Diff: --- gcc/builtin-types.def | 3 ++ gcc/fortran/types.def | 3 ++ gcc/omp-builtins.def | 6 +++ gcc/omp-expand.cc | 66 +++++++++++++++++++++----------- gcc/testsuite/c-c++-common/gomp/for-8.c | 67 +++++++++++++++++++++++++++++++++ libgomp/config/gcn/teams.c | 12 ++++++ libgomp/config/nvptx/teams.c | 11 ++++++ libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 2 + libgomp/parallel.c | 12 ++++++ libgomp/teams.c | 12 ++++++ 11 files changed, 174 insertions(+), 22 deletions(-) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index 365badca2aa9..55ed6f0cae10 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -61,6 +61,7 @@ DEF_PRIMITIVE_TYPE (BT_VOID, void_type_node) DEF_PRIMITIVE_TYPE (BT_BOOL, boolean_type_node) DEF_PRIMITIVE_TYPE (BT_INT, integer_type_node) DEF_PRIMITIVE_TYPE (BT_UINT, unsigned_type_node) +DEF_PRIMITIVE_TYPE (BT_COMPLEX_INT, complex_integer_type_node) DEF_PRIMITIVE_TYPE (BT_LONG, long_integer_type_node) DEF_PRIMITIVE_TYPE (BT_ULONG, long_unsigned_type_node) DEF_PRIMITIVE_TYPE (BT_LONGLONG, long_long_integer_type_node) @@ -235,7 +236,9 @@ DEF_FUNCTION_TYPE_0 (BT_FN_CONST_STRING, BT_CONST_STRING) DEF_FUNCTION_TYPE_0 (BT_FN_PID, BT_PID) DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_INT) DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT) +DEF_FUNCTION_TYPE_0 (BT_FN_COMPLEX_INT, BT_COMPLEX_INT) DEF_FUNCTION_TYPE_0 (BT_FN_ULONG, BT_ULONG) +DEF_FUNCTION_TYPE_0 (BT_FN_ULONGLONG, BT_ULONGLONG) DEF_FUNCTION_TYPE_0 (BT_FN_FLOAT, BT_FLOAT) DEF_FUNCTION_TYPE_0 (BT_FN_DOUBLE, BT_DOUBLE) /* For "long double" we use LONGDOUBLE (not LONG_DOUBLE) to diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index 32c533bc04fc..6f98e3e48ffa 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -47,6 +47,7 @@ along with GCC; see the file COPYING3. If not see DEF_PRIMITIVE_TYPE (BT_VOID, void_type_node) DEF_PRIMITIVE_TYPE (BT_BOOL, boolean_type_node) DEF_PRIMITIVE_TYPE (BT_INT, integer_type_node) +DEF_PRIMITIVE_TYPE (BT_COMPLEX_INT, complex_integer_type_node) DEF_PRIMITIVE_TYPE (BT_UINT, unsigned_type_node) DEF_PRIMITIVE_TYPE (BT_LONG, long_integer_type_node) DEF_PRIMITIVE_TYPE (BT_ULONGLONG, long_long_unsigned_type_node) @@ -76,8 +77,10 @@ DEF_POINTER_TYPE (BT_PTR_PTR, BT_PTR) DEF_FUNCTION_TYPE_0 (BT_FN_BOOL, BT_BOOL) DEF_FUNCTION_TYPE_0 (BT_FN_PTR, BT_PTR) DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_INT) +DEF_FUNCTION_TYPE_0 (BT_FN_COMPLEX_INT, BT_COMPLEX_INT) DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT) DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID) +DEF_FUNCTION_TYPE_0 (BT_FN_ULONGLONG, BT_ULONGLONG) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR) diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index c7ddea9880ad..82736e773407 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -499,3 +499,9 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ERROR, "GOMP_error", DEF_GOMP_BUILTIN (BUILT_IN_GOMP_HAS_MASKED_THREAD_NUM, "GOMP_has_masked_thread_num", BT_FN_BOOL_INT, ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING, + "GOMP_loop_static_worksharing", BT_FN_COMPLEX_INT, + ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING, + "GOMP_distribute_static_worksharing", BT_FN_COMPLEX_INT, + ATTR_CONST_NOTHROW_LEAF_LIST) diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 4d5e81f3fb17..de8889a96496 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -5190,27 +5190,38 @@ expand_omp_for_static_nochunk (struct omp_region *region, release_ssa_name (gimple_assign_lhs (g)); } } + /* Fetch the thread/team id and the number of threads/teams in a single + call to GOMP_loop_static_worksharing or GOMP_distribute_static_worksharing. + The helper returns both values packed into one complex int, with + the id as the imaginary part and the count as the real part. Returning + (rather than writing through pointers) keeps both values as plain SSA + names, which lets later passes - notably IPA-CP propagating constants + into the outlined kernel - reason about them. */ + tree decl; switch (gimple_omp_for_kind (fd->for_stmt)) { case GF_OMP_FOR_KIND_FOR: - nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS); - threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM); + decl = builtin_decl_explicit (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING); break; case GF_OMP_FOR_KIND_DISTRIBUTE: - nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS); - threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM); + decl = builtin_decl_explicit (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING); break; default: gcc_unreachable (); } - nthreads = build_call_expr (nthreads, 0); - nthreads = fold_convert (itype, nthreads); - nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE, - true, GSI_SAME_STMT); - threadid = build_call_expr (threadid, 0); - threadid = fold_convert (itype, threadid); - threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, + { + tree packed = build_call_expr (decl, 0); + packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE, true, GSI_SAME_STMT); + threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed); + threadid = fold_convert (itype, threadid); + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, + true, GSI_SAME_STMT); + nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed); + nthreads = fold_convert (itype, nthreads); + nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE, + true, GSI_SAME_STMT); + } n1 = fd->loop.n1; n2 = fd->loop.n2; @@ -5944,27 +5955,38 @@ expand_omp_for_static_chunk (struct omp_region *region, release_ssa_name (gimple_assign_lhs (g)); } } + /* Fetch the thread/team id and the number of threads/teams in a single + call to GOMP_loop_static_worksharing or GOMP_distribute_static_worksharing. + The helper returns both values packed into one complex int, with + the id as the imaginary part and the count as the real part. Returning + (rather than writing through pointers) keeps both values as plain SSA + names, which lets later passes - notably IPA-CP propagating constants + into the outlined kernel - reason about them. */ + tree decl; switch (gimple_omp_for_kind (fd->for_stmt)) { case GF_OMP_FOR_KIND_FOR: - nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS); - threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM); + decl = builtin_decl_explicit (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING); break; case GF_OMP_FOR_KIND_DISTRIBUTE: - nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS); - threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM); + decl = builtin_decl_explicit (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING); break; default: gcc_unreachable (); } - nthreads = build_call_expr (nthreads, 0); - nthreads = fold_convert (itype, nthreads); - nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE, - true, GSI_SAME_STMT); - threadid = build_call_expr (threadid, 0); - threadid = fold_convert (itype, threadid); - threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, + { + tree packed = build_call_expr (decl, 0); + packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE, true, GSI_SAME_STMT); + threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed); + threadid = fold_convert (itype, threadid); + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, + true, GSI_SAME_STMT); + nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed); + nthreads = fold_convert (itype, nthreads); + nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE, + true, GSI_SAME_STMT); + } n1 = fd->loop.n1; n2 = fd->loop.n2; diff --git a/gcc/testsuite/c-c++-common/gomp/for-8.c b/gcc/testsuite/c-c++-common/gomp/for-8.c new file mode 100644 index 000000000000..a894aaaaed11 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/for-8.c @@ -0,0 +1,67 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-ompexp" } */ + +/* Check that the static loop and distribute expanders fetch the thread/team + id and count through a single GOMP_loop_static_worksharing / + GOMP_distribute_static_worksharing call rather than separate + omp_get_thread_num/omp_get_num_threads or omp_get_team_num/omp_get_num_teams + calls. */ + +void bar (int); + +/* Static schedule without a chunk size goes through + expand_omp_for_static_nochunk. */ + +void +f1 (int n) +{ + int i; + #pragma omp for schedule(static) + for (i = 0; i < n; ++i) + bar (i); +} + +/* Static schedule with a chunk size goes through + expand_omp_for_static_chunk. */ + +void +f2 (int n) +{ + int i; + #pragma omp for schedule(static, 4) + for (i = 0; i < n; ++i) + bar (i); +} + +/* Distribute without a chunk size goes through + expand_omp_for_static_nochunk. */ + +void +f3 (int n) +{ + int i; +#pragma omp teams +#pragma omp distribute + for (i = 0; i < n; ++i) + bar (i); +} + +/* Distribute with a chunk size goes through + expand_omp_for_static_chunk. */ + +void +f4 (int n) +{ + int i; +#pragma omp teams +#pragma omp distribute dist_schedule(static, 4) + for (i = 0; i < n; ++i) + bar (i); +} + +/* { dg-final { scan-tree-dump "GOMP_loop_static_worksharing" "ompexp" } } */ +/* { dg-final { scan-tree-dump "GOMP_distribute_static_worksharing" "ompexp" } } */ +/* { dg-final { scan-tree-dump-not "omp_get_num_threads" "ompexp" } } */ +/* { dg-final { scan-tree-dump-not "omp_get_thread_num" "ompexp" } } */ +/* { dg-final { scan-tree-dump-not "omp_get_num_teams" "ompexp" } } */ +/* { dg-final { scan-tree-dump-not "omp_get_team_num" "ompexp" } } */ diff --git a/libgomp/config/gcn/teams.c b/libgomp/config/gcn/teams.c index 793d4bb78388..5aa8b15502aa 100644 --- a/libgomp/config/gcn/teams.c +++ b/libgomp/config/gcn/teams.c @@ -38,6 +38,18 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams, (void) thread_limit; } +/* For a distribute construct with static schedule, return the team ID and + number of teams packed into a single complexvalue. */ + +_Complex int +GOMP_distribute_static_worksharing (void) +{ + int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM; + unsigned tid = *gomp_team_num; + unsigned nteams = gomp_num_teams_var + 1; + return nteams + tid * 1I; +} + int omp_get_num_teams (void) { diff --git a/libgomp/config/nvptx/teams.c b/libgomp/config/nvptx/teams.c index af6bf6c57107..9763f6f41084 100644 --- a/libgomp/config/nvptx/teams.c +++ b/libgomp/config/nvptx/teams.c @@ -41,6 +41,17 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams, (void) thread_limit; } +/* For a distribute construct with static schedule, return the team ID and + number of teams packed into a single complexvalue. */ + +_Complex int +GOMP_distribute_static_worksharing (void) +{ + unsigned tid = __gomp_team_num; + unsigned nteams = gomp_num_teams_var + 1; + return nteams + tid * 1I; +} + int omp_get_num_teams (void) { diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 5faeb27b74b0..4aa518ecdae7 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -485,6 +485,8 @@ GOMP_6.0.1 { GOMP_6.0.2 { global: GOMP_has_masked_thread_num; + GOMP_loop_static_worksharing; + GOMP_distribute_static_worksharing; } GOMP_6.0.1; OACC_2.0 { diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index d36d9d360204..6dc1fd46fcea 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -294,6 +294,7 @@ extern unsigned GOMP_parallel_reductions (void (*) (void *), void *, unsigned, extern bool GOMP_cancel (int, bool); extern bool GOMP_cancellation_point (int); extern bool GOMP_has_masked_thread_num (int); +extern __complex__ int GOMP_loop_static_worksharing (void); /* task.c */ @@ -374,6 +375,7 @@ extern void *GOMP_target_map_indirect_ptr (void *); extern void GOMP_teams_reg (void (*) (void *), void *, unsigned, unsigned, unsigned); +extern __complex__ int GOMP_distribute_static_worksharing (void); /* allocator.c */ diff --git a/libgomp/parallel.c b/libgomp/parallel.c index 175554de4c9a..92472ecb4f0e 100644 --- a/libgomp/parallel.c +++ b/libgomp/parallel.c @@ -271,6 +271,18 @@ GOMP_cancel (int which, bool do_cancel) return true; } +/* For a worksharing-loop construct with static schedule, return the thread ID + and number of threads packed into a single complex value. */ + +_Complex int +GOMP_loop_static_worksharing (void) +{ + struct gomp_team *team = gomp_thread ()->ts.team; + unsigned tid = gomp_thread ()->ts.team_id; + unsigned nthreads = team ? team->nthreads : 1; + return nthreads + tid * 1I; +} + /* Return true if the current thread number equals TID. Used to implement the masked construct's filter clause. */ diff --git a/libgomp/teams.c b/libgomp/teams.c index bc80eb4f2ca3..7038623b2dfb 100644 --- a/libgomp/teams.c +++ b/libgomp/teams.c @@ -58,6 +58,18 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams, } } +/* For a distribute construct with static schedule, return the team ID and + number of teams packed into a single complex value. */ + +_Complex int +GOMP_distribute_static_worksharing (void) +{ + struct gomp_thread *thr = gomp_thread (); + unsigned tid = thr->team_num; + unsigned nteams = thr->num_teams + 1; + return nteams + tid * 1I; +} + int omp_get_num_teams (void) {
