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=53632d9a514707e40bb4ddc891635b32c52f7dae;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. (cherry picked from openacc-gcc-9-branch commit 99a4260fe264624136eda666363464eba48bdc3e) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 3c516287c92e..f23ad4ed2202 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.c (nvptx_propagate_unified): New. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index a17e54aa7935..ca37ddfb66e4 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -668,7 +668,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; @@ -699,7 +700,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 & 3) == 1 && omp_is_reference (var)) type = TREE_TYPE (type); @@ -1155,10 +1160,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; @@ -1505,7 +1512,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) && !OMP_CLAUSE_MAP_IN_REDUCTION (c)) install_var_local (decl, ctx); @@ -2612,6 +2620,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 dynamic array map clauses are placed at the very front of the chain. */ @@ -2664,6 +2734,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) omp_context *ctx; tree name; bool offloaded = is_gimple_omp_offloaded (stmt); + tree clauses = gimple_omp_target_clauses (stmt); ctx = new_omp_context (stmt, outer_ctx); ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0); @@ -2676,21 +2747,28 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; + 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"); } /* If OpenACC construct, put dynamic array clauses (if any) in front of clause chain. The runtime can then test the first to see if the additional map processing for them is required. */ if (is_gimple_omp_oacc (stmt)) - reorder_dynamic_array_clauses (gimple_omp_target_clauses_ptr (stmt)); + { + reorder_dynamic_array_clauses (gimple_omp_target_clauses_ptr (stmt)); + clauses = gimple_omp_target_clauses (stmt); + } - tree clauses = gimple_omp_target_clauses (stmt); - - 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)