diff mbox

[openacc] acc loop updates in fortran

Message ID 5639764A.2000209@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Nov. 4, 2015, 3:06 a.m. UTC
This patch updates the fortran front end so that it supports the acc
loop clauses in a similar manner to both the c and c++ front ends in
addition to addressing a couple of other loose ends. Specifically:

 * the tile clause now supports a list of arguments like c and c++

 * made the default clause parser aware that openacc only supports
  'none'

 * error when no data clauses are used in an update construct

 * corrects the way that reduction errors are detected for duplicate
   variables

 * add support for the gang static argument and the auto clause

Besides for that, a number of tests have either been updated or added.

Is this OK for trunk?

Cesar

Comments

Jakub Jelinek Nov. 4, 2015, 10:30 a.m. UTC | #1
On Tue, Nov 03, 2015 at 07:06:50PM -0800, Cesar Philippidis wrote:
> @@ -856,13 +857,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>        if ((mask & OMP_CLAUSE_DEFAULT)
>  	  && c->default_sharing == OMP_DEFAULT_UNKNOWN)
>  	{
> -	  if (gfc_match ("default ( shared )") == MATCH_YES)
> +	  if (!openacc && gfc_match ("default ( shared )") == MATCH_YES)
>  	    c->default_sharing = OMP_DEFAULT_SHARED;
> -	  else if (gfc_match ("default ( private )") == MATCH_YES)
> +	  else if (!openacc && gfc_match ("default ( private )") == MATCH_YES)
>  	    c->default_sharing = OMP_DEFAULT_PRIVATE;
>  	  else if (gfc_match ("default ( none )") == MATCH_YES)
>  	    c->default_sharing = OMP_DEFAULT_NONE;
> -	  else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
> +	  else if (!openacc
> +		   && gfc_match ("default ( firstprivate )") == MATCH_YES)
>  	    c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;

I'd rewrite this as:
	  if (gfc_match ("default ( none )") == MATCH_YES)
	    c->default_sharing = OMP_DEFAULT_NONE;
	  else if (openacc)
	    /* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
	  else if (gfc_match ("default ( shared )") == MATCH_YES)
	    ...

>  	  if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
>  	    continue;
> @@ -1304,10 +1306,19 @@ match
>  gfc_match_oacc_update (void)
>  {
>    gfc_omp_clauses *c;
> +  locus here = gfc_current_locus;
> +
>    if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
>        != MATCH_YES)
>      return MATCH_ERROR;
>  
> +  if (!c->lists[OMP_LIST_MAP])
> +    {
> +      gfc_error ("%<acc update%> must contain at least one "
> +		 "%<device%> or %<host/self%> clause at %L", &here);

There is no host/self clause I'd guess, so you should spell those
separately.

Otherwise, LGTM.

	Jakub
Thomas Schwinge Nov. 4, 2015, 5:15 p.m. UTC | #2
Hi Cesar!

On Tue, 3 Nov 2015 19:06:50 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch updates the fortran front end so that it supports the acc
> loop clauses in a similar manner to both the c and c++ front ends in
> addition to addressing a couple of other loose ends.

> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c

> @@ -3028,6 +3015,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  	n->sym->mark = 1;
>      }
>  
> +  /* OpenACC reductions.  */
> +  if (openacc)
> +    {
> +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> +	n->sym->mark = 0;

Maybe I'm just confugsed, but if setting all these to zero here...

> +
> +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> +	{
> +	  if (n->sym->mark)
> +	    gfc_error ("Symbol %qs present on multiple clauses at %L",
> +		       n->sym->name, &n->where);
> +	  else
> +	    n->sym->mark = 1;

... won't this just always run into the "else" branch?

> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c

> @@ -3449,16 +3478,28 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
>  	      sizeof (construct_clauses));
>        loop_clauses.collapse = construct_clauses.collapse;
> [...]
> -      construct_clauses.collapse = 0;

Again I'm confused by this, why this is being removed, as earlier in
<http://news.gmane.org/find-root.php?message_id=%3C87vb9r45qw.fsf%40kepler.schwinge.homeip.net%3E>.

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90

I suggest you also merge the existing
gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 into your new test case
file (consistent naming, with the other combined-directives* files).

> @@ -0,0 +1,152 @@
> +! Exercise combined OpenACC directives.
> +
> +! { dg-do compile }
> +! { dg-options "-fopenacc -fdump-tree-gimple" }
> +
> +! { dg-prune-output "sorry, unimplemented" }

What's still unimplemented here?  Please add a comment, or put the
dg-prune-output directive next to the offending OpenACC directive, so
we'll be sure to remove it later on.

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
> @@ -0,0 +1,363 @@
> +! { dg-do compile }
> +! { dg-additional-options "-fmax-errors=100" }
> +
> +! { dg-prune-output "sorry, unimplemented" }

Likewise.

> +! { dg-prune-output "Error: work-sharing region" }

What's the intention of this?  If we're expecting this error, place
dg-error directives where they belong?

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
> @@ -0,0 +1,80 @@
> +! { dg-do compile }
> +! { dg-additional-options "-fmax-errors=100" }
> +
> +! { dg-prune-output "sorry, unimplemented" }

Likewise.

> +! { dg-prune-output "Error: work-sharing region" }

Likewise.

> --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
> @@ -3,6 +3,9 @@
>  
>  ! test for tree-dump-original and spaces-commas
>  
> +! { dg-prune-output "sorry, unimplemented" }

Likewise.

> +! { dg-prune-output "Error: work-sharing region" }

Likewise.

> --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
> +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
> @@ -37,4 +37,3 @@ end program test
>  
>  ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
>  ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } 
> -! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } } 

