From aadfc984391c8f28464d448b9bc936344e53e92a Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Thu, 21 Jan 2021 23:04:26 +0800 Subject: [PATCH] Target mapping C++ members inside member functions This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562467.html This patch fixes some problems with target mapping when inside C++ member functions: 1. Allow deref '->' in map clauses. 2. Allow this[X] in map clauses. 3. Create map(this->member) from map(member), when encountering member's FIELD_DECL. This may possibly reverted/updated when a final patch is approved for mainline. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_clause_map): Adjust call to cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true. * semantics.c (handle_omp_array_sections_1): Add handling to create 'this->member' from 'member' FIELD_DECL. (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP map clauses. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-3.C: New test. * g++.dg/gomp/this-2.C: Adjust testcase. --- gcc/cp/parser.c | 2 +- gcc/cp/semantics.c | 31 +++++++++++++++++++++--- gcc/testsuite/g++.dg/gomp/target-3.C | 36 ++++++++++++++++++++++++++++ gcc/testsuite/g++.dg/gomp/this-2.C | 24 +++++++++---------- 4 files changed, 77 insertions(+), 16 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/target-3.C diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 274a1426a4af..c5192278ca8f 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -36728,7 +36728,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, - NULL); + NULL, C_ORT_OMP, true); for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 6d8519c0f619..0d8b1298f56f 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4754,6 +4754,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) + && TREE_CODE (t) == FIELD_DECL) + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); ret = t; if (TREE_CODE (t) == COMPONENT_REF && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -4779,8 +4782,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF) - t = TREE_OPERAND (t, 0); + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) + && TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } } if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); @@ -4800,6 +4807,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } else if (ort == C_ORT_OMP + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM && TREE_CODE (t) == PARM_DECL && DECL_ARTIFICIAL (t) && DECL_NAME (t) == this_identifier) @@ -7361,6 +7371,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_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))) @@ -7475,6 +7490,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) goto handle_map_references; } } + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) + && !processing_template_decl + && TREE_CODE (t) == FIELD_DECL) + { + OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE, + NULL_TREE); + break; + } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) @@ -7501,7 +7524,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (ort != C_ORT_ACC && t == current_class_ptr) + else if (ort != C_ORT_ACC + && ort != C_ORT_OMP + && t == current_class_ptr) { error_at (OMP_CLAUSE_LOCATION (c), "% allowed in OpenMP only in %" diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C new file mode 100644 index 000000000000..fe2e38b46a36 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -0,0 +1,36 @@ +// { dg-do compile } +// { dg-options "-fopenmp -fdump-tree-gimple" } + +struct S +{ + int a, b; + void bar (int); +}; + +void +S::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template +struct T +{ + int a, b; + void bar (int); +}; + +template +void +T::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template struct T<0>; + +/* { dg-final { scan-tree-dump-times "map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C index d03b8a0728ef..b521a4faf5e0 100644 --- a/gcc/testsuite/g++.dg/gomp/this-2.C +++ b/gcc/testsuite/g++.dg/gomp/this-2.C @@ -9,14 +9,14 @@ struct S void S::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template @@ -29,14 +29,14 @@ template void T::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template struct T<0>; -- 2.47.2