This patch tries to allow map(A->ptr) to be properly handled the same way as
map(B.ptr) expressions. map(struct:*A) clauses are now produced during
gimplify.
Julian, I'm CCing you since IIRC you seemed to be the author of this area of
code. Would appreciate if you gave a look if you have time, though I've already
went ahead and pushed to OG10 after testing results looked okay.
Thanks,
Chung-Lin
gcc/ChangeLog:
* gimplify.c ("tree-hash-traits.h"): Add include.
(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
hash_map<tree_operand, tree> *. Adjust struct map handling to handle
cases of *A and A->B expressions.
(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
exit data directives code to earlier position.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-3.C: Adjust testcase gimple scanning.
* g++.dg/gomp/target-this-2.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-23.C: New testcase.
From bf8605f14ec33ea31233a3567f3184fee667b695 Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <clt...@codesourcery.com>
Date: Mon, 8 Feb 2021 07:53:55 -0800
Subject: [PATCH] Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF
(INDIRECT_REF ...)) map clauses.
This patch tries to allow map(A->ptr) to be properly handled the same way as
map(B.ptr) expressions. map(struct:*A) clauses are now produced during
gimplify.
This patch, as of time of commit, is only pushed to devel/omp/gcc-10, not yet
submitted as mainline patch to upstream.
2021-02-08 Chung-Lin Tang <clt...@codesourcery.com>
gcc/ChangeLog:
* gimplify.c ("tree-hash-traits.h"): Add include.
(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
hash_map<tree_operand, tree> *. Adjust struct map handling to handle
cases of *A and A->B expressions.
(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
exit data directives code to earlier position.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-3.C: Adjust testcase gimple scanning.
* g++.dg/gomp/target-this-2.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-23.C: New testcase.
---
gcc/gimplify.c | 51 +++++++++++++++++++++++--------
gcc/testsuite/g++.dg/gomp/target-3.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-3.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +--
libgomp/testsuite/libgomp.c++/target-23.C | 34 +++++++++++++++++++++
6 files changed, 78 insertions(+), 17 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c++/target-23.C
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index b90ba5b..ba19017 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -53,6 +53,7 @@ along with GCC; see the file COPYING3. If not see
#include "langhooks.h"
#include "tree-cfg.h"
#include "tree-ssa.h"
+#include "tree-hash-traits.h"
#include "omp-general.h"
#include "omp-low.h"
#include "gimple-low.h"
@@ -8514,7 +8515,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
- hash_map<tree, tree> *struct_map_to_clause = NULL;
+ hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
@@ -9082,12 +9083,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
&& TREE_CODE (decl) == INDIRECT_REF
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE))
+ == REFERENCE_TYPE)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
{
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
bool indir_p = false;
+ bool component_ref_p = false;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9098,6 +9102,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
while (TREE_CODE (decl) == COMPONENT_REF)
{
decl = TREE_OPERAND (decl, 0);
+ component_ref_p = true;
if (((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
@@ -9117,8 +9122,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
}
}
}
- else if (TREE_CODE (decl) == COMPONENT_REF)
+ else if (TREE_CODE (decl) == COMPONENT_REF
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
{
+ component_ref_p = true;
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
if (TREE_CODE (decl) == INDIRECT_REF
@@ -9188,7 +9196,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
if (code == OACC_UPDATE
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- if (DECL_P (decl)
+ if ((DECL_P (decl)
+ || (component_ref_p && INDIRECT_REF_P (decl)))
&& 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
@@ -9245,7 +9254,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
gcc_assert (base == decl);
splay_tree_node n
- = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+ = (DECL_P (decl)
+ ? splay_tree_lookup (ctx->variables,
+ (splay_tree_key) decl)
+ : NULL);
bool ptr = (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_ALWAYS_POINTER);
bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
@@ -9271,7 +9283,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
- if (n == NULL || (n->value & GOVD_MAP) == 0)
+ if ((DECL_P (decl)
+ && (n == NULL || (n->value & GOVD_MAP) == 0))
+ || (!DECL_P (decl)
+ && (!struct_map_to_clause
+ || struct_map_to_clause->get (decl) == NULL)))
{
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
@@ -9290,7 +9306,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
: TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
if (struct_map_to_clause == NULL)
- struct_map_to_clause = new hash_map<tree, tree>;
+ struct_map_to_clause
+ = new hash_map<tree_operand_hash, tree>;
struct_map_to_clause->put (decl, l);
if (ptr || attach_detach)
{
@@ -9324,7 +9341,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
flags |= GOVD_SEEN;
if (has_attachments)
flags |= GOVD_MAP_HAS_ATTACHMENTS;
- goto do_add_decl;
+ if (DECL_P (decl))
+ goto do_add_decl;
}
else if (struct_map_to_clause)
{
@@ -9433,6 +9451,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
}
else if (*sc != c)
{
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue)
+ == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
*list_p = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *sc;
*sc = c;
@@ -10811,6 +10836,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
}
}
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+ && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA))
+ {
+ remove = true;
+ break;
+ }
if (!DECL_P (decl))
{
if ((ctx->region_type & ORT_TARGET) != 0
@@ -10857,10 +10888,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
= OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
}
}
- else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && (code == OMP_TARGET_EXIT_DATA
- || code == OACC_EXIT_DATA))
- remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C
b/gcc/testsuite/g++.dg/gomp/target-3.C
index fe2e38b..f4d40ec 100644
--- a/gcc/testsuite/g++.dg/gomp/target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -33,4 +33,4 @@ T<N>::bar (int x)
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" } } */
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\)
map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len:
\[0-9\]+\\\]\\)" 4 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C
b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index a5e8321..679c85a 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -46,4 +46,4 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+
\[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias:
0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len:
[0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" }
} */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C
b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 208ea07..a450b37 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign,
bias: 0\]\) map\(alloc:\*this->refptr \[pointer assign, zero-length array
section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped
\[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple"
} } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign,
bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section,
bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len:
[0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign,
bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\)
map\(from:mapped \[len: [0-9]+\]\) map\(alloc:MEM.* \[len: 0\]\)
firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C
b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index f42cf38..af23cbb 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer
assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)
map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)
map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias:
0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer
assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\)
map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len:
[0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)
map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)
map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias:
0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer
assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias:
0\]\) map\(attach:_3->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len:
0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer
assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\)
map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len:
[0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\)
map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len:
0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C
b/libgomp/testsuite/libgomp.c++/target-23.C
new file mode 100644
index 0000000..d4f9ff3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -0,0 +1,34 @@
+extern "C" void abort ();
+
+struct S
+{
+ int *data;
+};
+
+int
+main (void)
+{
+ #define SZ 10
+ S *s = new S ();
+ s->data = new int[SZ];
+
+ for (int i = 0; i < SZ; i++)
+ s->data[i] = 0;
+
+ #pragma omp target enter data map(to: s)
+ #pragma omp target enter data map(to: s->data[:SZ])
+ #pragma omp target
+ {
+ for (int i = 0; i < SZ; i++)
+ s->data[i] = i;
+ }
+ #pragma omp target exit data map(from: s->data[:SZ])
+ #pragma omp target exit data map(from: s)
+
+ for (int i = 0; i < SZ; i++)
+ if (s->data[i] != i)
+ abort ();
+
+ return 0;
+}
+
--
2.8.1