Which of your source code changes does this change related to?


Grüße
 Thomas
Thomas Schwinge Nov. 4, 2015, 5:20 p.m. UTC | #3
Hi Jakub!

On Wed, 4 Nov 2015 11:30:28 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> >  gfc_match_oacc_update (void)
> >  {
> >    gfc_omp_clauses *c;
> > +  locus here = gfc_current_locus;
> > +
> >    if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
> >        != MATCH_YES)
> >      return MATCH_ERROR;
> >  
> > +  if (!c->lists[OMP_LIST_MAP])
> > +    {
> > +      gfc_error ("%<acc update%> must contain at least one "
> > +		 "%<device%> or %<host/self%> clause at %L", &here);
> 
> There is no host/self clause I'd guess, so you should spell those
> separately.

That's the same language as used in the C and C++ front ends; the host
and self clauses are synonymous.

    $ git grep -C2 host/self upstream/trunk -- gcc/c*/
    upstream/trunk:gcc/c/c-parser.c-      error_at (loc,
    upstream/trunk:gcc/c/c-parser.c-                "%<#pragma acc update%> must contain at least one "
    upstream/trunk:gcc/c/c-parser.c:                "%<device%> or %<host/self%> clause");
    upstream/trunk:gcc/c/c-parser.c-      return;
    upstream/trunk:gcc/c/c-parser.c-    }
    --
    upstream/trunk:gcc/cp/parser.c-      error_at (pragma_tok->location,
    upstream/trunk:gcc/cp/parser.c-         "%<#pragma acc update%> must contain at least one "
    upstream/trunk:gcc/cp/parser.c:         "%<device%> or %<host/self%> clause");
    upstream/trunk:gcc/cp/parser.c-      return NULL_TREE;
    upstream/trunk:gcc/cp/parser.c-    }


Grüße
 Thomas
Jakub Jelinek Nov. 4, 2015, 5:21 p.m. UTC | #4
On Wed, Nov 04, 2015 at 06:15:14PM +0100, Thomas Schwinge wrote:
> > --- a/gcc/fortran/openmp.c
> > +++ b/gcc/fortran/openmp.c
> 
> > @@ -3028,6 +3015,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
> >  	n->sym->mark = 1;
> >      }
> >  
> > +  /* OpenACC reductions.  */
> > +  if (openacc)
> > +    {
> > +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> > +	n->sym->mark = 0;
> 
> Maybe I'm just confugsed, but if setting all these to zero here...
> 
> > +
> > +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> > +	{
> > +	  if (n->sym->mark)
> > +	    gfc_error ("Symbol %qs present on multiple clauses at %L",
> > +		       n->sym->name, &n->where);
> > +	  else
> > +	    n->sym->mark = 1;
> 
> ... won't this just always run into the "else" branch?

The point is to check if some symbol isn't present multiple times in
reduction clause(s).  So the first loop clears the flag as it could have
arbitrary values, and the second loop will diagnose an error if n->sym
is present multiple times in the list.  reduction(+: a, b, a) and the like.
In C/C++ FE we use bitmaps for this, in Fortran FE we have mark
field for those purposes.

	Jakub
Jakub Jelinek Nov. 4, 2015, 5:22 p.m. UTC | #5
On Wed, Nov 04, 2015 at 06:20:06PM +0100, Thomas Schwinge wrote:
> Hi Jakub!
> 
> On Wed, 4 Nov 2015 11:30:28 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > >  gfc_match_oacc_update (void)
> > >  {
> > >    gfc_omp_clauses *c;
> > > +  locus here = gfc_current_locus;
> > > +
> > >    if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
> > >        != MATCH_YES)
> > >      return MATCH_ERROR;
> > >  
> > > +  if (!c->lists[OMP_LIST_MAP])
> > > +    {
> > > +      gfc_error ("%<acc update%> must contain at least one "
> > > +		 "%<device%> or %<host/self%> clause at %L", &here);
> > 
> > There is no host/self clause I'd guess, so you should spell those
> > separately.
> 
> That's the same language as used in the C and C++ front ends; the host
> and self clauses are synonymous.

Then it should be %<host%>/%<self%> at least, or better
%<host%> or %<self%>.
Both in Fortran and in the C/C++ FEs.

	Jakub
Cesar Philippidis Nov. 4, 2015, 5:39 p.m. UTC | #6
On 11/04/2015 09:15 AM, Thomas Schwinge wrote:

>> --- a/gcc/fortran/trans-openmp.c
>> +++ b/gcc/fortran/trans-openmp.c
> 
>> @@ -3449,16 +3478,28 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
>>  	      sizeof (construct_clauses));
>>        loop_clauses.collapse = construct_clauses.collapse;
>> [...]
>> -      construct_clauses.collapse = 0;
> 
> Again I'm confused by this, why this is being removed, as earlier in
> <http://news.gmane.org/find-root.php?message_id=%3C87vb9r45qw.fsf%40kepler.schwinge.homeip.net%3E>.

I'm not sure why, but gfc_trans_omp_do needs it. It's probably an openmp
thing. If you look at gfc_trans_omp_do, you'll see that two sets of
clauses are passed into it. code->ext.omp_clauses corresponds to the
combined construct clauses and do_clauses are the filtered out ones.

So in order to get collapse to work as expected in combined loops, I
can't zero out construct_clauses.collapse.

>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
> 
> I suggest you also merge the existing
> gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 into your new test case
> file (consistent naming, with the other combined-directives* files).

OK, but it depends on what type of things combined_loop.f90 is checking.
If it's scanning gimple, it may have to be a separate file.

>> @@ -0,0 +1,152 @@
>> +! Exercise combined OpenACC directives.
>> +
>> +! { dg-do compile }
>> +! { dg-options "-fopenacc -fdump-tree-gimple" }
>> +
>> +! { dg-prune-output "sorry, unimplemented" }
> 
> What's still unimplemented here?  Please add a comment, or put the
> dg-prune-output directive next to the offending OpenACC directive, so
> we'll be sure to remove it later on.

I was still seeing those sorry messages. I'll put a comment on them.

>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
>> @@ -0,0 +1,363 @@
>> +! { dg-do compile }
>> +! { dg-additional-options "-fmax-errors=100" }
>> +
>> +! { dg-prune-output "sorry, unimplemented" }
> 
> Likewise.
> 
>> +! { dg-prune-output "Error: work-sharing region" }
> 
> What's the intention of this?  If we're expecting this error, place
> dg-error directives where they belong?

Trunk is missing some acc loop nesting verification code in omp-low.c
that's present in gomp4. I'm not sure who's going to port that to trunk.
I'll add a comment in this test to remove it with the sorry messages
when appropriate.

>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
>> @@ -0,0 +1,80 @@
>> +! { dg-do compile }
>> +! { dg-additional-options "-fmax-errors=100" }
>> +
>> +! { dg-prune-output "sorry, unimplemented" }
> 
> Likewise.
> 
>> +! { dg-prune-output "Error: work-sharing region" }
> 
> Likewise.
> 
>> --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
>> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
>> @@ -3,6 +3,9 @@
>>  
>>  ! test for tree-dump-original and spaces-commas
>>  
>> +! { dg-prune-output "sorry, unimplemented" }
> 
> Likewise.
> 
>> +! { dg-prune-output "Error: work-sharing region" }
> 
> Likewise.
> 
>> --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
>> +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
>> @@ -37,4 +37,3 @@ end program test
>>  
>>  ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
>>  ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } 
>> -! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } } 
> 
> Which of your source code changes does this change related to?

