Hi!

On 2021-11-06T00:51:59+0800, Chung-Lin Tang <cltang@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