2015-08-28  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* omp-low.c (extract_oacc_routine_gwv): New function.
	(build_outer_var_ref): Handle refs inside acc routines.
	(scan_omp_for): Check nested parallelism inside acc routines.
	(scan_omp_1_stmt): Check for compatible parallelism when calling
	routines.
	(execute_oacc_transform): Remove dead variable.

	gcc/testsuite/
	* c-c++-common/goacc/routine-6.c: New test.
	* c-c++-common/goacc/routine-7.c: New test.
	* gfortran.dg/goacc/routine-4.f90: New test.
	* gfortran.dg/goacc/routine-5.f90: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/routine-4.c: Fix calls to
	acc routines.
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/vector-routine.f90: Likewise.

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 4312a60..e8d7513 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -415,6 +415,35 @@ is_combined_parallel (struct omp_region *region)
   return region->is_combined_parallel;
 }
 
+/* Return the gang, worker and vector attributes from associated with
+   FNDECL.  Returns a GOMP_DIM for the lowest level of parallelism beginning
+   with GOMP_DIM_GANG, or -1 if the routine is a SEQ. Otherwise, return 0 if
+   the FNDECL is not an acc routine.
+*/
+
+static int
+extract_oacc_routine_gwv (tree fndecl)
+{
+  tree attrs = get_oacc_fn_attrib (fndecl);
+  tree pos;
+  unsigned gwv = 0;
+  int i;
+  int ret = 0;
+
+  if (attrs != NULL_TREE)
+    {
+      for (i = 0, pos = TREE_VALUE (attrs);
+	   gwv == 0 && i != GOMP_DIM_MAX;
+	   i++, pos = TREE_CHAIN (pos))
+	if (TREE_PURPOSE (pos) != boolean_false_node)
+	  return 1 << i;
+
+      ret = -1;
+    }
+
+  return ret;
+}
+
 
 /* Extract the header elements of parallel loop FOR_STMT and store
    them into *FD.  */
@@ -1227,7 +1256,8 @@ build_outer_var_ref (tree var, omp_context *ctx)
       else
 	x = lookup_decl (var, ctx->outer);
     }
-  else if (is_reference (var) || is_oacc_parallel (ctx))
+  else if (is_reference (var) || is_oacc_parallel (ctx)
+	   || extract_oacc_routine_gwv (current_function_decl) != 0)
     /* This can happen with orphaned constructs.  If var is reference, it is
        possible it is shared and as such valid.  */
     x = var;
@@ -2578,9 +2608,16 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
   bool gwv_clause = false;
   bool auto_clause = false;
   bool seq_clause = false;
+  int gwv_routine = 0;
 
   if (outer_ctx)
     outer_type = gimple_code (outer_ctx->stmt);
+  else
+    {
+      gwv_routine = extract_oacc_routine_gwv (current_function_decl);
+      if (gwv_routine > 0)
+	gwv_routine = gwv_routine >> 1;
+    }
 
   ctx = new_omp_context (stmt, outer_ctx);
 
@@ -2699,6 +2736,12 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
 	       && ctx->gwv_this > ctx->gwv_below)
 	error_at (gimple_location (stmt),
 		  "gang, worker and vector must occur in this order in a loop nest");
+      else if (!outer_ctx && ctx->gwv_this != 0 && gwv_routine != 0
+	       && ((ffs (ctx->gwv_this) <= gwv_routine)
+		   || gwv_routine < 0))
+	error_at (gimple_location (stmt),
+		  "invalid parallelism inside acc routine");
+
       if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
 	outer_ctx->gwv_below |= ctx->gwv_below;
     }
@@ -3287,6 +3330,16 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	      default:
 		break;
 	      }
+	  else if (ctx && is_gimple_omp_oacc (ctx->stmt)
+		   && !is_oacc_parallel (ctx))
+	    {
+	      /* Is this a call to an acc routine?  */
+	      int gwv = extract_oacc_routine_gwv (fndecl);
+
+	      if (gwv > 0 && ffs (ctx->gwv_this) >= ffs (gwv))
+		error_at (gimple_location (stmt),
+			  "incompatible parallelism with acc routine");
+	    }
 	}
     }
   if (remove)
