Hi!
On 2021-11-06T00:51:59+0800, Chung-Lin Tang <clt...@codesourcery.com> wrote:
> On 2021/6/24 11:55 PM, Jakub Jelinek wrote:
>> On Fri, May 14, 2021 at 09:20:25PM +0800, Chung-Lin Tang wrote:
>>> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
>>> index e790f08b23f..69c4a8e0a0a 100644
>>> --- a/gcc/gimplify.c
>>> +++ b/gcc/gimplify.c
>>> @@ -10374,6 +10374,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,
>>> void *data)
>>> gcc_unreachable ();
>>> }
>>> OMP_CLAUSE_SET_MAP_KIND (clause, kind);
>>> + OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1;
>>> if (DECL_SIZE (decl)
>>> && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
>>> {
>> Also as Thomas mentioned, it should be restricted to non-OpenACC,
> Agreed, I've adjusted the patch to only to this implicit setting for OpenMP.
> This reduces a lot of the originally needed scan test adjustment for existing
> OpenACC testcases.
..., but not all, because this piece is still effective:
>>> @@ -10971,9 +10972,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
>>> gimple_seq body, tree *list_p,
>>> list_p = &OMP_CLAUSE_CHAIN (c);
>>> }
>>>
>>> - /* Add in any implicit data sharing. */
>>> + /* Add in any implicit data sharing. Implicit clauses are added at the
>>> start
>>> + of the clause list, but after any non-map clauses. */
>>> struct gimplify_adjust_omp_clauses_data data;
>>> - data.list_p = list_p;
>>> + tree *implicit_add_list_p = orig_list_p;
>>> + while (*implicit_add_list_p
>>> + && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
>>> + implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
..., which effects changes such as:
> --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> -/* { dg-final { scan-tree-dump-times "omp target oacc_parallel
> reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "omp target oacc_parallel
> reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
> --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel
> map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel
> map.tofrom:s .len: 32.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
> --- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
..., and you've changed:
> --- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
> @@ -419,12 +419,7 @@ vla (int array_li)
> copyout (array_so)
> /* The gimplifier has created an implicit 'firstprivate' clause for the
> array
> length.
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower {
> target { ! c++ } } } }
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } }
> }
> - (C++ computes an intermediate value, so can't scan for
> 'firstprivate(array_li)'.) */
> - /* For C, non-LP64, the gimplifier has also created a mapping for the array
> - itself; PR90859.
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)
> map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array
> \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } }
> */
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel
> firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower
> } } */
> {
> array_so = sizeof array;
> }
..., however the clauses reordering alone isn't going to fix PR90859
"[OMP] Mappings for VLA different depending on 'target { c && { ! lp64 } }'",
so it's not correct to just remove that testing/documentation here -- this
change gave rise to PR103244
"c-c++-common/goacc/firstprivate-mappings-1.c fails in certain
configurations since g:b7e20480630e3eeb9eed8b3941da3b3f0c22c969". To
resolve that, and until we properly and deliberately look into also for
OpenACC enabling your "Implement relaxation of implicit map vs. existing
device mappings" (we certainly should!), I've now pushed to master branch
commit fdd34569e7a9fc2b6c638a7ef62b965ed7e832ce "Restore previous OpenACC
implicit data clauses ordering [PR103244]", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
>From fdd34569e7a9fc2b6c638a7ef62b965ed7e832ce Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Mon, 22 Nov 2021 17:29:09 +0100
Subject: [PATCH] Restore previous OpenACC implicit data clauses ordering
[PR103244]
Follow-up for recent commit b7e20480630e3eeb9eed8b3941da3b3f0c22c969
"openmp: Relax handling of implicit map vs. existing device mappings".
As discussed, we likely also for OpenACC ought to use
'OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P' and do the appropriate implicit clauses
ordering -- but that's for a separate step.
gcc/
PR middle-end/103244
* gimplify.c (gimplify_adjust_omp_clauses): Restore previous
OpenACC behavior.
gcc/testsuite/
PR middle-end/103244
* c-c++-common/goacc/combined-reduction.c: Revert/expect previous
OpenACC behavior.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
---
gcc/gimplify.c | 22 ++++++++++++-------
.../c-c++-common/goacc/combined-reduction.c | 2 +-
.../goacc/firstprivate-mappings-1.c | 7 +++++-
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +-
.../g++.dg/goacc/firstprivate-mappings-1.C | 2 +-
5 files changed, 23 insertions(+), 12 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 816cdaf8a18..8624f8221fd 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -11501,15 +11501,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
- /* Add in any implicit data sharing. Implicit clauses are added at the start
- of the clause list, but after any non-map clauses. */
+ /* Add in any implicit data sharing. */
struct gimplify_adjust_omp_clauses_data data;
- tree *implicit_add_list_p = orig_list_p;
- while (*implicit_add_list_p
- && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
- implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
-
- data.list_p = implicit_add_list_p;
+ if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
+ {
+ /* OpenMP. Implicit clauses are added at the start of the clause list,
+ but after any non-map clauses. */
+ tree *implicit_add_list_p = orig_list_p;
+ while (*implicit_add_list_p
+ && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
+ implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
+ data.list_p = implicit_add_list_p;
+ }
+ else
+ /* OpenACC. */
+ data.list_p = list_p;
data.pre_p = pre_p;
splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
index 74ab05bc856..ecf23f59d66 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
@@ -23,7 +23,7 @@ main ()
return 0;
}
-/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
index 5134ef6ed6c..7987beaed9a 100644
--- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
@@ -419,7 +419,12 @@ vla (int array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower } } */
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } } }
+ (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
+ /* For C, non-LP64, the gimplifier has also created a mapping for the array
+ itself; PR90859.
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */
{
array_so = sizeof array;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 0a123bec58f..c2b8dc6c880 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -45,7 +45,7 @@ t1 ()
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.tofrom:s .len: 32.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
index 99a3bd472f7..1b1badb1a90 100644
--- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
+++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
@@ -416,7 +416,7 @@ vla (int &array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
{
array_so = sizeof array;
--
2.33.0
