Update OpenACC test cases

	gcc/testsuite/
	* c-c++-common/goacc/combined-directives.c: Clean up dg-*
	directives.
	* c-c++-common/goacc/loop-clauses.c: Likewise.
	* g++.dg/goacc/template.C: Likewise.
	* gfortran.dg/goacc/combined-directives.f90: Likewise.
	* gfortran.dg/goacc/loop-1.f95: Likewise.
	* gfortran.dg/goacc/loop-5.f95: Likewise.
	* gfortran.dg/goacc/loop-6.f95: Likewise.
	* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
	* c-c++-common/goacc-gomp/nesting-1.c: Update.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/clauses-fail.c: Likewise.
	* c-c++-common/goacc/parallel-1.c: Likewise.
	* c-c++-common/goacc/reduction-1.c: Likewise.
	* c-c++-common/goacc/reduction-2.c: Likewise.
	* c-c++-common/goacc/reduction-3.c: Likewise.
	* c-c++-common/goacc/reduction-4.c: Likewise.
	* c-c++-common/goacc/routine-3.c: Likewise.
	* c-c++-common/goacc/routine-4.c: Likewise.
	* c-c++-common/goacc/routine-5.c: Likewise.
	* c-c++-common/goacc/tile.c: Likewise.
	* g++.dg/goacc/template.C: Likewise.
	* gfortran.dg/goacc/combined-directives.f90: Likewise.
	* c-c++-common/goacc/nesting-1.c: Move dg-error test cases into...
	* c-c++-common/goacc/nesting-fail-1.c: ... this file.  Update.
	* c-c++-common/goacc/kernels-1.c: Update.  Incorporate...
	* c-c++-common/goacc/kernels-empty.c: ... this file, and...
	* c-c++-common/goacc/kernels-eternal.c: ... this file, and...
	* c-c++-common/goacc/kernels-noreturn.c: ... this file.
	* c-c++-common/goacc/host_data-1.c: New file.  Incorporate...
	* c-c++-common/goacc/use_device-1.c: ... this file.
	* c-c++-common/goacc/host_data-2.c: New file.  Incorporate...
	* c-c++-common/goacc/host_data-5.c: ... this file, and...
	* c-c++-common/goacc/host_data-6.c: ... this file.
	* c-c++-common/goacc/loop-2-kernels.c: New file.
	* c-c++-common/goacc/loop-2-parallel.c: Likewise.
	* c-c++-common/goacc/loop-3.c: Likewise.
	* g++.dg/goacc/reference.C: Likewise.
	* g++.dg/goacc/routine-1.C: Likewise.
	* g++.dg/goacc/routine-2.C: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Update.
	* testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/vector-loop.c: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/declare-1.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise.
	XFAIL.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
	Incorporate...
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: ... this
	file.
	* testsuite/libgomp.oacc-c++/template-reduction.C: New file.
	* testsuite/libgomp.oacc-c-c++-common/gang-static-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/gang-static-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-loop-clauses.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/private-variables.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise.
	* testsuite/libgomp.oacc-fortran/clauses-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/firstprivate-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/gang-static-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/implicit-firstprivate-ref.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/pr68813.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Merge this
	file...
	* testsuite/libgomp.oacc-c-c++-common/parallel-1.c: ..., and this
	file into...
	* testsuite/libgomp.oacc-c-c++-common/data-clauses.h: ... this new
	file.  Update.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-2.c: Rename to...
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c:
	... this new file.  Update.
	* testsuite/libgomp.oacc-c-c++-common/parallel-2.c: Rename to...
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c:
	... this new file.  Update.
	* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: New
	file.  Incorporate...
	* testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c: ... this
	file, and...
	* testsuite/libgomp.oacc-c-c++-common/worker-single-4.c: ... this
	file, and...
	* testsuite/libgomp.oacc-c-c++-common/worker-single-6.c: ... this
	file.
	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Remove file.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com>
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r234575

37 files changed, 1357 insertions, 339 deletions

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 658e6c5..f4a73a7 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,52 @@
+2016-03-30  Thomas Schwinge  <thomas@codesourcery.com>
+	    Julian Brown  <julian@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Tom de Vries  <tom@codesourcery.com>
+	    Nathan Sidwell  <nathan@codesourcery.com>
+
+	* c-c++-common/goacc/combined-directives.c: Clean up dg-*
+	directives.
+	* c-c++-common/goacc/loop-clauses.c: Likewise.
+	* g++.dg/goacc/template.C: Likewise.
+	* gfortran.dg/goacc/combined-directives.f90: Likewise.
+	* gfortran.dg/goacc/loop-1.f95: Likewise.
+	* gfortran.dg/goacc/loop-5.f95: Likewise.
+	* gfortran.dg/goacc/loop-6.f95: Likewise.
+	* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
+	* c-c++-common/goacc-gomp/nesting-1.c: Update.
+	* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
+	* c-c++-common/goacc/clauses-fail.c: Likewise.
+	* c-c++-common/goacc/parallel-1.c: Likewise.
+	* c-c++-common/goacc/reduction-1.c: Likewise.
+	* c-c++-common/goacc/reduction-2.c: Likewise.
+	* c-c++-common/goacc/reduction-3.c: Likewise.
+	* c-c++-common/goacc/reduction-4.c: Likewise.
+	* c-c++-common/goacc/routine-3.c: Likewise.
+	* c-c++-common/goacc/routine-4.c: Likewise.
+	* c-c++-common/goacc/routine-5.c: Likewise.
+	* c-c++-common/goacc/tile.c: Likewise.
+	* g++.dg/goacc/template.C: Likewise.
+	* gfortran.dg/goacc/combined-directives.f90: Likewise.
+	* c-c++-common/goacc/nesting-1.c: Move dg-error test cases into...
+	* c-c++-common/goacc/nesting-fail-1.c: ... this file.  Update.
+	* c-c++-common/goacc/kernels-1.c: Update.  Incorporate...
+	* c-c++-common/goacc/kernels-empty.c: ... this file, and...
+	* c-c++-common/goacc/kernels-eternal.c: ... this file, and...
+	* c-c++-common/goacc/kernels-noreturn.c: ... this file.
+	* c-c++-common/goacc/host_data-1.c: New file.  Incorporate...
+	* c-c++-common/goacc/use_device-1.c: ... this file.
+	* c-c++-common/goacc/host_data-2.c: New file.  Incorporate...
+	* c-c++-common/goacc/host_data-5.c: ... this file, and...
+	* c-c++-common/goacc/host_data-6.c: ... this file.
+	* c-c++-common/goacc/loop-2-kernels.c: New file.
+	* c-c++-common/goacc/loop-2-parallel.c: Likewise.
+	* c-c++-common/goacc/loop-3.c: Likewise.
+	* g++.dg/goacc/reference.C: Likewise.
+	* g++.dg/goacc/routine-1.C: Likewise.
+	* g++.dg/goacc/routine-2.C: Likewise.
+
 2016-03-30  Richard Biener  <rguenther@suse.de>
 
 	PR middle-end/70450
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index dabba8c..aaf0e7a 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -20,12 +20,12 @@ f_acc_kernels (void)
   }
 }
 
