Hi Jakub,
this patch is to fix some of the 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 actually may not be the last related patch, since this patch is only for
C++,
while the deref feature parity probably is also needed in C. Also the convention
required "map(struct:) map(<element>)..." sequence still needs further
middle-end
gimplify modifications to generate. However this is already useful for compiling
some programs, and enter/exit data usage.
Tested C++ and libgomp without regressions, is this okay? Probably will need to
wait
till stage1 to commit.
Thanks,
Chung-Lin
2020-12-23 Chung-Lin Tang <clt...@codesourcery.com>
gcc/cp/
* 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/
* g++.dg/gomp/target-3.C: New test.
* g++.dg/gomp/this-2.C: Adjust testcase.
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7ea8c28830e..1d5a508dd71 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37720,7 +37720,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, 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 92e32c8e0ad..d21b58df954 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4808,6 +4808,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree>
&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
@@ -4833,8 +4836,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 && 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);
@@ -4854,6 +4861,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree>
&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)
@@ -7441,6 +7451,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)))
@@ -7555,6 +7570,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)
@@ -7581,7 +7604,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),
"%<this%> allowed in OpenMP only in %<declare simd%>"
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 00000000000..fe2e38b46a3
--- /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 <int N>
+struct T
+{
+ int a, b;
+ void bar (int);
+};
+
+template <int N>
+void
+T<N>::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 d03b8a0728e..b521a4faf5e 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 <int N>
@@ -29,14 +29,14 @@ template <int N>
void
T<N>::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>;
