public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] Reinstate kernels-restrict behaviour
@ 2022-06-29 14:34 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:34 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:90d4c72130b639d27509d749a2c272b5315dd5c2
commit 90d4c72130b639d27509d749a2c272b5315dd5c2
Author: Julian Brown <julian@codesourcery.com>
Date: Wed Mar 6 14:44:56 2019 -0800
Reinstate kernels-restrict behaviour
This patch contains a small fix for upstream churn relative to the last version
posted.
2018-09-20 Cesar Philippidis <cesar@codesourcery.com>
Julian Brown <julian@codesourcery.com>
* omp-low.c (install_var_field): New base_pointer_restrict
argument.
(scan_sharing_clauses): Update call to install_var_field.
(omp_target_base_pointers_restrict_p): New function.
(scan_omp_target): Update call to install_var_field.
Diff:
---
gcc/ChangeLog.omp | 9 ++++++
gcc/omp-low.cc | 88 +++++++++++++++++++++++++++++++++++++++++++++++++++----
2 files changed, 91 insertions(+), 6 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d1a9bcef523..6c0729986fa 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2018-09-20 Cesar Philippidis <cesar@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * omp-low.c (install_var_field): New base_pointer_restrict
+ argument.
+ (scan_sharing_clauses): Update call to install_var_field.
+ (omp_target_base_pointers_restrict_p): New function.
+ (scan_omp_target): Update call to install_var_field.
+
2018-10-30 Cesar Philippidis <cesar@codesourcery.com>
* config/nvptx/nvptx.cc (nvptx_propagate_unified): New.
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 70a6dc3c399..edc22e25e91 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -777,7 +777,8 @@ build_sender_ref (tree var, omp_context *ctx)
BASE_POINTERS_RESTRICT, declare the field with restrict. */
static void
-install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
+ bool base_pointers_restrict = false)
{
tree field, type, sfield = NULL_TREE;
splay_tree_key key = (splay_tree_key) var;
@@ -816,7 +817,11 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
type = build_pointer_type (build_pointer_type (type));
}
else if (by_ref)
- type = build_pointer_type (type);
+ {
+ type = build_pointer_type (type);
+ if (base_pointers_restrict)
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+ }
else if ((mask & (32 | 3)) == 1
&& omp_privatize_by_reference (var))
type = TREE_TYPE (type);
@@ -1256,10 +1261,12 @@ fixup_child_record_type (omp_context *ctx)
}
/* Instantiate decls as necessary in CTX to satisfy the data sharing
- specified by CLAUSES. */
+ specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with
+ restrict. */
static void
-scan_sharing_clauses (tree clauses, omp_context *ctx)
+scan_sharing_clauses (tree clauses, omp_context *ctx,
+ bool base_pointers_restrict = false)
{
tree c, decl;
bool scan_array_reductions = false;
@@ -1796,7 +1803,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
else
- install_var_field (decl, true, 3, ctx);
+ install_var_field (decl, true, 3, ctx,
+ base_pointers_restrict);
if (is_gimple_omp_offloaded (ctx->stmt)
&& !(is_gimple_omp_oacc (ctx->stmt)
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
@@ -3197,6 +3205,68 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
+/* Return true if the CLAUSES of an omp target guarantee that the base pointers
+ used in the corresponding offloaded function are restrict. */
+
+static bool
+omp_target_base_pointers_restrict_p (tree clauses)
+{
+ /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only
+ used by OpenACC. */
+ if (flag_openacc == 0)
+ return false;
+
+ /* I. Basic example:
+
+ void foo (void)
+ {
+ unsigned int a[2], b[2];
+
+ #pragma acc kernels \
+ copyout (a) \
+ copyout (b)
+ {
+ a[0] = 0;
+ b[0] = 1;
+ }
+ }
+
+ After gimplification, we have:
+
+ #pragma omp target oacc_kernels \
+ map(force_from:a [len: 8]) \
+ map(force_from:b [len: 8])
+ {
+ a[0] = 0;
+ b[0] = 1;
+ }
+
+ Because both mappings have the force prefix, we know that they will be
+ allocated when calling the corresponding offloaded function, which means we
+ can mark the base pointers for a and b in the offloaded function as
+ restrict. */
+
+ tree c;
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return false;
+
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ break;
+ default:
+ return false;
+ }
+ }
+
+ return true;
+}
+
/* Reorder clauses so that non-contiguous array map clauses are placed at the very
front of the chain. */
@@ -3269,13 +3339,19 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
tree clauses = gimple_omp_target_clauses (stmt);
+ bool base_pointers_restrict = false;
if (offloaded)
{
create_omp_child_function (ctx, false);
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+
+ base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses);
+ if (base_pointers_restrict && dump_file && (dump_flags & TDF_DETAILS))
+ fprintf (dump_file,
+ "Base pointers in offloaded function are restrict\n");
}
- scan_sharing_clauses (clauses, ctx);
+ scan_sharing_clauses (clauses, ctx, base_pointers_restrict);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:34 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:34 [gcc/devel/omp/gcc-12] Reinstate kernels-restrict behaviour Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).