Re: [PR middle-end/71373] Handle more OMP_CLAUSE_* in nested function decomposition

[Thomas Schwinge <thomas@codesourcery.com>]

Hi!

On Wed, 1 Jun 2016 17:12:17 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Jun 01, 2016 at 05:06:42PM +0200, Thomas Schwinge wrote:
> > Here are the OpenACC bits of <http://gcc.gnu.org/PR71373>.
> > 
> > As we're currently not paying attention to OpenACC tile clauses in the
> > middle end, and thus OMP_CLAUSE_TILE's arguments are not to be considered
> > stable, I opted to simply discard them early, simplifying their
> > gcc/tree-nested.c handling.  Everything else should be self-explanatory.
> > 
> > OK for trunk and gcc-6-branch?
> 
> LGTM for both, but please as a follow-up try to work also on a C testcase
> with nested functions that covers all the clauses (both referencing
> vars/expressions that are defined in the current function and used by a
> nested function, and for vars/expressions that are defined in parent
> function and used in clauses inside of nested function.

OK, I translated gcc/testsuite/gfortran.dg/goacc/subroutines.f90 from
Fortran to C: gcc/testsuite/gcc.dg/goacc/nested.c.  For amusement ;-) I'm
also including the test case that originally made us aware of the
problem, gcc/testsuite/gcc.dg/goacc/pr71373.c.  Oh, and I just remembered
<http://news.gmane.org/find-root.php?message_id=%3C5459732B.1010101%40codesourcery.com%3E>,
so after re-testing, I'll also include these test cases, as far as still
relevant.  Nested function decomposition is not applicable to C++, so we
don't need any C++ test cases, right?

During the translation of gcc/testsuite/gfortran.dg/goacc/subroutines.f90
from Fortran to C, I stumbled upon <https://gcc.gnu.org/PR71381> "C/C++
OpenACC cache directive rejects valid syntax",
<http://news.gmane.org/find-root.php?message_id=%3C877fe7sthf.fsf%40kepler.schwinge.homeip.net%3E>,
so that one will need to go in first, before I'll then commit the
following:

commit 7eff9da0e8fe5eda7d76b9a27dbb1ec4e6183844
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Wed Jun 1 17:01:35 2016 +0200

    [PR middle-end/71373] Handle more OMP_CLAUSE_* in nested function decomposition
    
    	gcc/
    	* gimplify.c (gimplify_adjust_omp_clauses): Discard
    	OMP_CLAUSE_TILE.
    	* omp-low.c (scan_sharing_clauses): Don't expect OMP_CLAUSE_TILE.
    	gcc/testsuite/
    	* c-c++-common/goacc/combined-directives.c: XFAIL tree scanning
    	for OpenACC tile clauses.
    	* gfortran.dg/goacc/combined-directives.f90: Likewise.
    
    	gcc/
    	PR middle-end/71373
    	* tree-nested.c (convert_nonlocal_omp_clauses)
    	(convert_local_omp_clauses): Handle OMP_CLAUSE_ASYNC,
    	OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
    	OMP_CLAUSE__CACHE_, OMP_CLAUSE_TILE.
    	gcc/testsuite/
    	PR middle-end/71373
    	* gcc.dg/goacc/nested.c: New file.
    	* gcc.dg/goacc/pr71373.c: Likewise.
    	* gfortran.dg/goacc/subroutines.f90: Update.
---
 gcc/gimplify.c                                     |   6 ++
 gcc/omp-low.c                                      |   4 +-
 .../c-c++-common/goacc/combined-directives.c       |   3 +-
 gcc/testsuite/gcc.dg/goacc/nested.c                | 100 +++++++++++++++++++++
 gcc/testsuite/gcc.dg/goacc/pr71373.c               |  41 +++++++++
 .../gfortran.dg/goacc/combined-directives.f90      |   3 +-
 gcc/testsuite/gfortran.dg/goacc/subroutines.f90    |  56 ++++++++----
 gcc/tree-nested.c                                  |  30 +++++++
 8 files changed, 221 insertions(+), 22 deletions(-)

diff --git gcc/gimplify.c gcc/gimplify.c
index f12c6a1..7c19cf3 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -8280,7 +8280,13 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_VECTOR:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	  break;
+
 	case OMP_CLAUSE_TILE:
+	  /* We're not yet making use of the information provided by OpenACC
+	     tile clauses.  Discard these here, to simplify later middle end
+	     processing.  */
+	  remove = true;
 	  break;
 
 	default:
diff --git gcc/omp-low.c gcc/omp-low.c
index 91d5fcf..c6ba31c 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2187,7 +2187,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
@@ -2201,6 +2200,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
+	case OMP_CLAUSE_TILE:
 	default:
 	  gcc_unreachable ();
 	}
@@ -2357,7 +2357,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
@@ -2365,6 +2364,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
+	case OMP_CLAUSE_TILE:
 	default:
 	  gcc_unreachable ();
 	}