I think Nathan made this change because he found a bug in the test or
something. I just included this test because trunk should be capable to
handle it now.

Cesar
Thomas Schwinge Nov. 4, 2015, 5:41 p.m. UTC | #7
Hi!

On Wed, 4 Nov 2015 18:21:36 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Nov 04, 2015 at 06:15:14PM +0100, Thomas Schwinge wrote:
> > > --- a/gcc/fortran/openmp.c
> > > +++ b/gcc/fortran/openmp.c
> > 
> > > @@ -3028,6 +3015,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
> > >  	n->sym->mark = 1;
> > >      }
> > >  
> > > +  /* OpenACC reductions.  */
> > > +  if (openacc)
> > > +    {
> > > +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> > > +	n->sym->mark = 0;
> > 
> > Maybe I'm just confugsed, but if setting all these to zero here...
> > 
> > > +
> > > +      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
> > > +	{
> > > +	  if (n->sym->mark)
> > > +	    gfc_error ("Symbol %qs present on multiple clauses at %L",
> > > +		       n->sym->name, &n->where);
> > > +	  else
> > > +	    n->sym->mark = 1;
> > 
> > ... won't this just always run into the "else" branch?
> 
> The point is to check if some symbol isn't present multiple times in
> reduction clause(s).  So the first loop clears the flag as it could have
> arbitrary values, and the second loop will diagnose an error if n->sym
> is present multiple times in the list.  reduction(+: a, b, a) and the like.
> In C/C++ FE we use bitmaps for this, in Fortran FE we have mark
> field for those purposes.

Thanks for the explanation -- I'd missed the fact that several "n" might
refer to the same "sym"...  ;-)


Grüße
 Thomas
Jakub Jelinek Nov. 4, 2015, 5:51 p.m. UTC | #8
On Wed, Nov 04, 2015 at 09:39:50AM -0800, Cesar Philippidis wrote:
> On 11/04/2015 09:15 AM, Thomas Schwinge wrote:
> 
> >> --- a/gcc/fortran/trans-openmp.c
> >> +++ b/gcc/fortran/trans-openmp.c
> > 
> >> @@ -3449,16 +3478,28 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
> >>  	      sizeof (construct_clauses));
> >>        loop_clauses.collapse = construct_clauses.collapse;
> >> [...]
> >> -      construct_clauses.collapse = 0;
> > 
> > Again I'm confused by this, why this is being removed, as earlier in
> > <http://news.gmane.org/find-root.php?message_id=%3C87vb9r45qw.fsf%40kepler.schwinge.homeip.net%3E>.
> 
> I'm not sure why, but gfc_trans_omp_do needs it. It's probably an openmp
> thing. If you look at gfc_trans_omp_do, you'll see that two sets of
> clauses are passed into it. code->ext.omp_clauses corresponds to the
> combined construct clauses and do_clauses are the filtered out ones.
> 
> So in order to get collapse to work as expected in combined loops, I
> can't zero out construct_clauses.collapse.

