From: Julian Brown Date: Wed, 6 Mar 2019 22:44:56 +0000 (-0800) Subject: Reinstate kernels-restrict behaviour X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=90d4c72130b639d27509d749a2c272b5315dd5c2;p=thirdparty%2Fgcc.git Reinstate kernels-restrict behaviour This patch contains a small fix for upstream churn relative to the last version posted. 2018-09-20 Cesar Philippidis Julian Brown * 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 --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index d1a9bcef5235..6c0729986fa1 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,12 @@ +2018-09-20 Cesar Philippidis + Julian Brown + + * 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 * config/nvptx/nvptx.cc (nvptx_propagate_unified): New. diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 70a6dc3c3995..edc22e25e91d 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)