diff --git gcc/testsuite/c-c++-common/goacc/combined-directives.c gcc/testsuite/c-c++-common/goacc/combined-directives.c
index c2a3c57..3fa800d 100644
--- gcc/testsuite/c-c++-common/goacc/combined-directives.c
+++ gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -111,6 +111,7 @@ test ()
 // { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
-// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
+// XFAILed: OpenACC tile clauses are discarded during gimplification.
+// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" { xfail *-*-* } } }
 // { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git gcc/testsuite/gcc.dg/goacc/nested.c gcc/testsuite/gcc.dg/goacc/nested.c
new file mode 100644
index 0000000..6e1f236
--- /dev/null
+++ gcc/testsuite/gcc.dg/goacc/nested.c
@@ -0,0 +1,100 @@
+/* Exercise how tree-nested.c handles OpenACC clauses.  */
+/* See gcc/testsuite/gfortran.dg/goacc/subroutines.f90 for the Fortran
+   version.  */
+
+int main ()
+{
+#define N 100
+  int nonlocal_arg;
+  int nonlocal_a[N];
+  int nonlocal_i;
+  int nonlocal_j;
+
+  for (int i = 0; i < N; ++i)
+    nonlocal_a[i] = 5;
+  nonlocal_arg = 5;
+
+  void local ()
+  {
+    int local_i;
+    int local_arg;
+    int local_a[N];
+    int local_j;
+
+    for (int i = 0; i < N; ++i)
+      local_a[i] = 5;
+    local_arg = 5;
+
+#pragma acc kernels loop \
+  gang(num:local_arg) worker(local_arg) vector(local_arg) \
+  wait async(local_arg)
+    for (local_i = 0; local_i < N; ++local_i)
+      {
+#pragma acc cache (local_a[local_i:5])
+	local_a[local_i] = 100;
+#pragma acc loop seq tile(*)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+#pragma acc loop auto independent tile(1)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+      }
+
+#pragma acc kernels loop \
+  gang(static:local_arg) worker(local_arg) vector(local_arg) \
+  wait(local_arg, local_arg + 1, local_arg + 2) async
+    for (local_i = 0; local_i < N; ++local_i)
+      {
+#pragma acc cache (local_a[local_i:4])
+	local_a[local_i] = 100;
+#pragma acc loop seq tile(1)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+#pragma acc loop auto independent tile(*)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+      }
+  }
+
+  void nonlocal ()
+  {
+    for (int i = 0; i < N; ++i)
+      nonlocal_a[i] = 5;
+    nonlocal_arg = 5;
+
+#pragma acc kernels loop \
+  gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
+  wait async(nonlocal_arg)
+    for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
+      {
+#pragma acc cache (nonlocal_a[nonlocal_i:3])
+	nonlocal_a[nonlocal_i] = 100;
+#pragma acc loop seq tile(2)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+#pragma acc loop auto independent tile(3)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+      }
+
+#pragma acc kernels loop \
+  gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
+  wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async
+    for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
+      {
+#pragma acc cache (nonlocal_a[nonlocal_i:2])
+	nonlocal_a[nonlocal_i] = 100;
+#pragma acc loop seq tile(*)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+#pragma acc loop auto independent tile(*)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+      }
+  }
+
+  local ();
+  nonlocal ();
+
+  return 0;
+}
diff --git gcc/testsuite/gcc.dg/goacc/pr71373.c gcc/testsuite/gcc.dg/goacc/pr71373.c
new file mode 100644
index 0000000..9381752
--- /dev/null
+++ gcc/testsuite/gcc.dg/goacc/pr71373.c
@@ -0,0 +1,41 @@
+/* Unintentional nested function usage.  */
+/* Due to missing right braces '}', the following functions are parsed as
+   nested functions.  This ran into an ICE.  */
+
+void foo (void)
+{
+  #pragma acc parallel
+  {
+    #pragma acc loop independent
+    for (int i = 0; i < 16; i++)
+      ;
+  // Note right brace '}' commented out here.
+  //}
+}
+void bar (void)
+{
+}
+
+// Adding right brace '}' here, to make this compile.
+}
+
+
+// ..., and the other way round:
+
+void BAR (void)
+{
+// Note right brace '}' commented out here.
+//}
+
+void FOO (void)
+{
+  #pragma acc parallel
+  {
+    #pragma acc loop independent
+    for (int i = 0; i < 16; i++)
+      ;
+  }
+}
+
+// Adding right brace '}' here, to make this compile.
+}
diff --git gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 42a447a..abb5e6b 100644
--- gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -143,7 +143,8 @@ end subroutine test
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
+! XFAILed: OpenACC tile clauses are discarded during gimplification.
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" { xfail *-*-* } } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } }
diff --git gcc/testsuite/gfortran.dg/goacc/subroutines.f90 gcc/testsuite/gfortran.dg/goacc/subroutines.f90
index 6cab798..479ef4f 100644
--- gcc/testsuite/gfortran.dg/goacc/subroutines.f90
+++ gcc/testsuite/gfortran.dg/goacc/subroutines.f90
@@ -1,6 +1,5 @@
-! Exercise how tree-nested.c handles gang, worker vector and seq.
-
-! { dg-do compile } 
+! Exercise how tree-nested.c handles OpenACC clauses.
+! See gcc/testsuite/c-c++-common/goacc/nested.c for the C version.
 
 program main
   integer, parameter :: N = 100
@@ -8,10 +7,10 @@ program main
   integer :: nonlocal_a(N)
   integer :: nonlocal_i
   integer :: nonlocal_j
-  
+
   nonlocal_a (:) = 5
   nonlocal_arg = 5
-  
+
   call local ()
   call nonlocal ()
 
@@ -22,24 +21,35 @@ contains
     integer :: local_arg
     integer :: local_a(N)
     integer :: local_j
-    
+
     local_a (:) = 5
     local_arg = 5
 
-    !$acc kernels loop gang(num:local_arg) worker(local_arg) vector(local_arg)
+    !$acc kernels loop &
+    !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) &
+    !$acc wait async(local_arg)
     do local_i = 1, N
+       !$acc cache (local_a(local_i:local_i + 5))
        local_a(local_i) = 100
-       !$acc loop seq
+       !$acc loop seq tile(*)
+       do local_j = 1, N
+       enddo
+       !$acc loop auto independent tile(1)
        do local_j = 1, N
        enddo
     enddo
     !$acc end kernels loop
 
-    !$acc kernels loop gang(static:local_arg) worker(local_arg) &
-    !$acc vector(local_arg)
+    !$acc kernels loop &
+    !$acc gang(static:local_arg) worker(local_arg) vector(local_arg) &
+    !$acc wait(local_arg, local_arg + 1, local_arg + 2) async
     do local_i = 1, N
+       !$acc cache (local_a(local_i:local_i + 4))
        local_a(local_i) = 100
-       !$acc loop seq
+       !$acc loop seq tile(1)
+       do local_j = 1, N
+       enddo
+       !$acc loop auto independent tile(*)
        do local_j = 1, N
        enddo
     enddo
@@ -49,22 +59,32 @@ contains
   subroutine nonlocal ()
     nonlocal_a (:) = 5
     nonlocal_arg = 5
-  
-    !$acc kernels loop gang(num:nonlocal_arg) worker(nonlocal_arg) &
-    !$acc vector(nonlocal_arg)
+
+    !$acc kernels loop &
+    !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
+    !$acc wait async(nonlocal_arg)
     do nonlocal_i = 1, N
+       !$acc cache (nonlocal_a(nonlocal_i:nonlocal_i + 3))
        nonlocal_a(nonlocal_i) = 100
-       !$acc loop seq
+       !$acc loop seq tile(2)
+       do nonlocal_j = 1, N
+       enddo
+       !$acc loop auto independent tile(3)
        do nonlocal_j = 1, N
        enddo
     enddo
     !$acc end kernels loop
 
-    !$acc kernels loop gang(static:nonlocal_arg) worker(nonlocal_arg) &
-    !$acc vector(nonlocal_arg)
+    !$acc kernels loop &
+    !$acc gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
+    !$acc wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async
     do nonlocal_i = 1, N
+       !$acc cache (nonlocal_a(nonlocal_i:nonlocal_i + 2))
        nonlocal_a(nonlocal_i) = 100
-       !$acc loop seq
+       !$acc loop seq tile(*)
+       do nonlocal_j = 1, N
+       enddo
+       !$acc loop auto independent tile(*)
        do nonlocal_j = 1, N
        enddo
     enddo
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 25a92aa..97d3c52 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1114,6 +1114,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  /* Several OpenACC clauses have optional arguments.  Check if they
 	     are present.  */
 	  if (OMP_CLAUSE_OPERAND (clause, 0))
@@ -1197,8 +1199,21 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_AUTO:
 	  break;
 
+	case OMP_CLAUSE__CACHE_:
+	  /* These clauses belong to the OpenACC cache directive, which is
+	     discarded during gimplification, so we don't expect to see
+	     anything here.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE_TILE:
+	  /* OpenACC tile clauses are discarded during gimplification, so we
+	     don't expect to see anything here.  */
+	  gcc_unreachable ();
+
 	default:
 	  gcc_unreachable ();
 	}
@@ -1790,6 +1805,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  /* Several OpenACC clauses have optional arguments.  Check if they
 	     are present.  */
 	  if (OMP_CLAUSE_OPERAND (clause, 0))
@@ -1878,8 +1895,21 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_AUTO:
 	  break;
 
+	case OMP_CLAUSE__CACHE_:
+	  /* These clauses belong to the OpenACC cache directive, which is
+	     discarded during gimplification, so we don't expect to see
+	     anything here.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE_TILE:
+	  /* OpenACC tile clauses are discarded during gimplification, so we
+	     don't expect to see anything here.  */
+	  gcc_unreachable ();
+
 	default:
 	  gcc_unreachable ();
 	}


Grüße
 Thomas