From: Chung-Lin Tang Date: Mon, 1 Feb 2021 11:16:47 +0000 (-0800) Subject: OpenMP 5.0: Allow multiple clauses mapping same variable X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=4b1205e579733ad9b777052df019cbb83ce696c3;p=thirdparty%2Fgcc.git OpenMP 5.0: Allow multiple clauses mapping same variable This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562081.html This patch now allows multiple clauses on the same construct to map the same variable, which was not valid in OpenMP 4.5, but now allowed in 5.0. This may possibly reverted/updated when a final patch is approved for mainline. 2021-02-01 Chung-Lin Tang gcc/cp/ChangeLog: * semantics.cc (finish_omp_clauses): Adjust to allow duplicate mapped variables for OpenMP. gcc/ChangeLog: * omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE' default parameter. Set splay-tree lookup key to key_expr instead of var if key_expr is non-NULL. Adjust call to install_parm_decl. Update comments. (scan_sharing_clauses): Use clause tree expression as splay-tree key for map/to/from and OpenACC firstprivate cases when installing the variable field into the send/receive record type. (maybe_lookup_field_in_outer_ctx): Add code to search through construct clauses instead of entirely based on splay-tree lookup. (lower_oacc_reductions): Adjust to find map-clause of reduction variable, then create receiver-ref. (lower_omp_target): Adjust to lookup var field using clause expression. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Adjust testcase. --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index b977f98af979..3d13da044f0b 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,18 @@ +2021-02-01 Chung-Lin Tang + + * omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE' + default parameter. Set splay-tree lookup key to key_expr instead of + var if key_expr is non-NULL. Adjust call to install_parm_decl. + Update comments. + (scan_sharing_clauses): Use clause tree expression as splay-tree key + for map/to/from and OpenACC firstprivate cases when installing the + variable field into the send/receive record type. + (maybe_lookup_field_in_outer_ctx): Add code to search through + construct clauses instead of entirely based on splay-tree lookup. + (lower_oacc_reductions): Adjust to find map-clause of reduction + variable, then create receiver-ref. + (lower_omp_target): Adjust to lookup var field using clause expression. + 2021-01-16 Andrew Stubbs * dwarf2out.cc (gen_subprogram_die): Check offload attributes only. diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index d8156d528c6b..1c7df9a15a03 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,8 @@ +2021-02-01 Chung-Lin Tang + + * semantics.cc (finish_omp_clauses): Adjust to allow duplicate + mapped variables for OpenMP. + 2020-08-19 Sandra Loosemore Add a "combined" flag for "acc kernels loop" etc directives. diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 2c0e1c7eaa0f..daef4fb581e9 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8271,7 +8271,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && !bitmap_bit_p (&map_field_head, DECL_UID (t))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t)) + && ort != C_ORT_OMP) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index b28d2b81493a..bac635fb58c0 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -773,25 +773,29 @@ build_sender_ref (tree var, omp_context *ctx) return build_sender_ref ((splay_tree_key) var, ctx); } -/* Add a new field for VAR inside the structure CTX->SENDER_DECL. If - BASE_POINTERS_RESTRICT, declare the field with restrict. */ - static void install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, + tree key_expr = NULL_TREE, bool base_pointers_restrict = false) { tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; - if ((mask & 16) != 0) - { - key = (splay_tree_key) &DECL_NAME (var); - gcc_checking_assert (key != (splay_tree_key) var); - } - if ((mask & 8) != 0) + if (key_expr) + /* Allow user to explicitly set the expression used as the key. */ + key = (splay_tree_key) key_expr; + else { - key = (splay_tree_key) &DECL_UID (var); - gcc_checking_assert (key != (splay_tree_key) var); + if ((mask & 16) != 0) + { + key = (splay_tree_key) &DECL_NAME (var); + gcc_checking_assert (key != (splay_tree_key) var); + } + if ((mask & 8) != 0) + { + key = (splay_tree_key) &DECL_UID (var); + gcc_checking_assert (key != (splay_tree_key) var); + } } gcc_assert ((mask & 1) == 0 || !splay_tree_lookup (ctx->field_map, key)); @@ -1488,8 +1492,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) { + /* OpenACC firstprivate clauses are later processed with same + code path as map clauses in lower_omp_target, so follow + the same convention of using the whole clause expression + as splay-tree key. */ + tree k = (is_oacc_parallel_or_serial (ctx) ? c : NULL_TREE); by_ref = !omp_privatize_by_reference (decl); - install_var_field (decl, by_ref, 3, ctx); + install_var_field (decl, by_ref, 3, ctx, k); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) { @@ -1778,7 +1787,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, t = TREE_TYPE (t); } - install_var_field (array_decl, by_ref, 3, ctx); + install_var_field (array_decl, by_ref, 3, ctx, c); install_var_local (array_decl, ctx); break; } @@ -1792,7 +1801,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); gcc_assert (DECL_P (decl2)); - install_var_field (decl2, true, 3, ctx); + install_var_field (decl2, true, 3, ctx, c); install_var_local (decl2, ctx); install_var_local (decl, ctx); } @@ -1802,9 +1811,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) - install_var_field (decl, true, 7, ctx); + install_var_field (decl, true, 7, ctx, c); else - install_var_field (decl, true, 3, ctx, + install_var_field (decl, true, 3, ctx, c, base_pointers_restrict); if (is_gimple_omp_offloaded (ctx->stmt) && !(is_gimple_omp_oacc (ctx->stmt) @@ -1840,7 +1849,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, FIELD_DECL, NULL_TREE, ptr_type_node); SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); insert_field_into_struct (ctx->record_type, field); - splay_tree_insert (ctx->field_map, (splay_tree_key) decl, + splay_tree_insert (ctx->field_map, (splay_tree_key) c, (splay_tree_value) field); } } @@ -4738,8 +4747,19 @@ maybe_lookup_field_in_outer_ctx (tree decl, omp_context *ctx) omp_context *up; for (up = ctx->outer; up; up = up->outer) - if (maybe_lookup_field (decl, up)) - return true; + { + for (tree c = gimple_omp_target_clauses (up->stmt); + c != NULL_TREE; c = OMP_CLAUSE_CHAIN (c)) + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM + || (is_oacc_parallel_or_serial (up) + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)) + && OMP_CLAUSE_DECL (c) == decl) + return true; + if (maybe_lookup_field (decl, up)) + return true; + } return false; } @@ -7736,6 +7756,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); tree orig = OMP_CLAUSE_DECL (c); + tree orig_clause; tree var; tree ref_to_res = NULL_TREE; tree incoming, outgoing; @@ -7818,8 +7839,18 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, do_lookup: /* This is the outermost construct with this reduction, see if there's a mapping for it. */ - if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET - && (maybe_lookup_field (orig, outer) || is_fpp) && !is_private) + orig_clause = NULL_TREE; + if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET) + for (tree cls = gimple_omp_target_clauses (outer->stmt); + cls; cls = OMP_CLAUSE_CHAIN (cls)) + if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_MAP + && orig == OMP_CLAUSE_DECL (cls) + && maybe_lookup_field (cls, outer)) + { + orig_clause = cls; + break; + } + if ((orig_clause != NULL_TREE || is_fpp) && !is_private) { tree type = TREE_TYPE (var); @@ -7831,7 +7862,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, } else { - ref_to_res = build_receiver_ref (orig, false, outer); + ref_to_res = build_receiver_ref (orig_clause, false, outer); if (omp_privatize_by_reference (orig)) ref_to_res = build_simple_mem_ref (ref_to_res); } @@ -13156,7 +13187,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) continue; } - if (!maybe_lookup_field (var, ctx)) + if (!maybe_lookup_field (c, ctx)) continue; /* Don't remap compute constructs' reduction variables, because the @@ -13176,7 +13207,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && TREE_CODE (var_type) != ARRAY_TYPE ? false : true); - x = build_receiver_ref (var, rcv_by_ref, ctx); + x = build_receiver_ref (c, rcv_by_ref, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE && (FLOAT_TYPE_P (inner_type) @@ -13430,7 +13461,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } else { - tree x = build_sender_ref (ovar, ctx); + tree x = build_sender_ref (c, ctx); tree v = ovar; if (in_reduction_clauses && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -13478,7 +13509,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (!maybe_lookup_field (ovar, ctx) + if (!maybe_lookup_field (c, ctx) && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))) @@ -13528,7 +13559,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } else if (nc) { - x = build_sender_ref (ovar, ctx); + x = build_sender_ref (nc, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -14351,7 +14382,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) new_var = decl2; type = TREE_TYPE (new_var); } - x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx); + x = build_receiver_ref (prev, false, ctx); x = fold_convert_loc (clause_loc, type, x); if (!integer_zerop (OMP_CLAUSE_SIZE (c))) { diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index a810edff4b5c..f298d9f838ec 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,7 @@ +2021-02-01 Chung-Lin Tang + + * c-c++-common/gomp/clauses-2.c: Adjust testcase. + 2020-08-30 Sandra Loosemore * c-c++-common/goacc/kernels-loop-annotation-21.c: New. diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c index bbc8fb4e32bf..f1ce73993108 100644 --- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c +++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c @@ -15,7 +15,7 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l) bar (p); #pragma omp target map (p) , map (p[0]) bar (p); - #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (q) map (q) bar (&q); #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */ bar (p);