The middle end requires that for composite constructs (i.e. where one
source collapsed loop is split among multiple levels of parallelism)
all those loop constructs contain the same collapse clause, and that all
but the innermost loop construct contain NULLs on the init/cond/step etc.
vectors.

	Jakub
diff mbox

Patch

2015-11-03  Cesar Philippidis  <cesar@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	gcc/fortran/
	* openmp.c (gfc_match_omp_clauses): Update support for the tile
	and default clauses in OpenACC.
	(gfc_match_oacc_update): Error when data clauses are supplied.
	(oacc_compatible_clauses): Delete.
	(resolve_omp_clauses): Give special care for OpenACC reductions.
	Also update error reporting for the tile clause.
	(resolve_oacc_loop_blocks): Update error reporting for the tile clause.
	* trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add
	OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument.
	(gfc_trans_oacc_combined_directive): Update the list of clauses which
	are split to acc loops.

2015-11-03  Cesar Philippidis  <cesar@codesourcery.com>
	    Tom de Vries  <tom@codesourcery.com>
	    Nathan Sidwell  <nathan@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	gcc/testsuite/
	* gfortran.dg/goacc/combined-directives.f90: New test.
	* gfortran.dg/goacc/default.f95: New test.
	* gfortran.dg/goacc/default_none.f95: New test.
	* gfortran.dg/goacc/firstprivate-1.f95: New test.
	* gfortran.dg/goacc/gang-static.f95: New test.
	* gfortran.dg/goacc/kernels-loop-inner.f95: New test.
	* gfortran.dg/goacc/kernels-loops-adjacent.f95: New test.
	* gfortran.dg/goacc/list.f95: Update test.
	* gfortran.dg/goacc/loop-2.f95: Likewise.
	* gfortran.dg/goacc/loop-4.f95: New test.
	* gfortran.dg/goacc/loop-5.f95: New test.
	* gfortran.dg/goacc/loop-6.f95: New test.
	* gfortran.dg/goacc/loop-tree-1.f90: Update test.
	* gfortran.dg/goacc/modules.f95: New test.
	* gfortran.dg/goacc/multi-clause.f90: New test.
	* gfortran.dg/goacc/parallel-tree.f95: Update test.
	* gfortran.dg/goacc/update.f95: New test.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 929a739..0a92541 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -703,6 +703,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 				       OMP_MAP_FORCE_FROM))
 	continue;
       if ((mask & OMP_CLAUSE_TILE)
+	  && !c->tile_list
 	  && match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES)
 	continue;
       if ((mask & OMP_CLAUSE_SEQ) && !c->seq
@@ -856,13 +857,14 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
       if ((mask & OMP_CLAUSE_DEFAULT)
 	  && c->default_sharing == OMP_DEFAULT_UNKNOWN)
 	{
-	  if (gfc_match ("default ( shared )") == MATCH_YES)
+	  if (!openacc && gfc_match ("default ( shared )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_SHARED;
-	  else if (gfc_match ("default ( private )") == MATCH_YES)
+	  else if (!openacc && gfc_match ("default ( private )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_PRIVATE;
 	  else if (gfc_match ("default ( none )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_NONE;
-	  else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
+	  else if (!openacc
+		   && gfc_match ("default ( firstprivate )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
 	  if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
 	    continue;
@@ -1304,10 +1306,19 @@  match
 gfc_match_oacc_update (void)
 {
   gfc_omp_clauses *c;
+  locus here = gfc_current_locus;
+
   if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
       != MATCH_YES)
     return MATCH_ERROR;
 
+  if (!c->lists[OMP_LIST_MAP])
+    {
+      gfc_error ("%<acc update%> must contain at least one "
+		 "%<device%> or %<host/self%> clause at %L", &here);
+      return MATCH_ERROR;
+    }
+
   new_st.op = EXEC_OACC_UPDATE;
   new_st.ext.omp_clauses = c;
   return MATCH_YES;
@@ -2846,30 +2857,6 @@  resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
   return copy;
 }
 
-/* Returns true if clause in list 'list' is compatible with any of
-   of the clauses in lists [0..list-1].  E.g., a reduction variable may
-   appear in both reduction and private clauses, so this function
-   will return true in this case.  */
-
-static bool
-oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
-			   gfc_symbol *sym, bool openacc)
-{
-  gfc_omp_namelist *n;
-
-  if (!openacc)
-    return false;
-
-  if (list != OMP_LIST_REDUCTION)
-    return false;
-
-  for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next)
-    if (n->sym == sym)
-      return true;
-
-  return false;
-}
-
 /* OpenMP directive resolving routines.  */
 
 static void
@@ -2975,11 +2962,11 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& list != OMP_LIST_DEPEND
 	&& (list != OMP_LIST_MAP || openacc)
 	&& list != OMP_LIST_FROM
-	&& list != OMP_LIST_TO)
+	&& list != OMP_LIST_TO
+	&& (list != OMP_LIST_REDUCTION || !openacc))
       for (n = omp_clauses->lists[list]; n; n = n->next)
 	{
-	  if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
-							n->sym, openacc))
+	  if (n->sym->mark)
 	    gfc_error ("Symbol %qs present on multiple clauses at %L",
 		       n->sym->name, &n->where);
 	  else
@@ -3028,6 +3015,22 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	n->sym->mark = 1;
     }
 
+  /* OpenACC reductions.  */
+  if (openacc)
+    {
+      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	n->sym->mark = 0;
+
+      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	{
+	  if (n->sym->mark)
+	    gfc_error ("Symbol %qs present on multiple clauses at %L",
+		       n->sym->name, &n->where);
+	  else
+	    n->sym->mark = 1;
+	}
+    }
+  
   for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
     n->sym->mark = 0;
   for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
@@ -4528,22 +4531,8 @@  resolve_oacc_loop_blocks (gfc_code *code)
       if (code->ext.omp_clauses->vector)
 	gfc_error ("Clause AUTO conflicts with VECTOR at %L", &code->loc);
     }
-  if (!code->ext.omp_clauses->tile_list)
-    {
-      if (code->ext.omp_clauses->gang)
-	{
-	  if (code->ext.omp_clauses->worker)
-	    gfc_error ("Clause GANG conflicts with WORKER at %L", &code->loc);
-	  if (code->ext.omp_clauses->vector)
-	    gfc_error ("Clause GANG conflicts with VECTOR at %L", &code->loc);
-	}
-      if (code->ext.omp_clauses->worker)
-	if (code->ext.omp_clauses->vector)
-	  gfc_error ("Clause WORKER conflicts with VECTOR at %L", &code->loc);
-    }
-  else if (code->ext.omp_clauses->gang
-	   && code->ext.omp_clauses->worker
-	   && code->ext.omp_clauses->vector)
+  if (code->ext.omp_clauses->tile_list && code->ext.omp_clauses->gang
+      && code->ext.omp_clauses->worker && code->ext.omp_clauses->vector)
     gfc_error ("Tiled loop cannot be parallelized across gangs, workers and "
 	       "vectors at the same time at %L", &code->loc);
 
@@ -4564,10 +4553,21 @@  resolve_oacc_loop_blocks (gfc_code *code)
 	{
 	  num++;
 	  if (el->expr == NULL)
-	    continue;
-	  resolve_oacc_positive_int_expr (el->expr, "TILE");
-	  if (el->expr->expr_type != EXPR_CONSTANT)
-	    gfc_error ("TILE requires constant expression at %L", &code->loc);
+	    {
+	      /* NULL expressions are used to represent '*' arguments.
+		 Convert those to a -1 expressions.  */
+	      el->expr = gfc_get_constant_expr (BT_INTEGER,
+						gfc_default_integer_kind,
+						&code->loc);
+	      mpz_set_si (el->expr->value.integer, -1);
+	    }
+	  else
+	    {
+	      resolve_oacc_positive_int_expr (el->expr, "TILE");
+	      if (el->expr->expr_type != EXPR_CONSTANT)
+		gfc_error ("TILE requires constant expression at %L",
+			   &code->loc);
+	    }
 	}
       resolve_oacc_nested_loops (code, code->block->next, num, "tiled");
     }
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 5f4c382..bf8cfa1 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2534,8 +2534,12 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
     }
   if (clauses->seq)
     {
-      c = build_omp_clause (where.lb->location, OMP_CLAUSE_ORDERED);
-      OMP_CLAUSE_ORDERED_EXPR (c) = NULL_TREE;
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_SEQ);
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+  if (clauses->par_auto)
+    {
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO);
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
   if (clauses->independent)
