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 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 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 + +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; +}