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 --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 365badca2aa96b7f4736991a92fcb4b642d4ccc7..55ed6f0cae10a1ecc3dbba581c622b1404701221 100644 (file)
--- 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 32c533bc04fc0946f4de40181d06af8fef39e88b..6f98e3e48ffaad78d82b56d4071fb7ccbba13179 100644 (file)
--- 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 c7ddea9880adc8ccd5e106b5f2e8d21b7c78a861..82736e773407b6c9bdd7b01b38ee0bfc273bc811 100644 (file)
--- 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 4d5e81f3fb17f00538aae9b3814d4830612f728d..de8889a96496f016e6dca7db103469c8c5d544bc 100644 (file)
--- 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 (file)
index 0000000..a894aaa
--- /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 793d4bb783885aacb9027e70ac413d33353ee9b3..5aa8b15502aa56ede257ade6294d6079a76de9e3 100644 (file)
--- 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 af6bf6c57107edcba7651e8970a63fa5f1329d96..9763f6f41084dac1d23990e465c2c963fab68732 100644 (file)
--- 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 5faeb27b74b053eda56ebf36bff02adea49b926c..4aa518ecdae72e567a803d1a51bb61e951918bb4 100644 (file)
--- 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 d36d9d36020483fe17046b8578a0fcc9e2edffb5..6dc1fd46fcea531b084c82cfc0aeca243e8f51f5 100644 (file)
--- 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 175554de4c9a5f72ffe2664ef3be31eeb5110c3d..92472ecb4f0e56785ff07ecf62669b5a2d660269 100644 (file)
--- 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 bc80eb4f2ca31c6225208829ca9ee841e74de9df..7038623b2dfb4160ed973827dcd702275167a438 100644 (file)
--- 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)
 {