@@ -2579,6 +2583,21 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = vector_length_var;
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
+  if (clauses->tile_list)
+    {
+      vec<tree, va_gc> *tvec;
+      gfc_expr_list *el;
+
+      vec_alloc (tvec, 4);
+
+      for (el = clauses->tile_list; el; el = el->next)
+	vec_safe_push (tvec, gfc_convert_expr_to_tree (block, el->expr));
+
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_TILE);
+      OMP_CLAUSE_TILE_LIST (c) = build_tree_list_vec (tvec);
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+      tvec->truncate (0);
+    }
   if (clauses->vector)
     {
       if (clauses->vector_expr)
@@ -2618,7 +2637,17 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  tree gang_var
 	    = gfc_convert_expr_to_tree (block, clauses->gang_expr);
 	  c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
-	  OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+	  if (clauses->gang_static)
+	    OMP_CLAUSE_GANG_STATIC_EXPR (c) = gang_var;
+	  else
+	    OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+	  omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+	}
+      else if (clauses->gang_static)
+	{
+	  /* This corresponds to gang (static: *).  */
+	  c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
+	  OMP_CLAUSE_GANG_STATIC_EXPR (c) = integer_minus_one_node;
 	  omp_clauses = gfc_trans_add_clause (c, omp_clauses);
 	}
       else
@@ -3449,16 +3478,28 @@  gfc_trans_oacc_combined_directive (gfc_code *code)
 	      sizeof (construct_clauses));
       loop_clauses.collapse = construct_clauses.collapse;
       loop_clauses.gang = construct_clauses.gang;
+      loop_clauses.gang_expr = construct_clauses.gang_expr;
+      loop_clauses.gang_static = construct_clauses.gang_static;
       loop_clauses.vector = construct_clauses.vector;
+      loop_clauses.vector_expr = construct_clauses.vector_expr;
       loop_clauses.worker = construct_clauses.worker;
