public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] openmp: Handle C/C++ array reference base-pointers in array sections
@ 2022-06-29 14:44 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:44 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:faf2c46547f42ca9d6e423b88c4728b3bca0f6f4
commit faf2c46547f42ca9d6e423b88c4728b3bca0f6f4
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Thu Feb 24 01:07:48 2022 -0800
openmp: Handle C/C++ array reference base-pointers in array sections
In cases where a program constructs its own deep-copying for arrays-of-pointers,
e.g:
#pragma omp target enter data map(to:level->vectors[:N])
for (i = 0; i < N; i++)
#pragma omp target enter data map(to:level->vectors[i][:N])
We need to treat the part of the array reference before the array section
as a base-pointer (here 'level->vectors[i]'), providing pointer-attachment
behavior.
This patch adds this inside handle_omp_array_sections(), tracing the whole
sequence of array dimensions, creating a whole base-pointer reference
iteratively using build_array_ref(). The conditions are that each of the
"absorbed" dimensions must be length==1, and the final reference must be
of pointer-type (so that pointer attachment makes sense).
Merged from:
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590658.html
2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/ChangeLog:
* c-typeck.cc (handle_omp_array_sections): Add handling for
creating array-reference base-pointer attachment clause.
gcc/cp/ChangeLog:
* semantics.cc (handle_omp_array_sections): Add handling for
creating array-reference base-pointer attachment clause.
gcc/ChangeLog:
* gimplify.cc (gimplify_scan_omp_clauses): Add case for
attach/detach map kind for ARRAY_REF of POINTER_TYPE.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-enter-data-1.c: Adjust testcase.
libgomp/ChangeLog:
* testsuite/libgomp.c-c++-common/ptr-attach-2.c: New test.
Diff:
---
gcc/ChangeLog.omp | 5 ++
gcc/c/ChangeLog.omp | 5 ++
gcc/c/c-typeck.cc | 27 +++++++++-
gcc/cp/ChangeLog.omp | 5 ++
gcc/cp/semantics.cc | 28 +++++++++-
gcc/gimplify.cc | 6 ++-
gcc/testsuite/ChangeLog.omp | 4 ++
.../c-c++-common/gomp/target-enter-data-1.c | 3 +-
libgomp/ChangeLog.omp | 4 ++
.../testsuite/libgomp.c-c++-common/ptr-attach-2.c | 60 ++++++++++++++++++++++
10 files changed, 143 insertions(+), 4 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d093c79efc6..1bfa4f6a664 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * gimplify.cc (gimplify_scan_omp_clauses): Add case for
+ attach/detach map kind for ARRAY_REF of POINTER_TYPE.
+
2022-02-11 Kwok Cheung Yeung <kcy@codesourcery.com>
* omp-general.cc (DELAY_METADIRECTIVES_AFTER_LTO): Check that cfun is
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 3e5fe5fd85e..11d0bbe7818 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * c-typeck.cc (handle_omp_array_sections): Add handling for
+ creating array-reference base-pointer attachment clause.
+
2022-02-18 Kwok Cheung Yeung <kcy@codesourcery.com>
* c-parser.cc (c_parser_omp_construct): Move handling of
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 32bf33e80d6..c974a609513 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13677,6 +13677,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
+ struct dim { tree low_bound, length; };
+ auto_vec<dim> dims (num);
+ dims.safe_grow (num);
+
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_CHAIN (t))
{
@@ -13798,6 +13802,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
else
size = size_binop (MULT_EXPR, size, l);
}
+
+ dim d = { low_bound, length };
+ dims[i] = d;
}
if (non_contiguous)
{
@@ -13845,6 +13852,23 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_DECL (c) = t;
return false;
}
+
+ tree aref = t;
+ for (i = 0; i < dims.length (); i++)
+ {
+ if (dims[i].length && integer_onep (dims[i].length))
+ {
+ tree lb = dims[i].low_bound;
+ aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb);
+ }
+ else
+ {
+ if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE)
+ t = aref;
+ break;
+ }
+ }
+
first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
@@ -13879,7 +13903,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
break;
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- if (TREE_CODE (t) == COMPONENT_REF)
+ if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF
+ || TREE_CODE (t) == INDIRECT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 251fc73f190..b0f96782d48 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * semantics.cc (handle_omp_array_sections): Add handling for
+ creating array-reference base-pointer attachment clause.
+
2022-02-18 Kwok Cheung Yeung <kcy@codesourcery.com>
* parser.cc (cp_parser_omp_construct): Move handling of
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 539bdca7eb0..49cdaf78d73 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5526,6 +5526,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (processing_template_decl && maybe_zero_len)
return false;
+ struct dim { tree low_bound, length; };
+ auto_vec<dim> dims (num);
+ dims.safe_grow (num);
+
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_CHAIN (t))
{
@@ -5645,6 +5649,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
else
size = size_binop (MULT_EXPR, size, l);
}
+
+ dim d = { low_bound, length };
+ dims[i] = d;
}
if (!processing_template_decl)
{
@@ -5696,6 +5703,24 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_DECL (c) = t;
return false;
}
+
+ tree aref = t;
+ for (i = 0; i < dims.length (); i++)
+ {
+ if (dims[i].length && integer_onep (dims[i].length))
+ {
+ tree lb = dims[i].low_bound;
+ aref = convert_from_reference (aref);
+ aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb);
+ }
+ else
+ {
+ if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE)
+ t = aref;
+ break;
+ }
+ }
+
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
return false;
@@ -5730,7 +5755,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
bool reference_always_pointer = true;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- if (TREE_CODE (t) == COMPONENT_REF)
+ if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF
+ || (TREE_CODE (t) == INDIRECT_REF && !REFERENCE_REF_P (t)))
{
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 105bd81b5b8..543b9445092 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10490,7 +10490,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|| (component_ref_p
&& (INDIRECT_REF_P (decl)
|| TREE_CODE (decl) == MEM_REF
- || TREE_CODE (decl) == ARRAY_REF)))))
+ || TREE_CODE (decl) == ARRAY_REF))
+ || (TREE_CODE (decl) == ARRAY_REF
+ && TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+ && (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_DETACH)))))
&& 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
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 72c03338b02..3b9e248f5fa 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * c-c++-common/gomp/target-enter-data-1.c: Adjust testcase.
+
2022-02-18 Kwok Cheung Yeung <kcy@codesourcery.com>
* c-c++-common/gomp/metadirective-1.c (f): Add test for
diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
index ce766d29e2d..3a1b488fa1f 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
@@ -21,4 +21,5 @@ void func (struct foo *f, int n, int m)
#pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
}
-/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\\*_\[0-9\]+ \\\[bias: \[^\]\]+\\\]\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 2 "gimple" } } */
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 5f4fb9362b5..865d7a83994 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,7 @@
+2022-02-24 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * testsuite/libgomp.c-c++-common/ptr-attach-2.c: New test.
+
2022-01-28 Kwok Cheung Yeung <kcy@codesourcery.com>
* testsuite/libgomp.c-c++-common/metadirective-2.c (main): Add
diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c
new file mode 100644
index 00000000000..889a4a253ae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c
@@ -0,0 +1,60 @@
+#include <stdlib.h>
+
+struct blk { int x, y; };
+struct L
+{
+ #define N 10
+ struct {
+ int num_blocks[N];
+ struct blk * blocks[N];
+ } m;
+};
+
+void foo (struct L *l)
+{
+ for (int i = 0; i < N; i++)
+ {
+ l->m.blocks[i] = (struct blk *) malloc (sizeof (struct blk) * N);
+ l->m.num_blocks[i] = N;
+ }
+
+ #pragma omp target enter data map(to:l[:1])
+ for (int i = 0; i < N; i++)
+ {
+ #pragma omp target enter data map(to:l->m.blocks[i][:l->m.num_blocks[i]])
+ }
+
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ l->m.blocks[i][j].x = i + j;
+ l->m.blocks[i][j].y = i * j;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ #pragma omp target exit data map(from:l->m.blocks[i][:l->m.num_blocks[i]])
+ }
+ #pragma omp target exit data map(from:l[:1])
+
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ if (l->m.blocks[i][j].x != i + j)
+ abort ();
+ if (l->m.blocks[i][j].y != i * j)
+ abort ();
+ }
+
+}
+
+int main (void)
+{
+ struct L l;
+ foo (&l);
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:44 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:44 [gcc/devel/omp/gcc-12] openmp: Handle C/C++ array reference base-pointers in array sections Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).