@@ -3352,6 +3405,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	    insert_decl_map (&ctx->cb, var, var);
       }
       break;
+
     default:
       *handled_ops_p = false;
       break;
@@ -14815,7 +14869,6 @@ execute_oacc_transform ()
 	    else if (gimple_call_internal_p (stmt))
 	      {
 		unsigned ifn_code = gimple_call_internal_fn (stmt);
-		int retval = 0;
 		switch (ifn_code)
 		  {
 		  default: break;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-6.c b/gcc/testsuite/c-c++-common/goacc/routine-6.c
new file mode 100644
index 0000000..72f035f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-6.c
@@ -0,0 +1,105 @@
+/* Test invalid calls to routines.  */
+/* { dg-do compile } */
+
+#pragma acc routine gang
+int
+gang ()
+{
+  return 1;
+}
+
+#pragma acc routine worker
+int
+worker ()
+{
+  return 1;
+}
+
+#pragma acc routine vector
+int
+vector ()
+{
+  return 1;
+}
+
+#pragma acc routine seq
+int
+seq ()
+{
+  return 1;
+}
+
+int
+main ()
+{
+  int red = 0;
+#pragma acc parallel copy (red)
+  {
+    /* Independent/seq loop tests.  */
+#pragma acc loop reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += gang ();
+
+#pragma acc loop reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += worker ();
+
+#pragma acc loop reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += vector ();
+
+    /* Gang routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+
+#pragma acc loop vector reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+
+    /* Worker routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += worker ();
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += worker (); // { dg-error "incompatible parallelism with acc routine" }
+
+#pragma acc loop vector reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += worker (); // { dg-error "incompatible parallelism with acc routine" }
+
+    /* Vector routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += vector ();
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += vector ();
+
+#pragma acc loop vector reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += vector (); // { dg-error "incompatible parallelism with acc routine" }
+
+    /* Seq routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += seq ();
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += seq ();
+
+#pragma acc loop vector reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += seq ();
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-7.c b/gcc/testsuite/c-c++-common/goacc/routine-7.c
new file mode 100644
index 0000000..96f3e3a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-7.c
@@ -0,0 +1,94 @@
+/* Test invalid intra-routine parallelism.  */
+/* { dg-do compile } */
+
+#pragma acc routine gang
+int
+gang (int red)
+{
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+  return 1;
+}
+
+#pragma acc routine worker
+int
+worker (int red)
+{
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+  return 1;
+}
+
+#pragma acc routine vector
+int
+vector (int red)
+{
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+  return 1;
+}
+
+#pragma acc routine seq
+int
+seq (int red)
+{
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+  return 1;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
new file mode 100644
index 0000000..2587ffd
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
@@ -0,0 +1,160 @@
+! Test invalid calls to routines.
+
+module param
+  integer, parameter :: N = 32
+end module param
+
+program main
+  use param
+  integer :: i
+  integer :: a(N)
+
+  do i = 1, N
+    a(i) = i
+  end do
+
+  !
+  ! Seq routine tests.
+  !
+
+  !$acc parallel copy (a)
+  !$acc loop
+  do i = 1, N
+     call seq (a)
+  end do
+
+  !$acc loop gang
+  do i = 1, N
+     call seq (a)
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     call seq (a)
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     call seq (a)
+  end do
+  !$acc end parallel
+
+  !
+  ! Gang routines loops.
+  !
+
+  !$acc parallel copy (a)
+  !$acc loop
+  do i = 1, N
+     call gang (a)
+  end do
+
+  !$acc loop gang
+  do i = 1, N
+     call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+  !$acc end parallel
+
+  !
+  ! Worker routines loops.
+  !
+
+  !$acc parallel copy (a)
+  !$acc loop
+  do i = 1, N
+     call worker (a)
+  end do
+
+  !$acc loop gang
+  do i = 1, N
+     call worker (a)
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     call worker (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     call worker (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+  !$acc end parallel
+
+  !
+  ! Vector routines loops.
+  !
+
+  !$acc parallel copy (a)
+  !$acc loop
+  do i = 1, N
+     call vector (a)
+  end do
+
+  !$acc loop gang
+  do i = 1, N
+     call vector (a)
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     call vector (a)
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     call vector (a) ! { dg-error "incompatible parallelism with acc routine" }
+  end do
+  !$acc end parallel
+contains
+
+  subroutine gang (a)
+    !$acc routine gang
+    integer, intent (inout) :: a(N)
+    integer :: i
+
+    do i = 1, N
+       a(i) = a(i) - a(i)
+    end do
+  end subroutine gang
+
+  subroutine worker (a)
+    !$acc routine worker
+    integer, intent (inout) :: a(N)
+    integer :: i
+
+    do i = 1, N
+       a(i) = a(i) - a(i)
+    end do
+  end subroutine worker
+
+  subroutine vector (a)
+    !$acc routine vector
+    integer, intent (inout) :: a(N)
+    integer :: i
+
+    do i = 1, N
+       a(i) = a(i) - a(i)
+    end do
+  end subroutine vector
+
+  subroutine seq (a)
+    !$acc routine seq
+    integer, intent (inout) :: a(N)
+    integer :: i
+
+    do i = 1, N
+       a(i) = a(i) - a(i)
+    end do
+  end subroutine seq
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-5.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-5.f90
new file mode 100644
index 0000000..f769334
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-5.f90
@@ -0,0 +1,109 @@
+! Test invalid intra-routine parallellism.
+
+module param
+  integer, parameter :: N = 32
+end module param
+
+subroutine gang (a)
+  !$acc routine gang
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop gang
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+end subroutine gang
+
+subroutine worker (a)
+  !$acc routine worker
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop worker
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+end subroutine worker
+
+subroutine vector (a)
+  !$acc routine vector
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop worker ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop vector
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+end subroutine vector
+
+subroutine seq (a)
+  !$acc routine seq
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop worker ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+
+  !$acc loop vector ! { dg-error "invalid parallelism inside acc routine" }
+  do i = 1, N
+     a(i) = a(i) - a(i)
+  end do
+end subroutine seq
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
index c73400c..e9cf4c4 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
@@ -23,10 +23,10 @@ worker (int *b)
 {
   int i, j;
 
-#pragma acc loop gang
+#pragma acc loop worker
   for (i = 0; i < N; i++)
     {
-#pragma acc loop worker
+#pragma acc loop vector
       for (j = 0; j < M; j++)
         b[i * M + j] += b[i  * M + j]; 
     }
@@ -112,7 +112,7 @@ main(int argc, char **argv)
 
 #pragma acc parallel copy (a[0:N])
   {
-#pragma acc loop vector
+#pragma acc loop
     for (i = 0; i < N; i++)
       vector (&a[0]);
   }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 354784e..7fc8169 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -41,7 +41,7 @@ program main
   end do
 
   !$acc parallel copy (b)
-  !$acc loop worker
+  !$acc loop
     do i = 1, N
       call worker (b)
     end do
@@ -56,7 +56,7 @@ program main
   end do
 
   !$acc parallel copy (a)
-  !$acc loop vector
+  !$acc loop
     do i = 1, N
       call vector (a)
     end do
@@ -85,9 +85,9 @@ subroutine worker (b)
   integer, intent (inout) :: b(M*N)
   integer :: i, j
 
-  !$acc loop gang
-  do i = 1, N
   !$acc loop worker
+  do i = 1, N
+  !$acc loop vector
     do j = 1, M
       b(j + ((i - 1) * M)) = b(j + ((i - 1) * M)) + 1
     end do
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90 b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90
index a8d078a..1edcee4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90
@@ -13,14 +13,8 @@ program main
     a(i) = i
   end do
 
-  !
-  ! Appears there's two bugs...
-  ! 1) loop with vector
-  ! 2) loop without vector
-  !
-
   !$acc parallel copy (a)
-  !$acc loop vector
+  !$acc loop worker
     do i = 1, N
       call vector (a)
     end do
@@ -37,6 +31,7 @@ contains
   integer, intent (inout) :: a(N)
   integer :: i
 
+  !$acc loop vector
   do i = 1, N
     a(i) = a(i) - a(i) 
   end do