+      loop_clauses.worker_expr = construct_clauses.worker_expr;
       loop_clauses.seq = construct_clauses.seq;
+      loop_clauses.par_auto = construct_clauses.par_auto;
       loop_clauses.independent = construct_clauses.independent;
-      construct_clauses.collapse = 0;
+      loop_clauses.tile_list = construct_clauses.tile_list;
       construct_clauses.gang = false;
+      construct_clauses.gang_expr = NULL;
+      construct_clauses.gang_static = NULL;
       construct_clauses.vector = false;
+      construct_clauses.vector_expr = NULL;
       construct_clauses.worker = false;
+      construct_clauses.worker_expr = NULL;
       construct_clauses.seq = false;
+      construct_clauses.par_auto = false;
+      construct_clauses.independent = false;
       construct_clauses.independent = false;
+      construct_clauses.tile_list = NULL;
       oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
 					    code->loc);
     }
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
new file mode 100644
index 0000000..620dfd2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -0,0 +1,152 @@ 
+! Exercise combined OpenACC directives.
+
+! { dg-do compile }
+! { dg-options "-fopenacc -fdump-tree-gimple" }
+
+! { dg-prune-output "sorry, unimplemented" }
+
+subroutine test
+  implicit none
+  integer a(100), i, j, z
+
+  ! PARALLEL
+  
+  !$acc parallel loop collapse (2)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+  
+  !$acc parallel loop gang
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop auto
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop tile (2, 3)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop independent
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop private (z)
+  do i = 1, 100
+     z = 0
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop reduction (+:z) copy (z)
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+
+  ! KERNELS
+
+  !$acc kernels loop collapse (2)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+  
+  !$acc kernels loop gang
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop worker
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop vector
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop seq
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop auto
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop tile (2, 3)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop independent
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop private (z)
+  do i = 1, 100
+     z = 0
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop reduction (+:z) copy (z)
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop gang" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } }
+! { 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" } }
+! { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:z .len: 4.. reduction..:z." 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/default.f95 b/gcc/testsuite/gfortran.dg/goacc/default.f95
new file mode 100644
index 0000000..c1fc52e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/default.f95
@@ -0,0 +1,17 @@ 
+! { dg-do compile }
+
+program tile
+  integer i, j, a
+
+  !$acc parallel default (shared) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+
+  !$acc parallel default (private) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+
+  !$acc parallel default (none)
+  !$acc end parallel
+
+  !$acc parallel default (firstprivate) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+end program tile
diff --git a/gcc/testsuite/gfortran.dg/goacc/default_none.f95 b/gcc/testsuite/gfortran.dg/goacc/default_none.f95
new file mode 100644
index 0000000..5ce66ae
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/default_none.f95
@@ -0,0 +1,59 @@ 
+! Ensure that the internal array variables, offset, lbound, etc., don't
+! trigger errors with default(none).
+
+! { dg-do compile }
+
+program main
+  implicit none
+  integer i
+  integer,parameter :: n = 100
+  integer,allocatable :: a1(:), a2(:,:)
+
+  allocate (a1 (n))
+  allocate (a2 (-n:n,-n:n))
+  a1 (:) = -1
+
+  !$acc parallel loop default(none) copy (a1(1:n))
+  do i = 1,n
+     a1(i) = i
+  end do
+  !$acc end parallel loop
+
+  call foo (a1)
+  call bar (a1, n)
+  call foobar (a2,n)
+
+contains
+
+  subroutine foo (da1)
+    integer :: da1(n)
+
+    !$acc parallel loop default(none) copy (da1(1:n))
+    do i = 1,n
+       da1(i) = i*2
+    end do
+    !$acc end parallel loop
+  end subroutine foo
+end program main
+
+subroutine bar (da2,n)
+  integer :: n, da2(n)
+  integer i
+
+  !$acc parallel loop default(none) copy (da2(1:n)) firstprivate(n)
+  do i = 1,n
+     da2(i) = i*3
+  end do
+  !$acc end parallel loop
+end subroutine bar
+
+subroutine foobar (da3,n)
+  integer :: n, da3(-n:n,-n:n)
+  integer i
+
+  !$acc parallel loop default(none) copy (da3(-n:n,-n:n)) firstprivate(n)
+  do i = 1,n
+     da3(i,0) = i*3
+  end do
+  !$acc end parallel loop
+end subroutine foobar
diff --git a/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95
new file mode 100644
index 0000000..fb92dee
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95
@@ -0,0 +1,11 @@ 
+! { dg-do compile }
+
+program test
+  integer a, b(100)
+
+  !$acc parallel firstprivate (a, b)
+  !$acc end parallel
+
+  !$acc parallel firstprivate (b(10:20)) ! { dg-error "Syntax error in OpenMP variable list" }
+  !$acc end parallel ! { dg-error "Unexpected !\\\$ACC END PARALLEL statement" }
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/gang-static.f95 b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95
new file mode 100644
index 0000000..4e46cf3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95
@@ -0,0 +1,68 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+program main
+  integer, parameter :: n = 100
+  integer i, a(n), b(n)
+
+  do i = 1, n
+     b(i) = i
+  end do
+
+  !$acc parallel loop gang (static:*) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 0
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 0, n)
+
+  !$acc parallel loop gang (static:1) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 1
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 1, n)
+
+  !$acc parallel loop gang (static:2) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 2
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 2, n)
+
+  !$acc parallel loop gang (static:5) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 5
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 5, n)
+
+  !$acc parallel loop gang (static:20) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 20
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 20, n)
+
+end program main
+
+subroutine test (a, b, sarg, n)
+  integer n
+  integer a (n), b(n), sarg
+  integer i
+
+  do i = 1, n
+     if (a(i) .ne. b(i) + sarg) call abort ()
+  end do
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "gang\\(static:\\\*\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:5\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:20\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95
new file mode 100644
index 0000000..4db3a50
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95
@@ -0,0 +1,23 @@ 
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=32" }
+
+program main
+   implicit none
+
+   integer :: a(100,100), b(100,100)
+   integer :: i, j, d
+
+   !$acc kernels
+   do i=1,100
+     do j=1,100
+       a(i,j) = 1
+       b(i,j) = 2
+       a(i,j) = a(i,j) + b(i,j)
+     end do
+   end do
+   !$acc end kernels
+
+   d = sum(a)
+
+   print *,d
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95
new file mode 100644
index 0000000..fef3d10
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95
@@ -0,0 +1,19 @@ 
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=10" }
+
+program main
+   implicit none
+
+   integer :: a(10000), b(10000)
+   integer :: d
+
+   !$acc kernels
+   a = 1
+   b = 2
+   a = a + b
+   !$acc end kernels
+
+   d = sum(a)
+
+   print *,d
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/list.f95 b/gcc/testsuite/gfortran.dg/goacc/list.f95
index 94fdadd..a8006bc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/list.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/list.f95
@@ -5,7 +5,7 @@  program test
   implicit none
 
   integer :: i, j, k, l, a(10)
