@@ -15700,7 +15700,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
}
}
- nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
+ nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
+ C_ORT_OMP, true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -12917,6 +12917,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == MEM_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
@@ -13778,6 +13784,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree ordered_clause = NULL_TREE;
tree schedule_clause = NULL_TREE;
bool oacc_async = false;
+ bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
tree *nogroup_seen = NULL;
@@ -14505,6 +14512,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
while (TREE_CODE (t) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == MEM_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14561,6 +14573,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bias) to zero here, so it is not set erroneously to the pointer
size later on in gimplify.c. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
+ indir_component_ref_p = false;
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+ {
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+ indir_component_ref_p = true;
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
@@ -14633,6 +14654,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
+ && !indir_component_ref_p
&& !c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -9197,7 +9197,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
if ((DECL_P (decl)
- || (component_ref_p && INDIRECT_REF_P (decl)))
+ || (component_ref_p
+ && (INDIRECT_REF_P (decl)
+ || TREE_CODE (decl) == MEM_REF)))
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9298,7 +9300,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (base_ref)
OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
else
- OMP_CLAUSE_DECL (l) = decl;
+ {
+ OMP_CLAUSE_DECL (l) = unshare_expr (decl);
+ if (!DECL_P (OMP_CLAUSE_DECL (l))
+ && (gimplify_expr (&OMP_CLAUSE_DECL (l),
+ pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue)
+ == GS_ERROR))
+ {
+ remove = true;
+ break;
+ }
+ }
OMP_CLAUSE_SIZE (l)
= (!attach
? size_int (1)
@@ -9341,6 +9354,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags |= GOVD_SEEN;
if (has_attachments)
flags |= GOVD_MAP_HAS_ATTACHMENTS;
+
+ /* If this is a *pointer-to-struct expression, make sure a
+ firstprivate map of the base-pointer exists. */
+ if (component_ref_p
+ && ((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ && DECL_P (TREE_OPERAND (decl, 0))
+ && !splay_tree_lookup (ctx->variables,
+ ((splay_tree_key)
+ TREE_OPERAND (decl, 0))))
+ {
+ decl = TREE_OPERAND (decl, 0);
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ enum gomp_map_kind mkind
+ = GOMP_MAP_FIRSTPRIVATE_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
+ }
+
if (DECL_P (decl))
goto do_add_decl;
}
new file mode 100644
@@ -0,0 +1,16 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+struct S
+{
+ int a, b;
+};
+
+void foo (struct S *s)
+{
+ #pragma omp target map (alloc: s->a, s->b)
+ ;
+ #pragma omp target enter data map (alloc: s->a, s->b)
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
This patch is a merge of parts from: https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562467.html and devel/omp/gcc-10 commit 36a1eb, which was a modified merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-November/558975.html to provide the equivalent front-end patches for support "map(A->B)" clauses for the C front-end (only the C++ front-end received such changes before). Some associated middle-end changes are also in this patch. Tested without regressions, and pushed to devel/omp/gcc-10. Chung-Lin From 08caada8efd8f35db634647bbda6091fb667b00d Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang <cltang@codesourcery.com> Date: Mon, 8 Mar 2021 15:56:52 +0800 Subject: [PATCH] Arrow operator handling for C front-end in OpenMP map clauses This patch merges some of the equivalent changes already done for the C++ front-end to the C parts. 2021-03-08 Chung-Lin Tang <cltang@codesourcery.com> gcc/c/ChangeLog: * c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in call to c_parser_omp_variable_list to 'true'. * c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in array base handling. (c_finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ChangeLog: * gimplify.c (gimplify_scan_omp_clauses): Add MEM_REF case when handling component_ref_p case. Add unshare_expr and gimplification when created GOMP_MAP_STRUCT is not a DECL. Add code to add firstprivate pointer for *pointer-to-struct case. gcc/testsuite/ChangeLog: * gcc.dg/gomp/target-3.c: New test. --- gcc/c/c-parser.c | 3 +- gcc/c/c-typeck.c | 22 +++++++++++++++ gcc/gimplify.c | 41 ++++++++++++++++++++++++++-- gcc/testsuite/gcc.dg/gomp/target-3.c | 16 +++++++++++ 4 files changed, 79 insertions(+), 3 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/gomp/target-3.c