[gcc r12-5498] Restore previous OpenACC implicit data clauses ordering [PR103244]

[Thomas Schwinge <tschwinge@gcc.gnu.org>]

https://gcc.gnu.org/g:fdd34569e7a9fc2b6c638a7ef62b965ed7e832ce

commit r12-5498-gfdd34569e7a9fc2b6c638a7ef62b965ed7e832ce
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Mon Nov 22 17:29:09 2021 +0100

    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.

Diff:
---
 gcc/gimplify.c                                     | 22 ++++++++++++++--------
 .../c-c++-common/goacc/combined-reduction.c        |  2 +-
 .../c-c++-common/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;