-  common /b/ j, k
+  common /b/ k
   real, pointer :: p1 => NULL()
   complex :: c, d(10)
 
@@ -108,4 +108,4 @@  program test
   !$acc host_data use_device(p1) ! { dg-error "POINTER" }
   !$acc end host_data
 
-end program test
\ No newline at end of file
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
index f85691e..b5e6368 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
@@ -66,7 +66,7 @@  program test
     !$acc loop seq worker ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang worker ! { dg-error "conflicts with" }
+    !$acc loop gang worker
     DO i = 1,10
     ENDDO
 
@@ -94,10 +94,10 @@  program test
     !$acc loop seq vector ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang vector ! { dg-error "conflicts with" }
+    !$acc loop gang vector
     DO i = 1,10
     ENDDO
-    !$acc loop worker vector ! { dg-error "conflicts with" }
+    !$acc loop worker vector
     DO i = 1,10
     ENDDO
 
@@ -239,7 +239,7 @@  program test
     !$acc loop seq worker ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang worker ! { dg-error "conflicts with" }
+    !$acc loop gang worker
     DO i = 1,10
     ENDDO
 
@@ -267,10 +267,10 @@  program test
     !$acc loop seq vector ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang vector ! { dg-error "conflicts with" }
+    !$acc loop gang vector
     DO i = 1,10
     ENDDO
-    !$acc loop worker vector ! { dg-error "conflicts with" }
+    !$acc loop worker vector
     DO i = 1,10
     ENDDO
 
@@ -392,7 +392,7 @@  program test
   !$acc kernels loop seq worker ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc kernels loop gang worker ! { dg-error "conflicts with" }
+  !$acc kernels loop gang worker
   DO i = 1,10
   ENDDO
 
@@ -420,10 +420,10 @@  program test
   !$acc kernels loop seq vector ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc kernels loop gang vector ! { dg-error "conflicts with" }
+  !$acc kernels loop gang vector
   DO i = 1,10
   ENDDO
-  !$acc kernels loop worker vector ! { dg-error "conflicts with" }
+  !$acc kernels loop worker vector
   DO i = 1,10
   ENDDO
 
@@ -544,7 +544,7 @@  program test
   !$acc parallel loop seq worker ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc parallel loop gang worker ! { dg-error "conflicts with" }
+  !$acc parallel loop gang worker
   DO i = 1,10
   ENDDO
 
@@ -572,10 +572,10 @@  program test
   !$acc parallel loop seq vector ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc parallel loop gang vector ! { dg-error "conflicts with" }
+  !$acc parallel loop gang vector
   DO i = 1,10
   ENDDO
-  !$acc parallel loop worker vector ! { dg-error "conflicts with" }
+  !$acc parallel loop worker vector
   DO i = 1,10
   ENDDO
 