+#pragma acc routine vector
 void
 f_acc_loop (void)
 {
   int i;
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 5e3f183..1a33242 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -1,4 +1,5 @@
 extern int i;
+#pragma acc declare create(i)
 
 void
 f_omp (void)
@@ -14,6 +15,9 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 
 #pragma omp for
@@ -358,85 +362,77 @@ f_acc_data (void)
   }
 }
 
+#pragma acc routine
 void
 f_acc_loop (void)
 {
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       for (i = 0; i < 3; i++)
 	;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       {
 	;
       }
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
     }
 
-#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
-#pragma omp target data map(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target data map(i) /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
       ;
-#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
     }
 }
 
diff --git a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index 661d364..853d010 100644
--- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -1,3 +1,5 @@
+/* Miscellaneous tests where clause parsing is expected to fail.  */
+
 void
 f (void)
 {
@@ -17,3 +19,13 @@ f (void)
   for (i = 0; i < 2; ++i)
     ;
 }
+
+
+void
+f2 (void)
+{
+  int a, b[100];
+
+#pragma acc parallel firstprivate (b[10:20]) /* { dg-error "expected ... before ... token" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
index c387285..c2a3c57 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-directives.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -1,10 +1,7 @@
-// { dg-do compile }
-// { dg-options "-fopenacc -fdump-tree-gimple" }
+// { dg-additional-options "-fdump-tree-gimple" }
 
-// This error is temporary.  Remove when support is added for these clauses
-// in the middle end.  Also remove the comments from the reduction test
+// Remove the comments from the reduction test
 // after the FE learns that reduction variables may appear in data clauses too.
-// { dg-prune-output "sorry, unimplemented" }
 
 void
 test ()
diff --git a/gcc/testsuite/c-c++-common/goacc/use_device-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
index 9a4f6d0..0c7a857 100644
--- a/gcc/testsuite/c-c++-common/goacc/use_device-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
@@ -1,4 +1,14 @@
-/* { dg-do compile } */
+/* Test valid use of host_data directive.  */
+
+int v1[3][3];
+
+void
+f (void)
+{
+#pragma acc host_data use_device(v1)
+  ;
+}
+
 
 void bar (float *, float *);
 
diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-2.c b/gcc/testsuite/c-c++-common/goacc/host_data-2.c
new file mode 100644
index 0000000..bdce424
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-2.c
@@ -0,0 +1,78 @@
+/* Test invalid use of host_data directive.  */
+
+int v0;
+#pragma acc host_data use_device(v0) /* { dg-error "expected declaration specifiers before" } */
+
+
+void
+f (void)
+{
+  int v2 = 3;
+#pragma acc host_data copy(v2) /* { dg-error ".copy. is not valid for ..pragma acc host_data." } */
+  ;
+
+#pragma acc host_data use_device(v2)
+  ;
+  /* { dg-error ".use_device_ptr. variable is neither a pointer nor an array" "" { target c } 14 } */
+  /* { dg-error ".use_device_ptr. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 14 } */
+  
+#pragma acc host_data use_device(v0)
+  ;
+  /* { dg-error ".use_device_ptr. variable is neither a pointer nor an array" "" { target c } 19 } */
+  /* { dg-error ".use_device_ptr. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 19 } */
+}
+
+
+void
+f2 (void)
+{
+  int x[100];
+
+#pragma acc enter data copyin (x)
+  /* Specifying an array index is not valid for host_data/use_device.  */
+#pragma acc host_data use_device (x[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+  ;
+#pragma acc exit data delete (x)
+}
+
+
+void
+f3 (void)
+{
+  int x[100];
+
+#pragma acc data copyin (x[25:50])
+  {
+    int *xp;
+#pragma acc host_data use_device (x)
+    {
+      /* This use of the present clause is undefined behavior for OpenACC.  */
+#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable .x. declared in enclosing .host_data. region" } */
+      {
+        xp = x;
+      }
+    }
+  }
+}
+
+
+void
+f4 (void)
+{
+  int x[50];
+
+#pragma acc data copyin (x[10:30])
+  {
+    int *xp;
+#pragma acc host_data use_device (x)
+    {
+      /* Here 'x' being implicitly firstprivate for the parallel region
+	 conflicts with it being declared as use_device in the enclosing
+	 host_data region.  */
+#pragma acc parallel copyout (xp)
+      {
+        xp = x; /* { dg-error "variable .x. declared in enclosing .host_data. region" } */
+      }
+    }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-5.c b/gcc/testsuite/c-c++-common/goacc/host_data-5.c
deleted file mode 100644
index a4206c8..0000000
--- a/gcc/testsuite/c-c++-common/goacc/host_data-5.c
+++ /dev/null
@@ -1,23 +0,0 @@
-/* { dg-do compile } */
-
-#define N 1024
-
-int main (int argc, char* argv[])
-{
-  int x[N];
-
-#pragma acc data copyin (x[0:N])
-  {
-    int *xp;
-#pragma acc host_data use_device (x)
-    {
-      /* This use of the present clause is undefined behavior for OpenACC.  */
-#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
-      {
-        xp = x;
-      }
-    }
-  }
-
-  return 0;
-}
diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-6.c b/gcc/testsuite/c-c++-common/goacc/host_data-6.c
deleted file mode 100644
index 8be7912..0000000
--- a/gcc/testsuite/c-c++-common/goacc/host_data-6.c
+++ /dev/null
@@ -1,25 +0,0 @@
-/* { dg-do compile } */
-
-#define N 1024
-
-int main (int argc, char* argv[])
-{
-  int x[N];
-
-#pragma acc data copyin (x[0:N])
-  {
-    int *xp;
-#pragma acc host_data use_device (x)
-    {
-      /* Here 'x' being implicitly firstprivate for the parallel region
-	 conflicts with it being declared as use_device in the enclosing
-	 host_data region.  */
-#pragma acc parallel copyout (xp)
-      {
-        xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
-      }
-    }
-  }
-
-  return 0;
-}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-1.c
index e91b81c..4fcf86e 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-1.c
@@ -1,6 +1,45 @@
-void
-foo (void)
+int
+kernels_empty (void)
 {
 #pragma acc kernels
   ;
+
+  return 0;
+}
+
+int
+kernels_eternal (void)
+{
+#pragma acc kernels
+  {
+    while (1)
+      ;
+  }
+
+  return 0;
+}
+
+int
+kernels_noreturn (void)
+{
+#pragma acc kernels
+  __builtin_abort ();
+
+  return 0;
+}
+
+
+float b[10][15][10];
+
+void
+kernels_loop_ptr_it (void)
+{
+  float *i;
+
+#pragma acc kernels
+  {
+#pragma acc loop
+    for (i = &b[0][0][0]; i < &b[0][0][10]; i++)
+      ;
+  }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-empty.c b/gcc/testsuite/c-c++-common/goacc/kernels-empty.c
deleted file mode 100644
index e91b81c..0000000
--- a/gcc/testsuite/c-c++-common/goacc/kernels-empty.c
+++ /dev/null
@@ -1,6 +0,0 @@
-void
-foo (void)
-{
-#pragma acc kernels
-  ;
-}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-eternal.c b/gcc/testsuite/c-c++-common/goacc/kernels-eternal.c
deleted file mode 100644
index edc17d2..0000000
--- a/gcc/testsuite/c-c++-common/goacc/kernels-eternal.c
+++ /dev/null
@@ -1,11 +0,0 @@
-int
-main (void)
-{
-#pragma acc kernels
-  {
-    while (1)
-      ;
-  }
-
-  return 0;
-}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-noreturn.c b/gcc/testsuite/c-c++-common/goacc/kernels-noreturn.c
deleted file mode 100644
index 1a8cc67..0000000
--- a/gcc/testsuite/c-c++-common/goacc/kernels-noreturn.c
+++ /dev/null
@@ -1,12 +0,0 @@
-int
-main (void)
-{
-
-#pragma acc kernels
-  {
-    __builtin_abort ();
-  }
-
-  return 0;
-}
-
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
new file mode 100644
index 0000000..01ad32d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
@@ -0,0 +1,189 @@
+void K(void)
+{
+  int i, j;
+
+#pragma acc kernels
+  {
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector 
+	for (j = 0; j < 10; j++)
+	  { }
+#pragma acc loop worker 
+	for (j = 0; j < 10; j++)
+	  { }
+#pragma acc loop gang // { dg-error "inner loop uses same" }
+	for (j = 0; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq gang // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop worker
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker(5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker(num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector 
+	for (j = 0; j < 10; j++)
+	  { }
+#pragma acc loop worker // { dg-error "inner loop uses same" }
+	for (j = 0; j < 10; j++)
+	  { }
+#pragma acc loop gang
+	for (j = 0; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq worker // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop vector
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector(5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector(length:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector // { dg-error "inner loop uses same" }
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop worker
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop gang
+	for (j = 1; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq vector // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang vector
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker vector
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq auto // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+#pragma acc kernels loop auto
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang(5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang(num:5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang(static:5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang(static:*)
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc kernels loop worker
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop worker(5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop worker(num:5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc kernels loop gang worker
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc kernels loop vector
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop vector(5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop vector(length:5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc kernels loop gang vector
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop worker vector
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc kernels loop auto
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
new file mode 100644
index 0000000..0ef5741
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
@@ -0,0 +1,162 @@
+void P(void)
+{
+  int i, j;
+
+#pragma acc parallel
+  {
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang // { dg-message "containing loop" }
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop worker 
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop gang // { dg-error "inner loop uses same" }
+	for (j = 1; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq gang // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop worker
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker // { dg-message "containing loop" 2 }
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector 
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop worker // { dg-error "inner loop uses same" }
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop gang // { dg-error "incorrectly nested" }
+	for (j = 1; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq worker // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop vector
+    for (i = 0; i < 10; i++)
+      { }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector // { dg-message "containing loop" 3 }
+    for (i = 0; i < 10; i++)
+      {
+#pragma acc loop vector // { dg-error "inner loop uses same" }
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop worker // { dg-error "incorrectly nested" }
+	for (j = 1; j < 10; j++)
+	  { }
+#pragma acc loop gang // { dg-error "incorrectly nested" }
+	for (j = 1; j < 10; j++)
+	  { }
+      }
+#pragma acc loop seq vector // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang vector
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker vector
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq auto // { dg-error "'seq' overrides" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
+    for (i = 0; i < 10; i++)
+      { }
+
+  }
+
+#pragma acc parallel loop auto
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop gang
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop gang(static:5)
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop gang(static:*)
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+
+#pragma acc parallel loop worker
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc parallel loop gang worker
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc parallel loop vector
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc parallel loop gang vector
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop worker vector
+  for (i = 0; i < 10; i++)
+    { }
+
+#pragma acc parallel loop auto
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
+    { }
+#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
+    { }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-3.c b/gcc/testsuite/c-c++-common/goacc/loop-3.c
new file mode 100644
index 0000000..44b65a8
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-3.c
@@ -0,0 +1,58 @@
+void par1 (void)
+{
+  int i, j;
+
+#pragma acc parallel
+  {
+#pragma acc loop gang(5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop gang(num:5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop worker(5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop worker(num:5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop vector(5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+#pragma acc loop vector(length:5) // { dg-error "argument not permitted" }
+    for (i = 0; i < 10; i++)
+      { }
+
+   }
+}
+
+void p2 (void)
+{
+  int i, j;
+
+#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+
+#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+
+#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" "" { target c } }
+  for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
+    { }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
index 97b8786..f3c7207 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
@@ -1,7 +1,3 @@
-/* { dg-do compile } */
-
-/* { dg-prune-output "sorry, unimplemented" } */
-
 int
 main ()
 {
diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c
index 3a8f838..cab4f98 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c
@@ -58,10 +58,6 @@ f_acc_data (void)
 
 #pragma acc exit data delete(i)
 
-#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
-    for (i = 0; i < 2; ++i)
-      ;
-
 #pragma acc data
     {
 #pragma acc parallel
@@ -92,10 +88,6 @@ f_acc_data (void)
 #pragma acc enter data copyin(i)
 
 #pragma acc exit data delete(i)
-
-#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
-      for (i = 0; i < 2; ++i)
-	;
     }
   }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 506a1ae..93a9111 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -38,6 +38,25 @@ f_acc_kernels (void)
   }
 }
 
+void
+f_acc_data (void)
+{
+  unsigned int i;
+#pragma acc data
+  {
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
+    for (i = 0; i < 2; ++i)
+      ;
+
+#pragma acc data
+    {
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
+      for (i = 0; i < 2; ++i)
+	;
+    }
+  }
+}
+
 #pragma acc routine
 void
 f_acc_routine (void)
@@ -45,3 +64,13 @@ f_acc_routine (void)
 #pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */
   ;
 }
+
+void
+f (void)
+{
+  int i, v = 0;
+
+#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
+  for (i = 0; i < 10; i++)
+    v++;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-1.c
index a860526..6c6cc88 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-1.c
@@ -1,6 +1,38 @@
-void
-foo (void)
+int
+parallel_empty (void)
 {
 #pragma acc parallel
   ;
+
+  return 0;
+}
+
+int
+parallel_eternal (void)
+{
+#pragma acc parallel
+  {
+    while (1)
+      ;
+  }
+
+  return 0;
+}
+
+int
+parallel_noreturn (void)
+{
+#pragma acc parallel
+  __builtin_abort ();
+
+  return 0;
+}
+
+int
+parallel_clauses (void)
+{
+  int a, b[100];
+
+#pragma acc parallel firstprivate (a, b)
+  ;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
index de97125b..3c1c2dd 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
@@ -1,70 +1,65 @@
-/* { dg-require-effective-target alloca } */
 /* Integer reductions.  */
 
-#define vl 32
+#define n 1000
 
 int
 main(void)
 {
-  const int n = 1000;
   int i;
   int result, array[n];
   int lresult;
 
   /* '+' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (+:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (+:result)
   for (i = 0; i < n; i++)
     result += array[i];
 
   /* '*' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (*:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (*:result)
   for (i = 0; i < n; i++)
     result *= array[i];
 
-//   result = 0;
-//   vresult = 0;
-// 
-//   /* 'max' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result > array[i] ? result : array[i];
-//
-//   /* 'min' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result < array[i] ? result : array[i];
+  /* 'max' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (max:result)
+  for (i = 0; i < n; i++)
+    result = result > array[i] ? result : array[i];
+
+  /* 'min' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (min:result)
+  for (i = 0; i < n; i++)
+    result = result < array[i] ? result : array[i];
 
   /* '&' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (&:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (&:result)
   for (i = 0; i < n; i++)
     result &= array[i];
 
   /* '|' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (|:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (|:result)
   for (i = 0; i < n; i++)
     result |= array[i];
 
   /* '^' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (^:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (^:result)
   for (i = 0; i < n; i++)
     result ^= array[i];
 
   /* '&&' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (&&:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (&&:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult && (result > array[i]);
 
   /* '||' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (||:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (||:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult || (result > array[i]);
 
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
index 2964236..c3105a2 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
@@ -1,49 +1,47 @@
-/* { dg-require-effective-target alloca } */
 /* float reductions.  */
 
-#define vl 32
+#define n 1000
 
 int
 main(void)
 {
-  const int n = 1000;
   int i;
   float result, array[n];
   int lresult;
 
   /* '+' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (+:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (+:result)
   for (i = 0; i < n; i++)
     result += array[i];
 
   /* '*' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (*:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (*:result)
   for (i = 0; i < n; i++)
     result *= array[i];
 
-//   /* 'max' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result > array[i] ? result : array[i];
-// 
-//   /* 'min' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result < array[i] ? result : array[i];
+  /* 'max' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (max:result)
+  for (i = 0; i < n; i++)
+    result = result > array[i] ? result : array[i];
+
+  /* 'min' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (min:result)
+  for (i = 0; i < n; i++)
+    result = result < array[i] ? result : array[i];
 
   /* '&&' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (&&:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (&&:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult && (result > array[i]);
 
   /* '||' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (||:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (||:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult || (result > array[i]);
 
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
index 34c51c2..4dbde04 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
@@ -1,49 +1,47 @@
-/* { dg-require-effective-target alloca } */
 /* double reductions.  */
 
-#define vl 32
+#define n 1000
 
 int
 main(void)
 {
-  const int n = 1000;
   int i;
   double result, array[n];
   int lresult;
 
   /* '+' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (+:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (+:result)
   for (i = 0; i < n; i++)
     result += array[i];
 
   /* '*' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (*:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (*:result)
   for (i = 0; i < n; i++)
     result *= array[i];
 
-//   /* 'max' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result > array[i] ? result : array[i];
-// 
-//   /* 'min' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result < array[i] ? result : array[i];
+  /* 'max' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (max:result)
+  for (i = 0; i < n; i++)
+    result = result > array[i] ? result : array[i];
+
+  /* 'min' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (min:result)
+  for (i = 0; i < n; i++)
+    result = result < array[i] ? result : array[i];
 
   /* '&&' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (&&:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (&&:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult && (result > array[i]);
 
   /* '||' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (||:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (||:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult || (result > array[i]);
 
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
index 328c0d4..c4572b9 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
@@ -1,51 +1,35 @@
-/* { dg-require-effective-target alloca } */
 /* complex reductions.  */
 
-#define vl 32
+#define n 1000
 
 int
 main(void)
 {
-  const int n = 1000;
   int i;
   __complex__ double result, array[n];
   int lresult;
 
   /* '+' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (+:result)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (+:result)
   for (i = 0; i < n; i++)
     result += array[i];
 
-  /* Needs support for complex multiplication.  */
-
-//   /* '*' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (*:result)
-//   for (i = 0; i < n; i++)
-//     result *= array[i];
-//
-//   /* 'max' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result > array[i] ? result : array[i];
-// 
-//   /* 'min' reductions.  */
-// #pragma acc parallel vector_length (vl)
-// #pragma acc loop reduction (+:result)
-//   for (i = 0; i < n; i++)
-//       result = result < array[i] ? result : array[i];
+  /* '*' reductions.  */
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (*:result)
+  for (i = 0; i < n; i++)
+    result *= array[i];
 
   /* '&&' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (&&:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (&&:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult && (__real__(result) > __real__(array[i]));
 
   /* '||' reductions.  */
-#pragma acc parallel vector_length (vl)
-#pragma acc loop reduction (||:lresult)
+#pragma acc parallel
+#pragma acc loop gang worker vector reduction (||:lresult)
   for (i = 0; i < n; i++)
     lresult = lresult || (__real__(result) > __real__(array[i]));
 
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-3.c b/gcc/testsuite/c-c++-common/goacc/routine-3.c
index e6f83bd..b322d26 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-3.c
@@ -1,52 +1,118 @@
+/* Test invalid calls to routines.  */
+
 #pragma acc routine gang
-void gang (void) /* { dg-message "declared here" 3 } */
+int
+gang () /* { dg-message "declared here" 3 } */
 {
+  #pragma acc loop gang worker vector
+  for (int i = 0; i < 10; i++)
+    {
+    }
+
+  return 1;
 }
 
 #pragma acc routine worker
-void worker (void) /* { dg-message "declared here" 2 } */
+int
+worker () /* { dg-message "declared here" 2 } */
 {
+  #pragma acc loop worker vector
+  for (int i = 0; i < 10; i++)
+    {
+    }
+
+  return 1;
 }
 
 #pragma acc routine vector
-void vector (void) /* { dg-message "declared here" 1 } */
+int
+vector () /* { dg-message "declared here" } */
 {
+  #pragma acc loop vector
+  for (int i = 0; i < 10; i++)
+    {
+    }
+
+  return 1;
 }
 
 #pragma acc routine seq
-void seq (void)
+int
+seq ()
 {
+  return 1;
 }
 
-int main ()
+int
+main ()
 {
-
-#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
+  int red = 0;
+#pragma acc parallel copy (red)
   {
-    #pragma acc loop gang /* { dg-message "loop here" 1 } */
-    for (int i = 0; i < 10; i++)
-      {
-	gang (); /*  { dg-error "routine call uses same" } */
-	worker ();
-	vector ();
-	seq ();
-      }
-    #pragma acc loop worker /* { dg-message "loop here" 2 } */
-    for (int i = 0; i < 10; i++)
-      {
-	gang (); /*  { dg-error "routine call uses same" } */
-	worker (); /*  { dg-error "routine call uses same" } */
-	vector ();
-	seq ();
-      }
-    #pragma acc loop vector /* { dg-message "loop here" 3 } */
-    for (int i = 0; i < 10; i++)
-      {
-	gang (); /*  { dg-error "routine call uses same" } */
-	worker (); /*  { dg-error "routine call uses same" } */
-	vector (); /*  { dg-error "routine call uses same" } */
-	seq ();
-      }
+    /* Independent/seq loop tests.  */
+#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning" }
+    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)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "routine call uses same" }
+
+#pragma acc loop worker reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "routine call uses same" }
+
+#pragma acc loop vector reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += gang (); // { dg-error "routine call uses same" }
+
+    /* Worker routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += worker ();
+
+#pragma acc loop worker reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += worker (); // { dg-error "routine call uses same" }
+
+#pragma acc loop vector reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += worker (); // { dg-error "routine call uses same" }
+
+    /* 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)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += vector (); // { dg-error "routine call uses same" }
+
+    /* 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-4.c b/gcc/testsuite/c-c++-common/goacc/routine-4.c
index 004d713..3e5fc4f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c
@@ -1,3 +1,4 @@
+/* Test invalid intra-routine parallelism.  */
 
 void gang (void);
 void worker (void);
@@ -14,6 +15,24 @@ void seq (void)
   worker ();  /* { dg-error "routine call uses" } */
   vector ();  /* { dg-error "routine call uses" } */
   seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
 }
 
 void vector (void) /* { dg-message "declared here" 1 } */
@@ -22,6 +41,24 @@ void vector (void) /* { dg-message "declared here" 1 } */
   worker ();  /* { dg-error "routine call uses" } */
   vector ();
   seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
 }
 
 void worker (void) /* { dg-message "declared here" 2 } */
@@ -30,6 +67,24 @@ void worker (void) /* { dg-message "declared here" 2 } */
   worker ();
   vector ();
   seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing 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 ++;
 }
 
 void gang (void) /* { dg-message "declared here" 3 } */
@@ -38,4 +93,22 @@ void gang (void) /* { dg-message "declared here" 3 } */
   worker ();
   vector ();
   seq ();
+
+  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 ++;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index c34838f..2a9db90 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -46,6 +46,21 @@ using namespace g;
   
 #pragma acc routine (c) /* { dg-error "does not refer to" } */
 
+
+void Bar ();
+
+void Foo ()
+{
+  Bar ();
+}
+
+#pragma acc routine (Bar) // { dg-error "must be applied before use" }
+
+#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
+
+#pragma acc routine (Baz) // { dg-error "not been declared" }
+
+
 int vb1;		/* { dg-error "directive for use" } */
 extern int vb2;		/* { dg-error "directive for use" } */
 static int vb3;		/* { dg-error "directive for use" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c
index 2a81427..8e70e71 100644
--- a/gcc/testsuite/c-c++-common/goacc/tile.c
+++ b/gcc/testsuite/c-c++-common/goacc/tile.c
@@ -1,5 +1,3 @@
-/* { dg-do compile } */
-
 int
 main ()
 {
@@ -71,3 +69,259 @@ main ()
 
   return 0;
 }
+
+
+void par (void)
+{
+  int i, j;
+
+#pragma acc parallel
+  {
+#pragma acc loop tile // { dg-error "expected" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile() // { dg-error "expected" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(1) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(2) 
+    for (i = 0; i < 10; i++)
+      {
+	for (j = 1; j < 10; j++)
+	  { }
+      }
+#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
+    for (i = 1; i < 10; i++)
+      { }
+#pragma acc loop tile(i)
+    for (i = 1; i < 10; i++)
+      { }
+#pragma acc loop tile(2, 2, 1)
+    for (i = 1; i < 3; i++)
+      {
+	for (j = 4; j < 6; j++)
+	  { }
+      } 
+#pragma acc loop tile(2, 2)
+    for (i = 1; i < 5; i+=2)
+      {
+	for (j = i + 1; j < 7; j+=i)
+	  { }
+      }
+#pragma acc loop vector tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector gang tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+  }
+}
+void p3 (void)
+{
+  int i, j;
+
+  
+#pragma acc parallel loop tile // { dg-error "expected" }
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop tile() // { dg-error "expected" }
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop tile(1) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop tile(*, 1) 
+  for (i = 0; i < 10; i++)
+    {
+      for (j = 1; j < 10; j++)
+	{ }
+    }
+#pragma acc parallel loop tile(-2) // { dg-warning "'tile' value must be positive" }
+  for (i = 1; i < 10; i++)
+    { }
+#pragma acc parallel loop tile(i)
+  for (i = 1; i < 10; i++)
+    { }
+#pragma acc parallel loop tile(2, 2, 1)
+  for (i = 1; i < 3; i++)
+    {
+      for (j = 4; j < 6; j++)
+        { }
+    }    
+#pragma acc parallel loop tile(2, 2)
+  for (i = 1; i < 5; i+=2)
+    {
+      for (j = i + 1; j < 7; j++)
+        { }
+    }
+#pragma acc parallel loop vector tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop gang tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop vector gang tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop vector worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc parallel loop gang worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+
+}
+
+
+void
+kern (void)
+{
+  int i, j;
+
+#pragma acc kernels
+  {
+#pragma acc loop tile // { dg-error "expected" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile() // { dg-error "expected" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(1)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(2)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(6-2) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(6+2) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(*, 1) 
+    for (i = 0; i < 10; i++)
+      {
+	for (j = 0; j < 10; i++)
+	  { }
+      }
+#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(i)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop tile(2, 2, 1)
+    for (i = 2; i < 4; i++)
+      for (i = 4; i < 6; i++)
+	{ }
+#pragma acc loop tile(2, 2)
+    for (i = 1; i < 5; i+=2)
+      for (j = i+1; j < 7; i++)
+	{ }
+#pragma acc loop vector tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector gang tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker tile(*) 
+    for (i = 0; i < 10; i++)
+      { }
+   }
+}
+
+
+void k3 (void)
+{
+  int i, j;
+
+#pragma acc kernels loop tile // { dg-error "expected" }
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop tile() // { dg-error "expected" }
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop tile(1) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop tile(*, 1) 
+  for (i = 0; i < 10; i++)
+    {
+      for (j = 1; j < 10; j++)
+	{ }
+    }
+#pragma acc kernels loop tile(-2) // { dg-warning "'tile' value must be positive" }
+  for (i = 1; i < 10; i++)
+    { }
+#pragma acc kernels loop tile(i)
+  for (i = 1; i < 10; i++)
+    { }
+#pragma acc kernels loop tile(2, 2, 1)
+  for (i = 1; i < 3; i++)
+    {
+      for (j = 4; j < 6; j++)
+	{ }
+    }    
+#pragma acc kernels loop tile(2, 2)
+  for (i = 1; i < 5; i++)
+    {
+      for (j = i + 1; j < 7; j += i)
+	{ }
+    }
+#pragma acc kernels loop vector tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop vector gang tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop vector worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+#pragma acc kernels loop gang worker tile(*) 
+  for (i = 0; i < 10; i++)
+    { }
+}
diff --git a/gcc/testsuite/g++.dg/goacc/reference.C b/gcc/testsuite/g++.dg/goacc/reference.C
new file mode 100644
index 0000000..b000668
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/reference.C
@@ -0,0 +1,39 @@
+int
+test1 (int &ref)
+{
+#pragma acc kernels copy (ref)
+  {
+    ref = 10;
+  }
+}
+
+int
+test2 (int &ref)
+{
+  int b;
+#pragma acc kernels copyout (b)
+  {
+    b = ref + 10;
+  }
+
+#pragma acc parallel copyout (b)
+  {
+    b = ref + 10;
+  }
+
+  ref = b;
+}
+
+int
+main()
+{
+  int a = 0;
+  int &ref_a = a;
+
+  #pragma acc parallel copy (a, ref_a)
+  {
+    ref_a = 5;
+  }
+
+  return a;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/routine-1.C b/gcc/testsuite/g++.dg/goacc/routine-1.C
new file mode 100644
index 0000000..a73a73d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/routine-1.C
@@ -0,0 +1,13 @@
+/* Test valid use of the routine directive.  */
+
+namespace N
+{
+  extern void foo1();
+  extern void foo2();
+#pragma acc routine (foo1)
+#pragma acc routine
+  void foo3()
+  {
+  }
+}
+#pragma acc routine (N::foo2)
diff --git a/gcc/testsuite/g++.dg/goacc/routine-2.C b/gcc/testsuite/g++.dg/goacc/routine-2.C
new file mode 100644
index 0000000..2d16466
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/routine-2.C
@@ -0,0 +1,42 @@
+/* Test invalid use of the routine directive.  */
+
+template <typename T>
+extern T one_d();
+#pragma acc routine (one_d) /* { dg-error "names a set of overloads" } */
+
+template <typename T>
+T
+one()
+{
+  return 1;
+}
+#pragma acc routine (one) /* { dg-error "names a set of overloads" } */
+
+int incr (int);
+float incr (float);
+int inc;
+
+#pragma acc routine (incr) /* { dg-error "names a set of overloads" } */
+
+#pragma acc routine (increment) /* { dg-error "has not been declared" } */
+
+#pragma acc routine (inc) /* { dg-error "does not refer to a function" } */
+
+#pragma acc routine (+) /* { dg-error "expected unqualified-id before '.' token" } */
+
+int sum (int, int);
+
+namespace foo {
+#pragma acc routine (sum)
+  int sub (int, int);
+}
+
+#pragma acc routine (foo::sub)
+
+/* It's strange to apply a routine directive to subset of overloaded
+   functions, but that is permissible in OpenACC 2.x.  */
+
+int decr (int a);
+
+#pragma acc routine
+float decr (float a);
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C
index f7a717b..f139dc2 100644
--- a/gcc/testsuite/g++.dg/goacc/template.C
+++ b/gcc/testsuite/g++.dg/goacc/template.C
@@ -1,8 +1,3 @@
-// This error is temporary.  Remove when support is added for these clauses
-// in the middle end.  Also remove the comments from the reduction test
-// after the FE learns that reduction variables may appear in data clauses too.
-// { dg-prune-output "sorry, unimplemented" }
-
 #pragma acc routine
 template <typename T> T
 accDouble(int val)
@@ -20,55 +15,62 @@ oacc_parallel_copy (T a)
   double z = 4;
 
 #pragma acc parallel num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
-  {
+#pragma acc loop gang worker vector
+  for (int i = 0; i < 1; i++)
     b = a;
-  }
 
 #pragma acc parallel num_gangs (a) copy (w, x, y, z)
-  {
-    w = accDouble<char>(w);
-    x = accDouble<int>(x);
-    y = accDouble<float>(y);
-    z = accDouble<double>(z);
-  }
+#pragma acc loop
+  for (int i = 0; i < 1; i++)
+    {
+      w = accDouble<char>(w);
+      x = accDouble<int>(x);
+      y = accDouble<float>(y);
+      z = accDouble<double>(z);
+    }
 
 #pragma acc parallel num_gangs (a) if (1)
   {
+#pragma acc loop independent collapse (2) gang
+    for (int i = 0; i < a; i++)
+      for (int j = 0; j < 5; j++)
+	b = a;
+
 #pragma acc loop auto tile (a, 3)
-  for (int i = 0; i < a; i++)
-    for (int j = 0; j < 5; j++)
-      b = a;
+    for (int i = 0; i < a; i++)
+      for (int j = 0; j < 5; j++)
+	b = a;
 
 #pragma acc loop seq
-  for (int i = 0; i < a; i++)
-    b = a;
+    for (int i = 0; i < a; i++)
+      b = a;
   }
 
   T c;
 
 #pragma acc parallel num_workers (10)
-  {
+#pragma acc loop worker
+  for (int i = 0; i < 1; i++)
+    {
 #pragma acc atomic capture
-    c = b++;
+      c = b++;
 
 #pragma atomic update
-    c++;
+      c++;
 
 #pragma acc atomic read
-    b = a;
+      b = a;
 
 #pragma acc atomic write
-    b = a;
-  }
+      b = a;
+    }
 
-//#pragma acc parallel reduction (+:c)
-//  {
-//    c = 1;
-//  }
+#pragma acc parallel reduction (+:c)
+  c = 1;
 
 #pragma acc data if (1) copy (b)
   {
-    #pragma acc parallel
+#pragma acc parallel
     {
       b = a;
     }
@@ -76,9 +78,9 @@ oacc_parallel_copy (T a)
 
 #pragma acc enter data copyin (b)
 #pragma acc parallel present (b)
-    {
-      b = a;
-    }
+  {
+    b = a;
+  }
 
 #pragma acc update host (b)
 #pragma acc update self (b)
@@ -109,11 +111,9 @@ oacc_kernels_copy (T a)
 #pragma acc kernels copyout (b) copyin (a)
   b = a;
 
-//#pragma acc kernels loop reduction (+:c)
-//  for (int i = 0; i < 10; i++)
-//    {
-//      c = 1;
-//    }
+#pragma acc kernels loop reduction (+:c)
+  for (int i = 0; i < 10; i++)
+    c = 1;
 
 #pragma acc data if (1) copy (b)
   {
@@ -125,9 +125,10 @@ oacc_kernels_copy (T a)
 
 #pragma acc enter data copyin (b)
 #pragma acc kernels present (b)
-    {
-      b = a;
-    }
+  {
+    b = a;
+  }
+
   return b;
 }
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 6977525..42a447a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -1,17 +1,10 @@
 ! Exercise combined OpenACC directives.
 
-! { dg-do compile }
-! { dg-options "-fopenacc -fdump-tree-gimple" }
-
-! This error is temporary.  Remove when support is added for these clauses
-! in the middle end.
-! { dg-prune-output "sorry, unimplemented" }
-
-! Update the reduction tests.
+! { dg-additional-options "-fdump-tree-gimple" }
 
 subroutine test
   implicit none
-  integer a(100), i, j, z
+  integer a(100), i, j, y, z
 
   ! PARALLEL
   
@@ -73,10 +66,10 @@ subroutine test
   end do
   !$acc end parallel loop
 
-!  !$acc parallel loop reduction (+:z) copy (z)
-!  do i = 1, 100
-!  end do
-!  !$acc end parallel loop
+  !$acc parallel loop reduction (+:y) copy (y)
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
 
   ! KERNELS
 
@@ -138,10 +131,10 @@ subroutine test
   end do
   !$acc end kernels loop
 
-!  !$acc kernels loop reduction (+:z) copy (z)
-!  do i = 1, 100
-!  end do
-!  !$acc end kernels loop
+  !$acc kernels loop reduction (+:y) copy (y)
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
 end subroutine test
 
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. collapse.2." 2 "gimple" } }
@@ -153,3 +146,5 @@ end subroutine test
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
 ! { 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" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-1.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
index 817039f..b5f9e03 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
@@ -1,5 +1,3 @@
-! { dg-do compile } 
-! { dg-additional-options "-fmax-errors=100" } 
 module test
   implicit none
 contains
@@ -29,14 +27,18 @@ subroutine test1
        i = i + 1
   end do
   !$acc loop
-  do 300 d = 1, 30, 6 ! { dg-error "integer" }
+  do 300 d = 1, 30, 6
       i = d
   300 a(i) = 1
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 30 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 30 }
   !$acc loop
-  do d = 1, 30, 5 ! { dg-error "integer" }
+  do d = 1, 30, 5
        i = d
       a(i) = 2
   end do
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 36 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 36 }
   !$acc loop
   do i = 1, 30
       if (i .eq. 16) exit ! { dg-error "EXIT statement" }
@@ -144,8 +146,10 @@ subroutine test1
     end do
     !$acc parallel loop collapse(2)
     do i = 1, 3
-        do r = 4, 6    ! { dg-error "integer" }
+        do r = 4, 6
         end do
+        ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 149 }
+        ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 149 }
     end do
 
     ! Both seq and independent are not allowed
@@ -167,4 +171,3 @@ subroutine test1
 
 end subroutine test1
 end module test
-! { dg-prune-output "Deleted" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
index 5cbd975..d059cf7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -1,9 +1,3 @@
-! { dg-do compile }
-! { dg-additional-options "-fmax-errors=100" }
-
-! { dg-prune-output "sorry, unimplemented" }
-! { dg-prune-output "Error: work-sharing region" }
-
 program test
   implicit none
   integer :: i, j
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
index e844468..d0855b4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
@@ -1,11 +1,3 @@
-! { dg-do compile }
-! { dg-additional-options "-fmax-errors=100" }
-
-! This error is temporary.  Remove when support is added for these clauses
-! in the middle end.
-! { dg-prune-output "sorry, unimplemented" }
-! { dg-prune-output "Error: work-sharing region" }
-
 program test
   implicit none
   integer :: i, j
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index 6cfd715..81bdc23 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -1,13 +1,7 @@
-! { dg-do compile } 
 ! { dg-additional-options "-fdump-tree-original -std=f2008" } 
 
 ! test for tree-dump-original and spaces-commas
 
-! This error is temporary.  Remove when support is added for these clauses
-! in the middle end.
-! { dg-prune-output "sorry, unimplemented" }
-! { dg-prune-output "Error: work-sharing region" }
-
 program test
   implicit none
   integer :: i, j, k, m, sum