Hi Jakub,
as encountered 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).
There's also a little patch in gimplify_scan_omp_clauses(), to make sure
the array-ref base-pointer goes down the right path.
This case was encountered when working to make 534.hpgmgfv_t from
SPEChpc 2021 properly compile. Tested without regressions on trunk.
Okay to go in once stage1 opens?
Thanks,
Chung-Lin
2022-02-21 Chung-Lin Tang <clt...@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/testsuite/ChangeLog:
* libgomp.c-c++-common/ptr-attach-2.c: New test.diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 3075c883548..4257e373557 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13649,6 +13649,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))
{
@@ -13763,6 +13767,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 (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
@@ -13802,6 +13809,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)
@@ -13836,7 +13860,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/semantics.cc b/gcc/cp/semantics.cc
index 0cb17a6a8ab..646f4883d66 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5497,6 +5497,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))
{
@@ -5604,6 +5608,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)
{
@@ -5647,6 +5654,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;
@@ -5681,7 +5706,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 f570daa015a..77b95cd8000 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9626,7 +9626,10 @@ 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/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/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;
+}