@@ -646,4 +646,4 @@  program test
   !$acc parallel loop gang worker tile(*) 
   DO i = 1,10
   ENDDO
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
new file mode 100644
index 0000000..7c53c02
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
@@ -0,0 +1,7 @@ 
+! Ensure that loops not affiliated with acc compute regions cause an error.
+
+subroutine test1
+    !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" }
+  DO i = 1,10
+  ENDDO
+end subroutine test1
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
new file mode 100644
index 0000000..5cbd975
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -0,0 +1,363 @@ 
+! { 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
+
+  !$acc kernels
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(num:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+      !$acc loop worker
+      DO j = 1,10
+      ENDDO
+    ENDDO
+
+    !$acc loop worker
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker(num:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop gang worker
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(length:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker vector
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop tile(1)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(6-2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(6+2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*, 1)
+    DO i = 1,10
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    do i = 1,10
+    enddo
+    !$acc loop vector tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang worker tile(*)
+    DO i = 1,10
+    ENDDO
+  !$acc end kernels
+
+
+  !$acc parallel
+    !$acc loop tile(1)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(2)
+    DO i = 1,10
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    do i = 1,10
+    enddo
+    !$acc loop vector tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang worker tile(*)
+    DO i = 1,10
+    ENDDO
+  !$acc end parallel
+
+  !$acc kernels loop auto
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(num:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(static:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(static:*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang
+  DO i = 1,10
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+
+  !$acc kernels loop worker
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker(num:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker
+  DO i = 1,10
+    !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop gang worker
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop vector
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector(length:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector
+  DO i = 1,10
+    !$acc kernels loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop gang vector
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker vector
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop auto
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop tile(1)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop tile(*, 1)
+  DO i = 1,10
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
+  do i = 1,10
+  enddo
+  !$acc kernels loop vector tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang worker tile(*)
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop auto
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang(static:5)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang(static:*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang
+  DO i = 1,10
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+
+  !$acc parallel loop worker
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker
+  DO i = 1,10
+    !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop gang worker
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop vector
+  DO i = 1,10
+    !$acc parallel loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop gang vector
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker vector
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop auto
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop tile(1)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop tile(*, 1)
+  DO i = 1,10
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
+  do i = 1,10
+  enddo
+  !$acc parallel loop vector tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang worker tile(*)
+  DO i = 1,10
+  ENDDO
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
new file mode 100644
index 0000000..c6236db
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
@@ -0,0 +1,80 @@ 
+! { 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
+
+  !$acc parallel
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+      !$acc loop worker
+      DO j = 1,10
+      ENDDO
+    ENDDO
+
+    !$acc loop worker
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop gang worker
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(5) ! { dg-error "argument not permitted" }
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(length:5) ! { dg-error "argument not permitted" }
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker vector
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+  !$acc end parallel
+
+  !$acc parallel loop vector
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector(5) ! { dg-error "argument not permitted" }
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector(length:5) ! { dg-error "argument not permitted" }
+  DO i = 1,10
+  ENDDO
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index d72bae4..c7f6fae 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -3,6 +3,9 @@ 
 
 ! test for tree-dump-original and spaces-commas
 
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
 program test
   implicit none
   integer :: i, j, k, m, sum
@@ -17,7 +20,7 @@  program test
 
   !$acc loop independent gang (3)
   DO i = 1,10
-    !$acc loop worker(3) ! { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+    !$acc loop worker(3)
     DO j = 1,10
       !$acc loop vector(5)
       DO k = 1,10
diff --git a/gcc/testsuite/gfortran.dg/goacc/modules.f95 b/gcc/testsuite/gfortran.dg/goacc/modules.f95
new file mode 100644
index 0000000..19a2abe
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/modules.f95
@@ -0,0 +1,55 @@ 
+! { dg-do compile } 
+
+MODULE reduction_test
+
+CONTAINS
+
+SUBROUTINE reduction_kernel(x_min,x_max,y_min,y_max,arr,sum)
+
+  IMPLICIT NONE
+
+  INTEGER      :: x_min,x_max,y_min,y_max
+  REAL(KIND=8), DIMENSION(x_min-2:x_max+2,y_min-2:y_max+2) :: arr
+  REAL(KIND=8) :: sum
+
+  INTEGER      :: j,k
+
+  sum=0.0
+
+!$ACC DATA PRESENT(arr) COPY(sum)
+!$ACC PARALLEL LOOP REDUCTION(+ : sum)
+  DO k=y_min,y_max
+    DO j=x_min,x_max
+      sum=sum*arr(j,k)
+    ENDDO
+  ENDDO
+!$ACC END PARALLEL LOOP
+!$ACC END DATA
+
+END SUBROUTINE reduction_kernel
+
+END MODULE reduction_test
+
+program main
+    use reduction_test
+
+    integer :: x_min,x_max,y_min,y_max
+    real(kind=8), dimension(1:10,1:10) :: arr
+    real(kind=8) :: sum
+
+    x_min = 5
+    x_max = 6
+    y_min = 5
+    y_max = 6
+
+    arr(:,:) = 1.0
+
+    sum = 1.0
+
+    !$acc data copy(arr)
+
+    call field_summary_kernel(x_min,x_max,y_min,y_max,arr,sum)
+
+    !$acc end data
+
+end program
diff --git a/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
new file mode 100644
index 0000000..2870076
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
@@ -0,0 +1,13 @@ 
+! Test if variable appearing in multiple clauses are errors.
+
+! { dg-compile }
+
+program combined
+  implicit none
+  integer a(100), i, j
+
+  !$acc parallel loop reduction (+:j) copy (j) copyout(j) ! { dg-error "Symbol 'j' present on multiple clauses" }
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+end program combined
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 4915744..9037f6c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -37,4 +37,3 @@  end program test
 
 ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } 
-! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } } 
diff --git a/gcc/testsuite/gfortran.dg/goacc/update.f95 b/gcc/testsuite/gfortran.dg/goacc/update.f95
new file mode 100644
index 0000000..ae23dfc
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/update.f95
@@ -0,0 +1,5 @@ 
+! { dg-do compile } 
+
+program foo
+  !$acc update ! { dg-error "must contain at least one 'device' or 'host/self' clause" }
+end program foo