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 an unsigned long long.
- 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_FN_ULONGLONG): New function type.
* 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_FN_ULONGLONG): 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.
* 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.
---
gcc/builtin-types.def | 1 +
gcc/fortran/types.def | 1 +
gcc/omp-builtins.def | 6 +++
gcc/omp-expand.cc | 72 +++++++++++++++++--------
gcc/testsuite/c-c++-common/gomp/for-8.c | 67 +++++++++++++++++++++++
libgomp/config/gcn/teams.c | 13 +++++
libgomp/config/nvptx/teams.c | 12 +++++
libgomp/libgomp.map | 2 +
libgomp/parallel.c | 13 +++++
libgomp/teams.c | 13 +++++
10 files changed, 178 insertions(+), 22 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/for-8.c
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 365badca2aa..6cdf1efb8bf 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -236,6 +236,7 @@ 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_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 32c533bc04f..8c18558dd49 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -78,6 +78,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_PTR, BT_PTR)
DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_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 c7ddea9880a..2578624be2d 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_ULONGLONG,
+ ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING,
+ "GOMP_distribute_static_worksharing", BT_FN_ULONGLONG,
+ ATTR_CONST_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index f7db4542b7d..ade4e0f5329 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -5190,27 +5190,41 @@ 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 unsigned long long, with
+ the id in the low 32 bits and the count in the high 32 bits. 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 utype = long_long_unsigned_type_node;
+ tree packed = build_call_expr (decl, 0);
+ packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
true, GSI_SAME_STMT);
+ threadid = fold_build2 (BIT_AND_EXPR, utype, packed,
+ build_int_cst (utype, 0xffffffffULL));
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
+ true, GSI_SAME_STMT);
+ nthreads = fold_build2 (RSHIFT_EXPR, utype, packed,
+ build_int_cst (integer_type_node, 32));
+ 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 +5958,41 @@ 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 unsigned long long, with
+ the id in the low 32 bits and the count in the high 32 bits. 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 utype = long_long_unsigned_type_node;
+ tree packed = build_call_expr (decl, 0);
+ packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
true, GSI_SAME_STMT);
+ threadid = fold_build2 (BIT_AND_EXPR, utype, packed,
+ build_int_cst (utype, 0xffffffffULL));
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
+ true, GSI_SAME_STMT);
+ nthreads = fold_build2 (RSHIFT_EXPR, utype, packed,
+ build_int_cst (integer_type_node, 32));
+ 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 00000000000..a894aaaaed1
--- /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 42570eb445e..3a937a47522 100644
--- a/libgomp/config/gcn/teams.c
+++ b/libgomp/config/gcn/teams.c
@@ -38,6 +38,19 @@ 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 value: the team ID in the low 32 bits
+ and the number of teams in the high 32 bits. */
+
+unsigned long long
+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 ((unsigned long long) nteams << 32) | tid;
+}
+
int
omp_get_num_teams (void)
{
diff --git a/libgomp/config/nvptx/teams.c b/libgomp/config/nvptx/teams.c
index 4af9aaef18a..b5b2a91c429 100644
--- a/libgomp/config/nvptx/teams.c
+++ b/libgomp/config/nvptx/teams.c
@@ -41,6 +41,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 value: the team ID in the low 32 bits
+ and the number of teams in the high 32 bits. */
+
+unsigned long long
+GOMP_distribute_static_worksharing (void)
+{
+ unsigned tid = __gomp_team_num;
+ unsigned nteams = gomp_num_teams_var + 1;
+ return ((unsigned long long) nteams << 32) | tid;
+}
+
int
omp_get_num_teams (void)
{
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 241fd45c4da..415d3883575 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -469,6 +469,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/parallel.c b/libgomp/parallel.c
index d45ea61ddaa..9eedcdb1eb4 100644
--- a/libgomp/parallel.c
+++ b/libgomp/parallel.c
@@ -271,6 +271,19 @@ 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 value: the thread ID in the low
+ 32 bits and the number of threads in the high 32 bits. */
+
+unsigned long long
+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 ((unsigned long long) nthreads << 32) | tid;
+}
+
/* 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 1dad41278e3..728b35df72e 100644
--- a/libgomp/teams.c
+++ b/libgomp/teams.c
@@ -58,6 +58,19 @@ 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 value: the team ID in the low 32 bits
+ and the number of teams in the high 32 bits. */
+
+unsigned long long
+GOMP_distribute_static_worksharing (void)
+{
+ struct gomp_thread *thr = gomp_thread ();
+ unsigned tid = thr->team_num;
+ unsigned nteams = thr->num_teams + 1;
+ return ((unsigned long long) nteams << 32) | tid;
+}
+
int
omp_get_num_teams (void)
{
--
2.53.0
