public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r12-1109] openmp: Fix up handling of reduction clause on constructs combined with target [PR99928]
@ 2021-05-28 9:34 Jakub Jelinek
0 siblings, 0 replies; only message in thread
From: Jakub Jelinek @ 2021-05-28 9:34 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:c94424b0ed786ec92b6904da69af8b5243b34fdc
commit r12-1109-gc94424b0ed786ec92b6904da69af8b5243b34fdc
Author: Jakub Jelinek <jakub@redhat.com>
Date: Fri May 28 11:26:48 2021 +0200
openmp: Fix up handling of reduction clause on constructs combined with target [PR99928]
The reduction clause should be copied as map (tofrom: ) to combined target if
present, but as we need different handling of array sections between map and reduction,
doing that during gimplification would be harder.
So, this patch adds them during splitting, and similarly to firstprivate adds them
with a new flag that they should be just ignored/removed if an explicit map clause
of the same list item is present.
The exact rules are to be decided in https://github.com/OpenMP/spec/issues/2766
so this patch just implements something that is IMHO reasonable and exact detailed
testcases for the cornercases will follow once it is clarified.
2021-05-28 Jakub Jelinek <jakub@redhat.com>
PR middle-end/99928
gcc/
* tree.h (OMP_CLAUSE_MAP_IMPLICIT): Define.
gcc/c-family/
* c-omp.c (c_omp_split_clauses): For reduction clause if combined with
target add a map tofrom clause with OMP_CLAUSE_MAP_IMPLICIT.
gcc/c/
* c-typeck.c (handle_omp_array_sections): Copy OMP_CLAUSE_MAP_IMPLICIT.
(c_finish_omp_clauses): Move not just OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT
marked clauses last, but also OMP_CLAUSE_MAP_IMPLICIT. Add
map_firstprivate_head bitmap, set it for GOMP_MAP_FIRSTPRIVATE_POINTER
maps and silently remove OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT if it is
present too. For OMP_CLAUSE_MAP_IMPLICIT silently remove the clause
if present in map_head, map_field_head or map_firstprivate_head
bitmaps.
gcc/cp/
* semantics.c (handle_omp_array_sections): Copy
OMP_CLAUSE_MAP_IMPLICIT.
(finish_omp_clauses): Move not just OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT
marked clauses last, but also OMP_CLAUSE_MAP_IMPLICIT. Add
map_firstprivate_head bitmap, set it for GOMP_MAP_FIRSTPRIVATE_POINTER
maps and silently remove OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT if it is
present too. For OMP_CLAUSE_MAP_IMPLICIT silently remove the clause
if present in map_head, map_field_head or map_firstprivate_head
bitmaps.
gcc/testsuite/
* c-c++-common/gomp/pr99928-8.c: Remove all xfails.
* c-c++-common/gomp/pr99928-9.c: Likewise.
* c-c++-common/gomp/pr99928-10.c: Likewise.
* c-c++-common/gomp/pr99928-16.c: New test.
Diff:
---
gcc/c-family/c-omp.c | 10 +++++
gcc/c/c-typeck.c | 63 +++++++++++++++++++++------
gcc/cp/semantics.c | 65 ++++++++++++++++++++++------
gcc/testsuite/c-c++-common/gomp/pr99928-10.c | 48 ++++++++++----------
gcc/testsuite/c-c++-common/gomp/pr99928-16.c | 16 +++++++
gcc/testsuite/c-c++-common/gomp/pr99928-8.c | 44 +++++++++----------
gcc/testsuite/c-c++-common/gomp/pr99928-9.c | 22 +++++-----
gcc/tree.h | 5 +++
8 files changed, 190 insertions(+), 83 deletions(-)
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 1b4a1124c25..28fbb1d0537 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -1989,6 +1989,16 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
"%<parallel for%>, %<parallel for simd%>");
OMP_CLAUSE_REDUCTION_INSCAN (clauses) = 0;
}
+ if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)) != 0)
+ {
+ c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+ OMP_CLAUSE_MAP_IMPLICIT (c) = 1;
+ OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c;
+ }
if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SCHEDULE)) != 0)
{
if (code == OMP_SIMD)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index bee01df1234..5f322874423 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13657,6 +13657,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !c_mark_addressable (t))
return false;
@@ -13927,7 +13928,8 @@ tree
c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
+ bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
+ bitmap_head oacc_reduction_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -13947,7 +13949,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
has been seen, -2 if mixed inscan/normal reduction diagnosed. */
int reduction_seen = 0;
bool allocate_seen = false;
- bool firstprivate_implicit_moved = false;
+ bool implicit_moved = false;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -13957,6 +13959,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack);
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
@@ -14389,17 +14392,25 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
- && !firstprivate_implicit_moved)
+ if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved)
{
- firstprivate_implicit_moved = true;
- /* Move firstprivate clauses with
- OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of
+ move_implicit:
+ implicit_moved = true;
+ /* Move firstprivate and map clauses with
+ OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of
clauses chain. */
- tree cl = NULL, *pc1 = pc, *pc2 = &cl;
+ tree cl1 = NULL_TREE, cl2 = NULL_TREE;
+ tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2;
while (*pc1)
if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE
&& OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1))
+ {
+ *pc3 = *pc1;
+ pc3 = &OMP_CLAUSE_CHAIN (*pc3);
+ *pc1 = OMP_CLAUSE_CHAIN (*pc1);
+ }
+ else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (*pc1))
{
*pc2 = *pc1;
pc2 = &OMP_CLAUSE_CHAIN (*pc2);
@@ -14407,10 +14418,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
pc1 = &OMP_CLAUSE_CHAIN (*pc1);
- *pc2 = NULL;
- *pc1 = cl;
- if (pc1 != pc)
- continue;
+ *pc3 = NULL;
+ *pc2 = cl2;
+ *pc1 = cl1;
+ continue;
}
t = OMP_CLAUSE_DECL (c);
need_complete = true;
@@ -14421,6 +14432,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qE is not a variable in clause %<firstprivate%>", t);
remove = true;
}
+ else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
+ && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c)
+ && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
+ remove = true;
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
@@ -14673,6 +14688,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
case OMP_CLAUSE_MAP:
+ if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
+ goto move_implicit;
+ /* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
@@ -14706,6 +14724,16 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
while (TREE_CODE (t) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (c)
+ && (bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || bitmap_bit_p (&map_firstprivate_head,
+ DECL_UID (t))))
+ {
+ remove = true;
+ break;
+ }
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14859,6 +14887,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (c)
+ && (bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))))
+ remove = true;
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
@@ -14880,7 +14914,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
}
else
- bitmap_set_bit (&generic_head, DECL_UID (t));
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
+ }
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& (ort != C_ORT_OMP
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 6fafd0648e0..fe370a21366 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5548,6 +5548,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !cxx_mark_addressable (t))
return false;
@@ -5574,6 +5575,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
+ OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
OMP_CLAUSE_DECL (c3) = ptr;
if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER
|| OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH)
@@ -6510,7 +6512,8 @@ tree
finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
+ bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
+ bitmap_head oacc_reduction_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -6527,7 +6530,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bool allocate_seen = false;
tree detach_seen = NULL_TREE;
bool mergeable_seen = false;
- bool firstprivate_implicit_moved = false;
+ bool implicit_moved = false;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -6537,6 +6540,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack);
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
@@ -6852,17 +6856,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
- && !firstprivate_implicit_moved)
+ if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved)
{
- firstprivate_implicit_moved = true;
- /* Move firstprivate clauses with
- OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of
+ move_implicit:
+ implicit_moved = true;
+ /* Move firstprivate and map clauses with
+ OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of
clauses chain. */
- tree cl = NULL, *pc1 = pc, *pc2 = &cl;
+ tree cl1 = NULL_TREE, cl2 = NULL_TREE;
+ tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2;
while (*pc1)
if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE
&& OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1))
+ {
+ *pc3 = *pc1;
+ pc3 = &OMP_CLAUSE_CHAIN (*pc3);
+ *pc1 = OMP_CLAUSE_CHAIN (*pc1);
+ }
+ else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (*pc1))
{
*pc2 = *pc1;
pc2 = &OMP_CLAUSE_CHAIN (*pc2);
@@ -6870,10 +6882,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
pc1 = &OMP_CLAUSE_CHAIN (*pc1);
- *pc2 = NULL;
- *pc1 = cl;
- if (pc1 != pc)
- continue;
+ *pc3 = NULL;
+ *pc2 = cl2;
+ *pc1 = cl1;
+ continue;
}
t = omp_clause_decl_field (OMP_CLAUSE_DECL (c));
if (t)
@@ -6904,6 +6916,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t);
remove = true;
}
+ else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
+ && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c)
+ && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
+ remove = true;
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
@@ -7620,6 +7636,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
case OMP_CLAUSE_MAP:
+ if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
+ goto move_implicit;
+ /* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
@@ -7651,6 +7670,16 @@ 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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (c)
+ && (bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || bitmap_bit_p (&map_firstprivate_head,
+ DECL_UID (t))))
+ {
+ remove = true;
+ break;
+ }
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -7831,6 +7860,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD is not a pointer variable", t);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (c)
+ && (bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || bitmap_bit_p (&map_firstprivate_head,
+ DECL_UID (t))))
+ remove = true;
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
@@ -7852,7 +7888,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
}
else
- bitmap_set_bit (&generic_head, DECL_UID (t));
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
+ }
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t)))
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
index d2987a1d9b4..6c44600b07e 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
@@ -95,22 +95,22 @@ bar (void)
#pragma omp section
r12[1]++;
}
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r13 \\+ 4" "gimple" } } */
#pragma omp target parallel reduction(+:r13[1:15])
r13[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME. */
#pragma omp target parallel for reduction(+:r14[1:16])
for (int i = 0; i < 64; i++)
r14[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME. */
@@ -118,31 +118,31 @@ bar (void)
#pragma omp target parallel for simd reduction(+:r15[1:17])
for (int i = 0; i < 64; i++)
r15[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization. */
/* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */
#pragma omp target parallel loop reduction(+:r16[1:18])
for (int i = 0; i < 64; i++)
r16[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r17 \\+ 4" "gimple" } } */
#pragma omp target teams reduction(+:r17[1:19])
r17[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */
#pragma omp target teams distribute reduction(+:r18[1:20])
for (int i = 0; i < 64; i++)
r18[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */
@@ -151,8 +151,8 @@ bar (void)
#pragma omp target teams distribute parallel for reduction(+:r19[1:21])
for (int i = 0; i < 64; i++)
r19[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */
@@ -162,8 +162,8 @@ bar (void)
#pragma omp target teams distribute parallel for simd reduction(+:r20[1:22])
for (int i = 0; i < 64; i++)
r20[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */
@@ -171,9 +171,9 @@ bar (void)
#pragma omp target teams distribute simd reduction(+:r21[1:23])
for (int i = 0; i < 64; i++)
r21[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization. */
/* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */
@@ -182,8 +182,8 @@ bar (void)
#pragma omp target teams loop reduction(+:r22[1:24])
for (int i = 0; i < 64; i++)
r22[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r23 \\+ 4" "gimple" } } */
#pragma omp target simd reduction(+:r23[1:25])
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-16.c b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c
new file mode 100644
index 00000000000..84cd85da7cf
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c
@@ -0,0 +1,16 @@
+/* PR middle-end/99928 */
+
+void
+foo (void)
+{
+ int a[6] = {};
+ #pragma omp target simd reduction(+:a[:3])
+ for (int i = 0; i < 6; i++)
+ a[0]++;
+ #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a)
+ for (int i = 0; i < 6; i++)
+ a[0]++;
+ #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a[:6])
+ for (int i = 0; i < 6; i++)
+ a[0]++;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
index fc57573eaf5..27e6ad1eb36 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
@@ -94,48 +94,48 @@ bar (void)
#pragma omp section
r12++;
}
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r13\\)" "gimple" } } */
#pragma omp target parallel reduction(+:r13)
r13++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME. */
#pragma omp target parallel for reduction(+:r14)
for (int i = 0; i < 64; i++)
r14++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME. */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } */
#pragma omp target parallel for simd reduction(+:r15)
for (int i = 0; i < 64; i++)
r15++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail. */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail. */
#pragma omp target parallel loop reduction(+:r16)
for (int i = 0; i < 64; i++)
r16++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r17\\)" "gimple" } } */
#pragma omp target teams reduction(+:r17)
r17++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */
#pragma omp target teams distribute reduction(+:r18)
for (int i = 0; i < 64; i++)
r18++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } *//* FIXME: This should be on for instead. */
@@ -143,8 +143,8 @@ bar (void)
#pragma omp target teams distribute parallel for reduction(+:r19)
for (int i = 0; i < 64; i++)
r19++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } *//* FIXME: This should be on for instead. */
@@ -153,16 +153,16 @@ bar (void)
#pragma omp target teams distribute parallel for simd reduction(+:r20)
for (int i = 0; i < 64; i++)
r20++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
#pragma omp target teams distribute simd reduction(+:r21)
for (int i = 0; i < 64; i++)
r21++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */
@@ -171,8 +171,8 @@ bar (void)
#pragma omp target teams loop reduction(+:r22)
for (int i = 0; i < 64; i++)
r22++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" { xfail *-*-* } } } */
- /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" } } */
+ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r23\\)" "gimple" } } */
#pragma omp target simd reduction(+:r23)
for (int i = 0; i < 64; i++)
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
index c049bb0d257..86235277667 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
@@ -94,19 +94,19 @@ bar (void)
#pragma omp section
r12[1]++;
}
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r13 \\+ 4" "gimple" } } */
#pragma omp target parallel reduction(+:r13[1:2])
r13[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME. */
#pragma omp target parallel for reduction(+:r14[1:2])
for (int i = 0; i < 64; i++)
r14[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */
/* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME. */
@@ -114,7 +114,7 @@ bar (void)
#pragma omp target parallel for simd reduction(+:r15[1:2])
for (int i = 0; i < 64; i++)
r15[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */
@@ -122,19 +122,19 @@ bar (void)
#pragma omp target parallel loop reduction(+:r16[1:2])
for (int i = 0; i < 64; i++)
r16[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r17 \\+ 4" "gimple" } } */
#pragma omp target teams reduction(+:r17[1:2])
r17[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */
#pragma omp target teams distribute reduction(+:r18[1:2])
for (int i = 0; i < 64; i++)
r18[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */
@@ -143,7 +143,7 @@ bar (void)
#pragma omp target teams distribute parallel for reduction(+:r19[1:2])
for (int i = 0; i < 64; i++)
r19[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */
@@ -153,7 +153,7 @@ bar (void)
#pragma omp target teams distribute parallel for simd reduction(+:r20[1:2])
for (int i = 0; i < 64; i++)
r20[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */
@@ -161,7 +161,7 @@ bar (void)
#pragma omp target teams distribute simd reduction(+:r21[1:2])
for (int i = 0; i < 64; i++)
r21[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */
@@ -171,7 +171,7 @@ bar (void)
#pragma omp target teams loop reduction(+:r22[1:2])
for (int i = 0; i < 64; i++)
r22[1]++;
- /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+ /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
/* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
/* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r23 \\+ 4" "gimple" } } */
#pragma omp target simd reduction(+:r23[1:2])
diff --git a/gcc/tree.h b/gcc/tree.h
index 37aca8963fe..5935d0e9f7c 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1654,6 +1654,11 @@ class auto_suppress_location_wrappers
variable. */
#define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero on map clauses added implicitly for reduction clauses on combined
+ or composite constructs. They shall be removed if there is an explicit
+ map clause. */
+#define OMP_CLAUSE_MAP_IMPLICIT(NODE) \
+ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag)
/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
clause. */
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-28 9:34 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-28 9:34 [gcc r12-1109] openmp: Fix up handling of reduction clause on constructs combined with target [PR99928] Jakub Jelinek
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).