diff mbox

[gomp4] Initial support of OpenACC loop directive in C front-end.

Message ID 87zjklkk2f.fsf@kepler.schwinge.homeip.net
State New
Headers show

Commit Message

Thomas Schwinge March 20, 2014, 2:42 p.m. UTC
Hi!

On Tue, 18 Mar 2014 14:50:44 +0100, I wrote:
> On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usmanov@samsung.com> wrote:
> > This patch introduces support of OpenACC loop directive (and combined 
> > directives) in C front-end up to GENERIC. Currently no clause is allowed.
> 
> Thanks!  I had worked on a simpler patch, not yet dealing with combined
> clauses.  Also, I have some work for the GIMPLE level, namely building on
> GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP.  I'll post this
> soon.

Here are the patches, committed in r208702..4 to gomp-4_0-branch.  Jakub,
are the first two fine for trunk, or shall I wait until stage 1?

commit 834daebdafa1cf4f8507fb932b7115ba3ebb02a3
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Thu Mar 20 14:39:30 2014 +0000

    Just enumerate all GF_OMP_FOR_KIND_* and GF_OMP_TARGET_KIND_*.
    
    	gcc/
    	* gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used
    	for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR,
    	GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD,
    	GF_OMP_FOR_KIND_CILKSIMD, GF_OMP_TARGET_KIND_MASK,
    	GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA,
    	GF_OMP_TARGET_KIND_UPDATE, GF_OMP_TARGET_KIND_OACC_DATA.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208702 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp |  7 +++++++
 gcc/gimple.h       | 20 ++++++++++----------
 2 files changed, 17 insertions(+), 10 deletions(-)



Grüße,
 Thomas

Comments

Jakub Jelinek March 20, 2014, 2:45 p.m. UTC | #1
On Thu, Mar 20, 2014 at 03:42:48PM +0100, Thomas Schwinge wrote:
> Here are the patches, committed in r208702..4 to gomp-4_0-branch.  Jakub,
> are the first two fine for trunk, or shall I wait until stage 1?

Stage1 IMHO.

	Jakub
Thomas Schwinge July 29, 2014, 9:07 a.m. UTC | #2
Hi Cesar!

On Thu, 20 Mar 2014 15:42:48 +0100, I wrote:
> On Tue, 18 Mar 2014 14:50:44 +0100, I wrote:
> > On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usmanov@samsung.com> wrote:
> > > This patch introduces support of OpenACC loop directive (and combined 
> > > directives) in C front-end up to GENERIC. Currently no clause is allowed.
> > 
> > Thanks!  I had worked on a simpler patch, not yet dealing with combined
> > clauses.  Also, I have some work for the GIMPLE level, namely building on
> > GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP.  I'll post this
> > soon.
> 
> Here are the patches, committed in r208702..4 to gomp-4_0-branch.

Cesar, I hope I'm not confusing things here, but I remember that you once
pointed out that in Fortran OpenACC, we have to explicitly specify the
loop iteration variable in a private clause, whereas the OpenACC
specification says this needs to happen automatically (predetermined data
attribute).  I see gcc/fortran/openmp.c:gfc_resolve_do_iterator add this
private clause, which I assume we're using also for OpenACC loops.
Looking at -fdump-tree-all output for the following two test cases
complied in OpenMP and OpenACC mode, I see that indeed an explicit
private clause is added for Fortran code, but not for C.  Why is that
required?  (I have not yet spent any time on figuring this out myself.)

    int
    main(void)
    {
      int i;
    
    #pragma acc parallel
    #pragma acc loop
    #pragma omp parallel for
      for (i = 0; i < 10; ++i)
        ;
    
      return 0;
    }

    program test
      implicit none
      integer :: i
    
      !$acc parallel
      !$acc loop
      !$omp parallel do
      DO i = 1, 10
      ENDDO
      !$omp end parallel do
      !$acc end parallel
    end

In light of this, please also review whether the following
gimplify_omp_for changes of mine (when I originally added the OACC_LOOP
support) are correct.  Evidently, this code still has TODO markers in it;
this was well before you added preliminary support for the private
clause.  That is, should we use GOVD_PRIVATE also for OACC_LOOP?

Why does this nevertheless currently work for C for loops without any
private clause for the loop iteration variable?  Maybe because in C, the
loop iteration variable ends up in a register and that is not shared
between threads?  I do see a private clause being added during
gimplification for the OpenMP C test case, but not for OpenACC -- which I
assume is precisely due to the code I'm quoting below.

Once this has been resolved, please also remove the explicit private
clauses from the existing test cases (Fortran and also C, as applicable)
to make sure that the predermined data attributes are tested there.

> --- gcc/gimplify.c
> +++ gcc/gimplify.c
> @@ -6683,14 +6683,36 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>    gimple_seq for_body, for_pre_body;
>    int i;
>    bool simd;
> +  enum gimplify_omp_var_data govd_private;
> +  enum omp_region_type ort;
>    bitmap has_decl_expr = NULL;
>  
>    orig_for_stmt = for_stmt = *expr_p;
>  
> -  simd = TREE_CODE (for_stmt) == OMP_SIMD
> -    || TREE_CODE (for_stmt) == CILK_SIMD;
> -  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
> -			     simd ? ORT_SIMD : ORT_WORKSHARE);
> +  switch (TREE_CODE (for_stmt))
> +    {
> +    case OMP_FOR:
> +    case OMP_DISTRIBUTE:
> +      simd = false;
> +      govd_private = GOVD_PRIVATE;
> +      ort = ORT_WORKSHARE;
> +      break;
> +    case OACC_LOOP:
> +      simd = false;
> +      govd_private = /* TODO */ GOVD_LOCAL;
> +      ort = /* TODO */ ORT_WORKSHARE;
> +      break;
> +    case OMP_SIMD:
> +    case CILK_SIMD:
> +      simd = true;
> +      govd_private = GOVD_PRIVATE;
> +      ort = ORT_SIMD;
> +      break;
> +    default:
> +      gcc_unreachable ();
> +    }
> +
> +  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
>  
>    /* Handle OMP_FOR_INIT.  */
>    for_pre_body = NULL;
> @@ -6722,6 +6744,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>  
>    if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
>      {
> +      gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
>        for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
>  			    NULL, NULL);
>        gcc_assert (for_stmt != NULL_TREE);
> @@ -6742,7 +6765,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>        gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
>  		  || POINTER_TYPE_P (TREE_TYPE (decl)));
>  
> -      /* Make sure the iteration variable is private.  */
> +      /* Make sure the iteration variable is some kind of private.  */
>        tree c = NULL_TREE;
>        if (orig_for_stmt != for_stmt)
>  	/* Do this only on innermost construct for combined ones.  */;
> @@ -6768,6 +6791,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>  	    }
>  	  else
>  	    {
> +	      gcc_assert (govd_private == GOVD_PRIVATE);
>  	      bool lastprivate
>  		= (!has_decl_expr
>  		   || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
> @@ -6785,7 +6809,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>        else if (omp_is_private (gimplify_omp_ctxp, decl, simd))
>  	omp_notice_variable (gimplify_omp_ctxp, decl, true);
>        else
> -	omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
> +	omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
>  
>        /* If DECL is not a gimple register, create a temporary variable to act
>  	 as an iteration counter.  This is valid, since DECL cannot be
> @@ -6799,7 +6823,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>  
>  	  gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
>  
> -	  omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
> +	  omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>  	}
>        else
>  	var = decl;
> @@ -6936,7 +6960,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>  	t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
>  	decl = TREE_OPERAND (t, 0);
>  	var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
> -	omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
> +	omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>  	TREE_OPERAND (t, 0) = var;
>  	t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
>  	TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
> @@ -6952,6 +6976,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>      case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
>      case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
>      case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
> +    case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
>      default:
>        gcc_unreachable ();
>      }


Grüße,
 Thomas
Cesar Philippidis Aug. 5, 2014, 3:22 p.m. UTC | #3
On 07/29/2014 02:07 AM, Thomas Schwinge wrote:

> On Thu, 20 Mar 2014 15:42:48 +0100, I wrote:
>> On Tue, 18 Mar 2014 14:50:44 +0100, I wrote:
>>> On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usmanov@samsung.com> wrote:
>>>> This patch introduces support of OpenACC loop directive (and combined 
>>>> directives) in C front-end up to GENERIC. Currently no clause is allowed.
>>>
>>> Thanks!  I had worked on a simpler patch, not yet dealing with combined
>>> clauses.  Also, I have some work for the GIMPLE level, namely building on
>>> GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP.  I'll post this
>>> soon.
>>
>> Here are the patches, committed in r208702..4 to gomp-4_0-branch.
> 
> Cesar, I hope I'm not confusing things here, but I remember that you once
> pointed out that in Fortran OpenACC, we have to explicitly specify the
> loop iteration variable in a private clause, whereas the OpenACC
> specification says this needs to happen automatically (predetermined data
> attribute).  I see gcc/fortran/openmp.c:gfc_resolve_do_iterator add this
> private clause, which I assume we're using also for OpenACC loops.
> Looking at -fdump-tree-all output for the following two test cases
> complied in OpenMP and OpenACC mode, I see that indeed an explicit
> private clause is added for Fortran code, but not for C.  Why is that
> required?  (I have not yet spent any time on figuring this out myself.)
> 
>     int
>     main(void)
>     {
>       int i;
>     
>     #pragma acc parallel
>     #pragma acc loop
>     #pragma omp parallel for
>       for (i = 0; i < 10; ++i)
>         ;
>     
>       return 0;
>     }
> 
>     program test
>       implicit none
>       integer :: i
>     
>       !$acc parallel
>       !$acc loop
>       !$omp parallel do
>       DO i = 1, 10
>       ENDDO
>       !$omp end parallel do
>       !$acc end parallel
>     end
> 
> In light of this, please also review whether the following
> gimplify_omp_for changes of mine (when I originally added the OACC_LOOP
> support) are correct.  Evidently, this code still has TODO markers in it;
> this was well before you added preliminary support for the private
> clause.  That is, should we use GOVD_PRIVATE also for OACC_LOOP?

OMP has both a private and a shared clause. And those neither of those
clauses can share variables, i.e. an index variable cannot be both
private and shared. By default, index variables are private in OMP, so
the gimplification pass makes those variables explicitly private so that
later passes can check for errors.

Because the fortran frontend needs to do a little more error handling
early on, it makes the induction variables private.

> Why does this nevertheless currently work for C for loops without any
> private clause for the loop iteration variable?  Maybe because in C, the
> loop iteration variable ends up in a register and that is not shared
> between threads?  I do see a private clause being added during
> gimplification for the OpenMP C test case, but not for OpenACC -- which I
> assume is precisely due to the code I'm quoting below.

OMP lowering pass makes the index variables local to the parallel
function/kernel. So each thread effectively has its private copy of the
index variable.

> Once this has been resolved, please also remove the explicit private
> clauses from the existing test cases (Fortran and also C, as applicable)
> to make sure that the predermined data attributes are tested there.

Considering that ACC lacks a shared clause, it should be harmless to
revert that patch below. Do you want to revert it with your
private/firstprivate patch?

Cesar

>> --- gcc/gimplify.c
>> +++ gcc/gimplify.c
>> @@ -6683,14 +6683,36 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>    gimple_seq for_body, for_pre_body;
>>    int i;
>>    bool simd;
>> +  enum gimplify_omp_var_data govd_private;
>> +  enum omp_region_type ort;
>>    bitmap has_decl_expr = NULL;
>>  
>>    orig_for_stmt = for_stmt = *expr_p;
>>  
>> -  simd = TREE_CODE (for_stmt) == OMP_SIMD
>> -    || TREE_CODE (for_stmt) == CILK_SIMD;
>> -  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
>> -			     simd ? ORT_SIMD : ORT_WORKSHARE);
>> +  switch (TREE_CODE (for_stmt))
>> +    {
>> +    case OMP_FOR:
>> +    case OMP_DISTRIBUTE:
>> +      simd = false;
>> +      govd_private = GOVD_PRIVATE;
>> +      ort = ORT_WORKSHARE;
>> +      break;
>> +    case OACC_LOOP:
>> +      simd = false;
>> +      govd_private = /* TODO */ GOVD_LOCAL;
>> +      ort = /* TODO */ ORT_WORKSHARE;
>> +      break;
>> +    case OMP_SIMD:
>> +    case CILK_SIMD:
>> +      simd = true;
>> +      govd_private = GOVD_PRIVATE;
>> +      ort = ORT_SIMD;
>> +      break;
>> +    default:
>> +      gcc_unreachable ();
>> +    }
>> +
>> +  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
>>  
>>    /* Handle OMP_FOR_INIT.  */
>>    for_pre_body = NULL;
>> @@ -6722,6 +6744,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>  
>>    if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
>>      {
>> +      gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
>>        for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
>>  			    NULL, NULL);
>>        gcc_assert (for_stmt != NULL_TREE);
>> @@ -6742,7 +6765,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>        gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
>>  		  || POINTER_TYPE_P (TREE_TYPE (decl)));
>>  
>> -      /* Make sure the iteration variable is private.  */
>> +      /* Make sure the iteration variable is some kind of private.  */
>>        tree c = NULL_TREE;
>>        if (orig_for_stmt != for_stmt)
>>  	/* Do this only on innermost construct for combined ones.  */;
>> @@ -6768,6 +6791,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>  	    }
>>  	  else
>>  	    {
>> +	      gcc_assert (govd_private == GOVD_PRIVATE);
>>  	      bool lastprivate
>>  		= (!has_decl_expr
>>  		   || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
>> @@ -6785,7 +6809,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>        else if (omp_is_private (gimplify_omp_ctxp, decl, simd))
>>  	omp_notice_variable (gimplify_omp_ctxp, decl, true);
>>        else
>> -	omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
>> +	omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
>>  
>>        /* If DECL is not a gimple register, create a temporary variable to act
>>  	 as an iteration counter.  This is valid, since DECL cannot be
>> @@ -6799,7 +6823,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>  
>>  	  gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
>>  
>> -	  omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
>> +	  omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>>  	}
>>        else
>>  	var = decl;
>> @@ -6936,7 +6960,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>  	t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
>>  	decl = TREE_OPERAND (t, 0);
>>  	var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
>> -	omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
>> +	omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>>  	TREE_OPERAND (t, 0) = var;
>>  	t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
>>  	TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
>> @@ -6952,6 +6976,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>      case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
>>      case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
>>      case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
>> +    case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
>>      default:
>>        gcc_unreachable ();
>>      }
> 
> 
> Grüße,
>  Thomas
>
diff mbox

Patch

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index f43452c..72828fb 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,12 @@ 
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used
+	for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR,
+	GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD,
+	GF_OMP_FOR_KIND_CILKSIMD, GF_OMP_TARGET_KIND_MASK,
+	GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA,
+	GF_OMP_TARGET_KIND_UPDATE, GF_OMP_TARGET_KIND_OACC_DATA.
+
 	* omp-low.c (check_omp_nesting_restrictions): Allow nesting of
 	OpenACC constructs inside of OpenACC data constructs.
 
diff --git gcc/gimple.h gcc/gimple.h
index 910072d..17441ac 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -91,18 +91,18 @@  enum gf_mask {
     GF_CALL_ALLOCA_FOR_VAR	= 1 << 5,
     GF_CALL_INTERNAL		= 1 << 6,
     GF_OMP_PARALLEL_COMBINED	= 1 << 0,
-    GF_OMP_FOR_KIND_MASK	= 3 << 0,
-    GF_OMP_FOR_KIND_FOR		= 0 << 0,
-    GF_OMP_FOR_KIND_DISTRIBUTE	= 1 << 0,
-    GF_OMP_FOR_KIND_SIMD	= 2 << 0,
-    GF_OMP_FOR_KIND_CILKSIMD	= 3 << 0,
+    GF_OMP_FOR_KIND_MASK	= (1 << 2) - 1,
+    GF_OMP_FOR_KIND_FOR		= 0,
+    GF_OMP_FOR_KIND_DISTRIBUTE	= 1,
+    GF_OMP_FOR_KIND_SIMD	= 2,
+    GF_OMP_FOR_KIND_CILKSIMD	= 3,
     GF_OMP_FOR_COMBINED		= 1 << 2,
     GF_OMP_FOR_COMBINED_INTO	= 1 << 3,
-    GF_OMP_TARGET_KIND_MASK	= 3 << 0,
-    GF_OMP_TARGET_KIND_REGION	= 0 << 0,
-    GF_OMP_TARGET_KIND_DATA	= 1 << 0,
-    GF_OMP_TARGET_KIND_UPDATE	= 2 << 0,
-    GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0,
+    GF_OMP_TARGET_KIND_MASK	= (1 << 2) - 1,
+    GF_OMP_TARGET_KIND_REGION	= 0,
+    GF_OMP_TARGET_KIND_DATA	= 1,
+    GF_OMP_TARGET_KIND_UPDATE	= 2,
+    GF_OMP_TARGET_KIND_OACC_DATA = 3,
 
     /* True on an GIMPLE_OMP_RETURN statement if the return does not require
        a thread synchronization via some sort of barrier.  The exact barrier

commit c32a48d3d47bbaa811991e2e5f42e62d9c715a60
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Thu Mar 20 14:39:42 2014 +0000

    GF_OMP_FOR_SIMD: Flag for SIMD variants of OMP_FOR kinds.
    
    	gcc/
    	* gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD.
    	* omp-low.c: Update accordingly.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208703 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp |  3 +++
 gcc/gimple.h       |  6 ++++--
 gcc/omp-low.c      | 18 +++++++++---------
 3 files changed, 16 insertions(+), 11 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 72828fb..1753d73 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@ 
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD.
+	* omp-low.c: Update accordingly.
+
 	* gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used
 	for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR,
 	GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD,
diff --git gcc/gimple.h gcc/gimple.h
index 17441ac..34a0bdb 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -94,8 +94,10 @@  enum gf_mask {
     GF_OMP_FOR_KIND_MASK	= (1 << 2) - 1,
     GF_OMP_FOR_KIND_FOR		= 0,
     GF_OMP_FOR_KIND_DISTRIBUTE	= 1,
-    GF_OMP_FOR_KIND_SIMD	= 2,
-    GF_OMP_FOR_KIND_CILKSIMD	= 3,
+    /* Flag for SIMD variants of OMP_FOR kinds.  */
+    GF_OMP_FOR_SIMD		= 1 << 1,
+    GF_OMP_FOR_KIND_SIMD	= GF_OMP_FOR_SIMD | 0,
+    GF_OMP_FOR_KIND_CILKSIMD	= GF_OMP_FOR_SIMD | 1,
     GF_OMP_FOR_COMBINED		= 1 << 2,
     GF_OMP_FOR_COMBINED_INTO	= 1 << 3,
     GF_OMP_TARGET_KIND_MASK	= (1 << 2) - 1,
diff --git gcc/omp-low.c gcc/omp-low.c
index 23a0dda..c3b3e95 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -298,7 +298,7 @@  extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd,
   int i;
   struct omp_for_data_loop dummy_loop;
   location_t loc = gimple_location (for_stmt);
-  bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_KIND_SIMD;
+  bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_SIMD;
   bool distribute = gimple_omp_for_kind (for_stmt)
 		    == GF_OMP_FOR_KIND_DISTRIBUTE;
 
@@ -1024,7 +1024,7 @@  build_outer_var_ref (tree var, omp_context *ctx)
       x = build_receiver_ref (var, by_ref, ctx);
     }
   else if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-	   && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+	   && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
     {
       /* #pragma omp simd isn't a worksharing construct, and can reference even
 	 private vars in its linear etc. clauses.  */
@@ -2451,7 +2451,7 @@  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
   if (ctx != NULL)
     {
       if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-	  && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+	  && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
 	{
 	  error_at (gimple_location (stmt),
 		    "OpenMP constructs may not be nested inside simd region");
@@ -2474,7 +2474,7 @@  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
   switch (gimple_code (stmt))
     {
     case GIMPLE_OMP_FOR:
-      if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_KIND_SIMD)
+      if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD)
 	return true;
       if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
 	{
@@ -2802,7 +2802,7 @@  scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	  if (setjmp_or_longjmp_p (fndecl)
 	      && ctx
 	      && gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-	      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+	      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
 	    {
 	      remove = true;
 	      error_at (gimple_location (stmt),
@@ -3225,7 +3225,7 @@  lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
   bool reduction_omp_orig_ref = false;
   int pass;
   bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-		  && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD);
+		  && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD);
   int max_vf = 0;
   tree lane = NULL_TREE, idx = NULL_TREE;
   tree ivar = NULL_TREE, lvar = NULL_TREE;
@@ -3969,7 +3969,7 @@  lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
     }
 
   if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
     {
       simduid = find_omp_clause (orig_clauses, OMP_CLAUSE__SIMDUID_);
       if (simduid)
@@ -4066,7 +4066,7 @@  lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
 
   /* SIMD reductions are handled in lower_rec_input_clauses.  */
   if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
     return;
 
   /* First see if there is exactly one reduction clause.  Use OMP_ATOMIC
@@ -7393,7 +7393,7 @@  expand_omp_for (struct omp_region *region, gimple inner_stmt)
        original loops from being detected.  Fix that up.  */
     loops_state_set (LOOPS_NEED_FIXUP);
 
-  if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_KIND_SIMD)
+  if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_SIMD)
     expand_omp_simd (region, &fd);
   else if (fd.sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
 	   && !fd.have_ordered)

commit f1d39706db8dccbc988e2c66552511cd54632257
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Thu Mar 20 14:40:01 2014 +0000

    Continue implementation of OpenACC loop construct.
    
    	gcc/
    	* gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP.
    	(is_gimple_omp_oacc_specifically): Handle it.
    	* gimple-pretty-print.c (dump_gimple_omp_for): Likewise.
    	* gimple.def (GIMPLE_OMP_FOR): Update for OpenACC loop.
    	* gimple.c (gimple_build_omp_for): Don't explicitly mention some
    	clauses.
    	(gimple_copy) <GIMPLE_OMP_FOR>: Handle GF_OMP_FOR_KIND_OACC_LOOP.
    	* omp-low.c (extract_omp_for_data, scan_sharing_clauses)
    	(check_omp_nesting_restrictions, lower_rec_input_clauses)
    	(lower_lastprivate_clauses, lower_reduction_clauses)
    	(expand_omp_for_generic, expand_omp_for_static_nochunk)
    	(expand_omp_for_static_chunk, maybe_add_implicit_barrier_cancel)
    	(lower_omp_for): Likewise.
    	* tree-inline.c (remap_gimple_stmt): Likewise.
    	* tree-nested.c (walk_gimple_omp_for)
    	(convert_nonlocal_reference_stmt, convert_local_reference_stmt)
    	(convert_gimple_call): Likewise.
    	* doc/gimple.texi (GIMPLE_OMP_FOR): Don't explicitly mention some
    	clauses.
    	* gimplify.c (gimplify_omp_for, gimplify_expr): Handle OACC_LOOP.
    	gcc/testsuite/
    	* c-c++-common/goacc-gomp/nesting-1.c: New file.
    	* c-c++-common/goacc-gomp/nesting-fail-1.c: Extend.
    	* c-c++-common/goacc/clauses-fail.c: Likewise.
    	* c-c++-common/goacc/nesting-1.c: Likewise.
    	* gcc.dg/goacc/sb-1.c: Likewise.
    	* gcc.dg/goacc/sb-3.c: New file.
    
    	gcc/c-family/
    	* c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update
    	comments.
    	* c-pragma.c (oacc_pragmas): Sort PRAGMA_OACC_LOOP alphabetically.
    	* c-pragma.h (enum pragma_kind): Likewise.
    	gcc/c/
    	* c-parser.c: Update comments.
    	(c_parser_oacc_loop): Move.
    	(c_parser_omp_for_loop): Catch some unsupported cases.
    	(c_parser_omp_construct) <case PRAGMA_OACC_LOOP>: Sort
    	alphabetically.
    	gcc/
    	* tree.def (OACC_LOOP): Sort after OMP_DISTRIBUTE.
    	* tree.h (OMP_LOOP_CHECK): Update accordingly.
    	* gimplify.c (is_gimple_stmt): Sort OACC_LOOP after
    	OMP_DISTRIBUTE.
    	* tree-pretty-print.c (dump_generic_node): Likewise.
    	* doc/generic.texi (OACC_LOOP): Sort after OACC_HOST_DATA.
    	(OMP_FOR): Fix and extend for OACC_LOOP.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208704 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                 |  29 ++++
 gcc/c-family/ChangeLog.gomp                        |   7 +
 gcc/c-family/c-omp.c                               |   4 +-
 gcc/c-family/c-pragma.c                            |   2 +-
 gcc/c-family/c-pragma.h                            |   2 +-
 gcc/c/ChangeLog.gomp                               |   8 ++
 gcc/c/c-parser.c                                   |  80 ++++++-----
 gcc/doc/generic.texi                               |  21 +--
 gcc/doc/gimple.texi                                |   5 +-
 gcc/gimple-pretty-print.c                          |  18 ++-
 gcc/gimple.c                                       |   4 +-
 gcc/gimple.def                                     |   5 +-
 gcc/gimple.h                                       | 117 ++++++++-------
 gcc/gimplify.c                                     |  45 ++++--
 gcc/omp-low.c                                      | 157 +++++++++++++++------
 gcc/testsuite/ChangeLog.gomp                       |   7 +
 gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c  |  12 ++
 .../c-c++-common/goacc-gomp/nesting-fail-1.c       |  98 +++++++++++++
 gcc/testsuite/c-c++-common/goacc/clauses-fail.c    |   6 +
 gcc/testsuite/c-c++-common/goacc/nesting-1.c       |  43 ++++++
 gcc/testsuite/gcc.dg/goacc/sb-1.c                  |  21 +++
 gcc/testsuite/gcc.dg/goacc/sb-3.c                  |  18 +++
 gcc/tree-inline.c                                  |   1 +
 gcc/tree-nested.c                                  |   5 +
 gcc/tree-pretty-print.c                            |   8 +-
 gcc/tree.def                                       |  10 +-
 gcc/tree.h                                         |   2 +-
 27 files changed, 560 insertions(+), 175 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 1753d73..1d35b58 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,34 @@ 
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP.
+	(is_gimple_omp_oacc_specifically): Handle it.
+	* gimple-pretty-print.c (dump_gimple_omp_for): Likewise.
+	* gimple.def (GIMPLE_OMP_FOR): Update for OpenACC loop.
+	* gimple.c (gimple_build_omp_for): Don't explicitly mention some
+	clauses.
+	(gimple_copy) <GIMPLE_OMP_FOR>: Handle GF_OMP_FOR_KIND_OACC_LOOP.
+	* omp-low.c (extract_omp_for_data, scan_sharing_clauses)
+	(check_omp_nesting_restrictions, lower_rec_input_clauses)
+	(lower_lastprivate_clauses, lower_reduction_clauses)
+	(expand_omp_for_generic, expand_omp_for_static_nochunk)
+	(expand_omp_for_static_chunk, maybe_add_implicit_barrier_cancel)
+	(lower_omp_for): Likewise.
+	* tree-inline.c (remap_gimple_stmt): Likewise.
+	* tree-nested.c (walk_gimple_omp_for)
+	(convert_nonlocal_reference_stmt, convert_local_reference_stmt)
+	(convert_gimple_call): Likewise.
+	* doc/gimple.texi (GIMPLE_OMP_FOR): Don't explicitly mention some
+	clauses.
+	* gimplify.c (gimplify_omp_for, gimplify_expr): Handle OACC_LOOP.
+
+	* tree.def (OACC_LOOP): Sort after OMP_DISTRIBUTE.
+	* tree.h (OMP_LOOP_CHECK): Update accordingly.
+	* gimplify.c (is_gimple_stmt): Sort OACC_LOOP after
+	OMP_DISTRIBUTE.
+	* tree-pretty-print.c (dump_generic_node): Likewise.
+	* doc/generic.texi (OACC_LOOP): Sort after OACC_HOST_DATA.
+	(OMP_FOR): Fix and extend for OACC_LOOP.
+
 	* gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD.
 	* omp-low.c: Update accordingly.
 
diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index b33b365..37ebfe9 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,10 @@ 
+2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update
+	comments.
+	* c-pragma.c (oacc_pragmas): Sort PRAGMA_OACC_LOOP alphabetically.
+	* c-pragma.h (enum pragma_kind): Likewise.
+
 2014-03-18  Ilmir Usmanov  <i.usmanov@samsung.com>
 
 	* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_LOOP.
diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index 06f5712..5a1fb6d 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -293,7 +293,7 @@  c_finish_omp_flush (location_t loc)
 }
 
 
-/* Check and canonicalize #pragma omp for increment expression.
+/* Check and canonicalize OMP_FOR increment expression.
    Helper function for c_finish_omp_for.  */
 
 static tree
@@ -381,7 +381,7 @@  c_omp_for_incr_canonicalize_ptr (location_t loc, tree decl, tree incr)
   return incr;
 }
 
-/* Validate and emit code for the OpenMP directive #pragma omp for.
+/* Validate and generate OMP_FOR.
    DECLV is a vector of iteration variables, for each collapsed loop.
    INITV, CONDV and INCRV are vectors containing initialization
    expressions, controlling predicates and increment expressions.
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index f99b087..aef4f10 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1171,8 +1171,8 @@  struct omp_pragma_def { const char *name; unsigned int id; };
 static const struct omp_pragma_def oacc_pragmas[] = {
   { "data", PRAGMA_OACC_DATA },
   { "kernels", PRAGMA_OACC_KERNELS },
-  { "parallel", PRAGMA_OACC_PARALLEL },
   { "loop", PRAGMA_OACC_LOOP },
+  { "parallel", PRAGMA_OACC_PARALLEL },
 };
 static const struct omp_pragma_def omp_pragmas[] = {
   { "atomic", PRAGMA_OMP_ATOMIC },
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index f4b5b80..bb9c367 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -29,8 +29,8 @@  typedef enum pragma_kind {
 
   PRAGMA_OACC_DATA,
   PRAGMA_OACC_KERNELS,
-  PRAGMA_OACC_PARALLEL,
   PRAGMA_OACC_LOOP,
+  PRAGMA_OACC_PARALLEL,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
   PRAGMA_OMP_CANCEL,
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 0358b3a..91978db 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,11 @@ 
+2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-parser.c: Update comments.
+	(c_parser_oacc_loop): Move.
+	(c_parser_omp_for_loop): Catch some unsupported cases.
+	(c_parser_omp_construct) <case PRAGMA_OACC_LOOP>: Sort
+	alphabetically.
+
 2014-03-18  Ilmir Usmanov  <i.usmanov@samsung.com>
 
 	* c-parser.c (c_parser_oacc_loop): New function.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 734d44e..90d0035 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1204,10 +1204,13 @@  static struct c_expr c_parser_expression_conv (c_parser *);
 static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
 					     vec<tree, va_gc> **, location_t *,
 					     tree *, vec<location_t> *);
+static tree c_parser_oacc_loop (location_t, c_parser *, char *);
 static void c_parser_omp_construct (c_parser *);
 static void c_parser_omp_threadprivate (c_parser *);
 static void c_parser_omp_barrier (c_parser *);
 static void c_parser_omp_flush (c_parser *);
+static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code,
+				   tree, tree *);
 static void c_parser_omp_taskwait (c_parser *);
 static void c_parser_omp_taskyield (c_parser *);
 static void c_parser_omp_cancel (c_parser *);
@@ -4778,6 +4781,7 @@  c_parser_label (c_parser *parser)
      parallel-construct
      kernels-construct
      data-construct
+     loop-construct
 
    parallel-construct:
      parallel-directive structured-block
@@ -4788,6 +4792,9 @@  c_parser_label (c_parser *parser)
    data-construct:
      data-directive structured-block
 
+   loop-construct:
+     loop-directive structured-block
+
    OpenMP:
 
    statement:
@@ -11557,8 +11564,6 @@  c_parser_oacc_data (location_t loc, c_parser *parser)
   return stmt;
 }
 
-static tree c_parser_oacc_loop (location_t, c_parser *, char *);
-
 /* OpenACC 2.0:
    # pragma acc kernels oacc-kernels-clause[optseq] new-line
      structured-block
@@ -11611,6 +11616,33 @@  c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
 }
 
 /* OpenACC 2.0:
+   # pragma acc loop oacc-loop-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_LOOP_CLAUSE_MASK						\
+	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE)
+
+static tree
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+{
+  tree stmt, clauses, block;
+
+  strcat (p_name, " loop");
+
+  clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+
+  block = c_begin_compound_stmt (true);
+  stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+  block = c_end_compound_stmt (loc, block, true);
+  add_stmt (block);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
    # pragma acc parallel oacc-parallel-clause[optseq] new-line
      structured-block
 
@@ -12120,10 +12152,11 @@  c_parser_omp_flush (c_parser *parser)
   c_finish_omp_flush (loc);
 }
 
-/* Parse the restricted form of the for statement allowed by OpenMP.
+/* Parse the restricted form of loop statements allowed by OpenACC and OpenMP.
    The real trick here is to determine the loop control variable early
    so that we can push a new decl if necessary to make it private.
-   LOC is the location of the OMP in "#pragma omp".  */
+   LOC is the location of the "acc" or "omp" in "#pragma acc" or "#pragma omp",
+   respectively.  */
 
 static tree
 c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
@@ -12138,7 +12171,10 @@  c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
 
   for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
     if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
+      {
+	gcc_assert (code != OACC_LOOP);
       collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
+      }
 
   gcc_assert (collapse >= 1);
 
@@ -12369,6 +12405,7 @@  c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
 	  if (cclauses != NULL
 	      && cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL)
 	    {
+	      gcc_assert (code != OACC_LOOP);
 	      tree *c;
 	      for (c = &cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL]; *c ; )
 		if (OMP_CLAUSE_CODE (*c) != OMP_CLAUSE_FIRSTPRIVATE
@@ -12433,33 +12470,6 @@  omp_split_clauses (location_t loc, enum tree_code code,
       cclauses[i] = c_finish_omp_clauses (cclauses[i]);
 }
 
-/* OpenACC 2.0:
-   # pragma acc loop oacc-loop-clause[optseq] new-line
-     structured-block
-
-   LOC is the location of the #pragma token.
-*/
-
-#define OACC_LOOP_CLAUSE_MASK	\
-	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE)
-
-static tree
-c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
-{
-  tree block, clauses, ret;
-
-  strcat (p_name, " loop");
-
-  clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
-
-  block = c_begin_compound_stmt (true);
-  ret = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
-  block = c_end_compound_stmt (loc, block, true);
-  add_stmt (block);
-
-  return ret;
-}
-
 /* OpenMP 4.0:
    #pragma omp simd simd-clause[optseq] new-line
      for-loop
@@ -13977,14 +13987,14 @@  c_parser_omp_construct (c_parser *parser)
       strcpy (p_name, "#pragma acc");
       stmt = c_parser_oacc_kernels (loc, parser, p_name);
       break;
-    case PRAGMA_OACC_PARALLEL:
-      strcpy (p_name, "#pragma acc");
-      stmt = c_parser_oacc_parallel (loc, parser, p_name);
-      break;
     case PRAGMA_OACC_LOOP:
       strcpy (p_name, "#pragma acc");
       stmt = c_parser_oacc_loop (loc, parser, p_name);
       break;
+    case PRAGMA_OACC_PARALLEL:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_parallel (loc, parser, p_name);
+      break;
     case PRAGMA_OMP_ATOMIC:
       c_parser_omp_atomic (loc, parser);
       return;
diff --git gcc/doc/generic.texi gcc/doc/generic.texi
index 0a77a86..7780fe8 100644
--- gcc/doc/generic.texi
+++ gcc/doc/generic.texi
@@ -2054,8 +2054,8 @@  edge.  Rethrowing the exception is represented using @code{RESX_EXPR}.
 @tindex OACC_PARALLEL
 @tindex OACC_KERNELS
 @tindex OACC_DATA
-@tindex OACC_LOOP
 @tindex OACC_HOST_DATA
+@tindex OACC_LOOP
 @tindex OACC_DECLARE
 @tindex OACC_UPDATE
 @tindex OACC_ENTER_DATA
@@ -2091,14 +2091,16 @@  Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
 
 Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
 
-@item OACC_LOOP
-
-Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}.
-
 @item OACC_HOST_DATA
 
 Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}.
 
+@item OACC_LOOP
+
+Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}.
+
+See the description of the @code{OMP_FOR} code.
+
 @item OACC_DECLARE
 
 Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}.
@@ -2150,8 +2152,8 @@  variables.
 
 @item OMP_FOR
 
-Represents @code{#pragma omp for [clause1 @dots{} clauseN]}.  It
-has 5 operands:
+Represents @code{#pragma omp for [clause1 @dots{} clauseN]}.  It has
+six operands:
 
 Operand @code{OMP_FOR_BODY} contains the loop body.
 
@@ -2241,10 +2243,9 @@  building code (@code{omp-low.c}).
 @item OMP_CONTINUE
 
 Similarly, this instruction does not represent an OpenMP
-directive, it is used by @code{OMP_FOR} and
+directive, it is used by @code{OACC_LOOP}, @code{OMP_FOR} as well as
 @code{OMP_SECTIONS} to mark the place where the code needs to
-loop to the next iteration (in the case of @code{OMP_FOR}) or
-the next section (in the case of @code{OMP_SECTIONS}).
+loop to the next iteration, or the next section, respectively.
 
 In some cases, @code{OMP_CONTINUE} is placed right before
 @code{OMP_RETURN}.  But if there are cleanups that need to
diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi
index 91748a6..fd6feae 100644
--- gcc/doc/gimple.texi
+++ gcc/doc/gimple.texi
@@ -1679,9 +1679,8 @@  Set @code{NAME} to be the name associated with @code{OMP} critical statement @co
 tree clauses, tree index, tree initial, tree final, tree incr, @
 gimple_seq pre_body, enum tree_code omp_for_cond)
 Build a @code{GIMPLE_OMP_FOR} statement. @code{BODY} is sequence of statements
-inside the for loop.  @code{CLAUSES}, are any of the @code{OMP} loop
-construct's clauses: private, firstprivate,  lastprivate,
-reductions, ordered, schedule, and nowait.  @code{PRE_BODY} is the
+inside the for loop.  @code{CLAUSES}, are any of the loop
+construct's clauses.  @code{PRE_BODY} is the
 sequence of statements that are loop invariant.  @code{INDEX} is the
 index variable.  @code{INITIAL} is the initial value of @code{INDEX}.  @code{FINAL} is
 final value of @code{INDEX}.  OMP_FOR_COND is the predicate used to
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index c62c517..f251060 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1116,15 +1116,18 @@  dump_gimple_omp_for (pretty_printer *buffer, gimple gs, int spc, int flags)
 	case GF_OMP_FOR_KIND_FOR:
 	  kind = "";
 	  break;
+	case GF_OMP_FOR_KIND_DISTRIBUTE:
+	  kind = " distribute";
+	  break;
+	case GF_OMP_FOR_KIND_OACC_LOOP:
+	  kind = " oacc_loop";
+	  break;
 	case GF_OMP_FOR_KIND_SIMD:
 	  kind = " simd";
 	  break;
 	case GF_OMP_FOR_KIND_CILKSIMD:
 	  kind = " cilksimd";
 	  break;
-	case GF_OMP_FOR_KIND_DISTRIBUTE:
-	  kind = " distribute";
-	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -1150,15 +1153,18 @@  dump_gimple_omp_for (pretty_printer *buffer, gimple gs, int spc, int flags)
 	case GF_OMP_FOR_KIND_FOR:
 	  pp_string (buffer, "#pragma omp for");
 	  break;
+	case GF_OMP_FOR_KIND_DISTRIBUTE:
+	  pp_string (buffer, "#pragma omp distribute");
+	  break;
+	case GF_OMP_FOR_KIND_OACC_LOOP:
+	  pp_string (buffer, "#pragma acc loop");
+	  break;
 	case GF_OMP_FOR_KIND_SIMD:
 	  pp_string (buffer, "#pragma omp simd");
 	  break;
 	case GF_OMP_FOR_KIND_CILKSIMD:
 	  pp_string (buffer, "#pragma simd");
 	  break;
-	case GF_OMP_FOR_KIND_DISTRIBUTE:
-	  pp_string (buffer, "#pragma omp distribute");
-	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git gcc/gimple.c gcc/gimple.c
index 1862de2..6580d10 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -853,8 +853,7 @@  gimple_build_omp_critical (gimple_seq body, tree name)
 
    BODY is sequence of statements inside the for loop.
    KIND is the `for' variant.
-   CLAUSES, are any of the OMP loop construct's clauses: private, firstprivate,
-   lastprivate, reductions, ordered, schedule, and nowait.
+   CLAUSES, are any of the loop construct's clauses.
    COLLAPSE is the collapse count.
    PRE_BODY is the sequence of statements that are loop invariant.  */
 
@@ -1694,6 +1693,7 @@  gimple_copy (gimple stmt)
           gcc_unreachable ();
 
 	case GIMPLE_OMP_FOR:
+	  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
 	  new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt));
 	  gimple_omp_for_set_pre_body (copy, new_seq);
 	  t = unshare_expr (gimple_omp_for_clauses (stmt));
diff --git gcc/gimple.def gcc/gimple.def
index c9756b7..e2e912c 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -267,6 +267,9 @@  DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL)
    for (INDEX = INITIAL; INDEX COND FINAL; INDEX {+=,-=} INCR)
    BODY
 
+   Likewise for:
+   #pragma acc loop [clause1 ... clauseN]
+
    BODY is the loop body.
 
    CLAUSES is the list of clauses.
@@ -293,7 +296,7 @@  DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL)
    INITIAL, FINAL and INCR are required to be loop invariant integer
    expressions that are evaluated without any synchronization.
    The evaluation order, frequency of evaluation and side-effects are
-   unspecified by the standard.  */
+   unspecified by the standards.  */
 DEFGSCODE(GIMPLE_OMP_FOR, "gimple_omp_for", GSS_OMP_FOR)
 
 /* GIMPLE_OMP_MASTER <BODY> represents #pragma omp master.
diff --git gcc/gimple.h gcc/gimple.h
index 34a0bdb..f059789 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -91,15 +91,16 @@  enum gf_mask {
     GF_CALL_ALLOCA_FOR_VAR	= 1 << 5,
     GF_CALL_INTERNAL		= 1 << 6,
     GF_OMP_PARALLEL_COMBINED	= 1 << 0,
-    GF_OMP_FOR_KIND_MASK	= (1 << 2) - 1,
+    GF_OMP_FOR_KIND_MASK	= (1 << 3) - 1,
     GF_OMP_FOR_KIND_FOR		= 0,
     GF_OMP_FOR_KIND_DISTRIBUTE	= 1,
+    GF_OMP_FOR_KIND_OACC_LOOP	= 2,
     /* Flag for SIMD variants of OMP_FOR kinds.  */
-    GF_OMP_FOR_SIMD		= 1 << 1,
+    GF_OMP_FOR_SIMD		= 1 << 2,
     GF_OMP_FOR_KIND_SIMD	= GF_OMP_FOR_SIMD | 0,
     GF_OMP_FOR_KIND_CILKSIMD	= GF_OMP_FOR_SIMD | 1,
-    GF_OMP_FOR_COMBINED		= 1 << 2,
-    GF_OMP_FOR_COMBINED_INTO	= 1 << 3,
+    GF_OMP_FOR_COMBINED		= 1 << 3,
+    GF_OMP_FOR_COMBINED_INTO	= 1 << 4,
     GF_OMP_TARGET_KIND_MASK	= (1 << 2) - 1,
     GF_OMP_TARGET_KIND_REGION	= 0,
     GF_OMP_TARGET_KIND_DATA	= 1,
@@ -4518,7 +4519,7 @@  gimple_omp_critical_set_name (gimple gs, tree name)
 }
 
 
-/* Return the kind of OMP for statemement.  */
+/* Return the kind of the OMP_FOR statemement G.  */
 
 static inline int
 gimple_omp_for_kind (const_gimple g)
@@ -4528,7 +4529,7 @@  gimple_omp_for_kind (const_gimple g)
 }
 
 
-/* Set the OMP for kind.  */
+/* Set the kind of the OMP_FOR statement G.  */
 
 static inline void
 gimple_omp_for_set_kind (gimple g, int kind)
@@ -4539,7 +4540,7 @@  gimple_omp_for_set_kind (gimple g, int kind)
 }
 
 
-/* Return true if OMP for statement G has the
+/* Return true if OMP_FOR statement G has the
    GF_OMP_FOR_COMBINED flag set.  */
 
 static inline bool
@@ -4550,8 +4551,8 @@  gimple_omp_for_combined_p (const_gimple g)
 }
 
 
-/* Set the GF_OMP_FOR_COMBINED field in G depending on the boolean
-   value of COMBINED_P.  */
+/* Set the GF_OMP_FOR_COMBINED field in the OMP_FOR statement G depending on
+   the boolean value of COMBINED_P.  */
 
 static inline void
 gimple_omp_for_set_combined_p (gimple g, bool combined_p)
@@ -4564,7 +4565,7 @@  gimple_omp_for_set_combined_p (gimple g, bool combined_p)
 }
 
 
-/* Return true if OMP for statement G has the
+/* Return true if the OMP_FOR statement G has the
    GF_OMP_FOR_COMBINED_INTO flag set.  */
 
 static inline bool
@@ -4575,8 +4576,8 @@  gimple_omp_for_combined_into_p (const_gimple g)
 }
 
 
-/* Set the GF_OMP_FOR_COMBINED_INTO field in G depending on the boolean
-   value of COMBINED_P.  */
+/* Set the GF_OMP_FOR_COMBINED_INTO field in the OMP_FOR statement G depending
+   on the boolean value of COMBINED_P.  */
 
 static inline void
 gimple_omp_for_set_combined_into_p (gimple g, bool combined_p)
@@ -4589,7 +4590,7 @@  gimple_omp_for_set_combined_into_p (gimple g, bool combined_p)
 }
 
 
-/* Return the clauses associated with OMP_FOR GS.  */
+/* Return the clauses associated with the OMP_FOR statement GS.  */
 
 static inline tree
 gimple_omp_for_clauses (const_gimple gs)
@@ -4600,7 +4601,8 @@  gimple_omp_for_clauses (const_gimple gs)
 }
 
 
-/* Return a pointer to the OMP_FOR GS.  */
+/* Return a pointer to the clauses associated with the OMP_FOR statement
+   GS.  */
 
 static inline tree *
 gimple_omp_for_clauses_ptr (gimple gs)
@@ -4611,7 +4613,8 @@  gimple_omp_for_clauses_ptr (gimple gs)
 }
 
 
-/* Set CLAUSES to be the list of clauses associated with OMP_FOR GS.  */
+/* Set CLAUSES to be the list of clauses associated with the OMP_FOR statement
+   GS.  */
 
 static inline void
 gimple_omp_for_set_clauses (gimple gs, tree clauses)
@@ -4622,7 +4625,7 @@  gimple_omp_for_set_clauses (gimple gs, tree clauses)
 }
 
 
-/* Get the collapse count of OMP_FOR GS.  */
+/* Get the collapse count of the OMP_FOR statement GS.  */
 
 static inline size_t
 gimple_omp_for_collapse (gimple gs)
@@ -4633,7 +4636,7 @@  gimple_omp_for_collapse (gimple gs)
 }
 
 
-/* Return the index variable for OMP_FOR GS.  */
+/* Return the index variable for the OMP_FOR statement GS.  */
 
 static inline tree
 gimple_omp_for_index (const_gimple gs, size_t i)
@@ -4645,7 +4648,7 @@  gimple_omp_for_index (const_gimple gs, size_t i)
 }
 
 
-/* Return a pointer to the index variable for OMP_FOR GS.  */
+/* Return a pointer to the index variable for the OMP_FOR statement GS.  */
 
 static inline tree *
 gimple_omp_for_index_ptr (gimple gs, size_t i)
@@ -4657,7 +4660,7 @@  gimple_omp_for_index_ptr (gimple gs, size_t i)
 }
 
 
-/* Set INDEX to be the index variable for OMP_FOR GS.  */
+/* Set INDEX to be the index variable for the OMP_FOR statement GS.  */
 
 static inline void
 gimple_omp_for_set_index (gimple gs, size_t i, tree index)
@@ -4669,7 +4672,7 @@  gimple_omp_for_set_index (gimple gs, size_t i, tree index)
 }
 
 
-/* Return the initial value for OMP_FOR GS.  */
+/* Return the initial value for the OMP_FOR statement GS.  */
 
 static inline tree
 gimple_omp_for_initial (const_gimple gs, size_t i)
@@ -4681,7 +4684,7 @@  gimple_omp_for_initial (const_gimple gs, size_t i)
 }
 
 
-/* Return a pointer to the initial value for OMP_FOR GS.  */
+/* Return a pointer to the initial value for the OMP_FOR statement GS.  */
 
 static inline tree *
 gimple_omp_for_initial_ptr (gimple gs, size_t i)
@@ -4693,7 +4696,7 @@  gimple_omp_for_initial_ptr (gimple gs, size_t i)
 }
 
 
-/* Set INITIAL to be the initial value for OMP_FOR GS.  */
+/* Set INITIAL to be the initial value for the OMP_FOR statement GS.  */
 
 static inline void
 gimple_omp_for_set_initial (gimple gs, size_t i, tree initial)
@@ -4705,7 +4708,7 @@  gimple_omp_for_set_initial (gimple gs, size_t i, tree initial)
 }
 
 
-/* Return the final value for OMP_FOR GS.  */
+/* Return the final value for the OMP_FOR statement GS.  */
 
 static inline tree
 gimple_omp_for_final (const_gimple gs, size_t i)
@@ -4717,7 +4720,7 @@  gimple_omp_for_final (const_gimple gs, size_t i)
 }
 
 
-/* Return a pointer to the final value for OMP_FOR GS.  */
+/* Return a pointer to the final value for the OMP_FOR statement GS.  */
 
 static inline tree *
 gimple_omp_for_final_ptr (gimple gs, size_t i)
@@ -4729,7 +4732,7 @@  gimple_omp_for_final_ptr (gimple gs, size_t i)
 }
 
 
-/* Set FINAL to be the final value for OMP_FOR GS.  */
+/* Set FINAL to be the final value for the OMP_FOR statement GS.  */
 
 static inline void
 gimple_omp_for_set_final (gimple gs, size_t i, tree final)
@@ -4741,7 +4744,32 @@  gimple_omp_for_set_final (gimple gs, size_t i, tree final)
 }
 
 
-/* Return the increment value for OMP_FOR GS.  */
+/* Set COND to be the condition code for the OMP_FOR statement GS.  */
+
+static inline void
+gimple_omp_for_set_cond (gimple gs, size_t i, enum tree_code cond)
+{
+  gimple_statement_omp_for *omp_for_stmt =
+    as_a <gimple_statement_omp_for> (gs);
+  gcc_gimple_checking_assert (TREE_CODE_CLASS (cond) == tcc_comparison
+			      && i < omp_for_stmt->collapse);
+  omp_for_stmt->iter[i].cond = cond;
+}
+
+
+/* Return the condition code associated with the OMP_FOR statement GS.  */
+
+static inline enum tree_code
+gimple_omp_for_cond (const_gimple gs, size_t i)
+{
+  const gimple_statement_omp_for *omp_for_stmt =
+    as_a <const gimple_statement_omp_for> (gs);
+  gcc_gimple_checking_assert (i < omp_for_stmt->collapse);
+  return omp_for_stmt->iter[i].cond;
+}
+
+
+/* Return the increment value for the OMP_FOR statement GS.  */
 
 static inline tree
 gimple_omp_for_incr (const_gimple gs, size_t i)
@@ -4753,7 +4781,7 @@  gimple_omp_for_incr (const_gimple gs, size_t i)
 }
 
 
-/* Return a pointer to the increment value for OMP_FOR GS.  */
+/* Return a pointer to the increment value for the OMP_FOR statement GS.  */
 
 static inline tree *
 gimple_omp_for_incr_ptr (gimple gs, size_t i)
@@ -4765,7 +4793,7 @@  gimple_omp_for_incr_ptr (gimple gs, size_t i)
 }
 
 
-/* Set INCR to be the increment value for OMP_FOR GS.  */
+/* Set INCR to be the increment value for the OMP_FOR statement GS.  */
 
 static inline void
 gimple_omp_for_set_incr (gimple gs, size_t i, tree incr)
@@ -5470,31 +5498,6 @@  gimple_omp_sections_set_control (gimple gs, tree control)
 }
 
 
-/* Set COND to be the condition code for OMP_FOR GS.  */
-
-static inline void
-gimple_omp_for_set_cond (gimple gs, size_t i, enum tree_code cond)
-{
-  gimple_statement_omp_for *omp_for_stmt =
-    as_a <gimple_statement_omp_for> (gs);
-  gcc_gimple_checking_assert (TREE_CODE_CLASS (cond) == tcc_comparison
-			      && i < omp_for_stmt->collapse);
-  omp_for_stmt->iter[i].cond = cond;
-}
-
-
-/* Return the condition code associated with OMP_FOR GS.  */
-
-static inline enum tree_code
-gimple_omp_for_cond (const_gimple gs, size_t i)
-{
-  const gimple_statement_omp_for *omp_for_stmt =
-    as_a <const gimple_statement_omp_for> (gs);
-  gcc_gimple_checking_assert (i < omp_for_stmt->collapse);
-  return omp_for_stmt->iter[i].cond;
-}
-
-
 /* Set the value being stored in an atomic store.  */
 
 static inline void
@@ -5811,6 +5814,14 @@  is_gimple_omp_oacc_specifically (const_gimple stmt)
     case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       return true;
+    case GIMPLE_OMP_FOR:
+      switch (gimple_omp_for_kind (stmt))
+	{
+	case GF_OMP_FOR_KIND_OACC_LOOP:
+	  return true;
+	default:
+	  return false;
+	}      
     case GIMPLE_OMP_TARGET:
       switch (gimple_omp_target_kind (stmt))
 	{
diff --git gcc/gimplify.c gcc/gimplify.c
index 0985bb2..9788f4c 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -4363,8 +4363,8 @@  is_gimple_stmt (tree t)
     case OMP_FOR:
     case OMP_SIMD:
     case CILK_SIMD:
-    case OACC_LOOP:
     case OMP_DISTRIBUTE:
+    case OACC_LOOP:
     case OMP_SECTIONS:
     case OMP_SECTION:
     case OMP_SINGLE:
@@ -6683,14 +6683,36 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
   gimple_seq for_body, for_pre_body;
   int i;
   bool simd;
+  enum gimplify_omp_var_data govd_private;
+  enum omp_region_type ort;
   bitmap has_decl_expr = NULL;
 
   orig_for_stmt = for_stmt = *expr_p;
 
-  simd = TREE_CODE (for_stmt) == OMP_SIMD
-    || TREE_CODE (for_stmt) == CILK_SIMD;
-  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
-			     simd ? ORT_SIMD : ORT_WORKSHARE);
+  switch (TREE_CODE (for_stmt))
+    {
+    case OMP_FOR:
+    case OMP_DISTRIBUTE:
+      simd = false;
+      govd_private = GOVD_PRIVATE;
+      ort = ORT_WORKSHARE;
+      break;
+    case OACC_LOOP:
+      simd = false;
+      govd_private = /* TODO */ GOVD_LOCAL;
+      ort = /* TODO */ ORT_WORKSHARE;
+      break;
+    case OMP_SIMD:
+    case CILK_SIMD:
+      simd = true;
+      govd_private = GOVD_PRIVATE;
+      ort = ORT_SIMD;
+      break;
+    default:
+      gcc_unreachable ();
+    }
+
+  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
 
   /* Handle OMP_FOR_INIT.  */
   for_pre_body = NULL;
@@ -6722,6 +6744,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 
   if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
     {
+      gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
       for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
 			    NULL, NULL);
       gcc_assert (for_stmt != NULL_TREE);
@@ -6742,7 +6765,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
 		  || POINTER_TYPE_P (TREE_TYPE (decl)));
 
-      /* Make sure the iteration variable is private.  */
+      /* Make sure the iteration variable is some kind of private.  */
       tree c = NULL_TREE;
       if (orig_for_stmt != for_stmt)
 	/* Do this only on innermost construct for combined ones.  */;
@@ -6768,6 +6791,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	    }
 	  else
 	    {
+	      gcc_assert (govd_private == GOVD_PRIVATE);
 	      bool lastprivate
 		= (!has_decl_expr
 		   || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
@@ -6785,7 +6809,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       else if (omp_is_private (gimplify_omp_ctxp, decl, simd))
 	omp_notice_variable (gimplify_omp_ctxp, decl, true);
       else
-	omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
+	omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
 
       /* If DECL is not a gimple register, create a temporary variable to act
 	 as an iteration counter.  This is valid, since DECL cannot be
@@ -6799,7 +6823,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 
 	  gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
 
-	  omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
+	  omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
 	}
       else
 	var = decl;
@@ -6936,7 +6960,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
 	decl = TREE_OPERAND (t, 0);
 	var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
-	omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
+	omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
 	TREE_OPERAND (t, 0) = var;
 	t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
 	TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
@@ -6952,6 +6976,7 @@  gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
     case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
     case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
+    case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
     default:
       gcc_unreachable ();
     }
@@ -8048,7 +8073,6 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case OACC_EXIT_DATA:
 	case OACC_WAIT:
 	case OACC_CACHE:
-	case OACC_LOOP:
 	  sorry ("directive not yet implemented");
 	  ret = GS_ALL_DONE;
 	  break;
@@ -8067,6 +8091,7 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case OMP_SIMD:
 	case CILK_SIMD:
 	case OMP_DISTRIBUTE:
+	case OACC_LOOP:
 	  ret = gimplify_omp_for (expr_p, pre_p);
 	  break;
 
diff --git gcc/omp-low.c gcc/omp-low.c
index c3b3e95..13373ca 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -177,6 +177,8 @@  typedef struct omp_context
   bool cancellable;
 } omp_context;
 
+/* A structure holding the elements of:
+   for (V = N1; V cond N2; V += STEP) [...] */
 
 struct omp_for_data_loop
 {
@@ -310,9 +312,9 @@  extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd,
   else
     fd->loops = &fd->loop;
 
-  fd->have_nowait = distribute || simd;
+  fd->have_nowait = (gimple_omp_for_kind (for_stmt) != GF_OMP_FOR_KIND_FOR);
   fd->have_ordered = false;
-  fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
+  fd->sched_kind = /* TODO: OACC_LOOP */ OMP_CLAUSE_SCHEDULE_STATIC;
   fd->chunk_size = NULL_TREE;
   collapse_iter = NULL;
   collapse_count = NULL;
@@ -1626,7 +1628,10 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
 	  if (ctx->outer)
+	    {
+	      gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
 	    scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
+	    }
 	  break;
 
 	case OMP_CLAUSE_TO:
@@ -2288,7 +2293,7 @@  scan_omp_task (gimple_stmt_iterator *gsi, omp_context *outer_ctx)
 }
 
 
-/* Scan an OpenMP loop directive.  */
+/* Scan a GIMPLE_OMP_FOR.  */
 
 static void
 scan_omp_for (gimple stmt, omp_context *outer_ctx)
@@ -2421,6 +2426,10 @@  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
   if (is_gimple_omp (stmt)
       && is_gimple_omp_oacc_specifically (stmt))
     {
+      /* Regular handling of OpenACC loop constructs.  */
+      if (gimple_code (stmt) == GIMPLE_OMP_FOR
+	  && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+	goto cont;
       /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different
 	 from an OpenACC data construct.  */
       for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
@@ -2447,6 +2456,7 @@  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 	    return false;
 	  }
     }
+ cont:
 
   if (ctx != NULL)
     {
@@ -2626,6 +2636,8 @@  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 		      "of work-sharing, critical, ordered, master or explicit "
 		      "task region");
 	    return false;
+	  case GIMPLE_OACC_KERNELS:
+	  case GIMPLE_OACC_PARALLEL:
 	  case GIMPLE_OMP_PARALLEL:
 	    return true;
 	  default:
@@ -3217,8 +3229,6 @@  static void
 lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
 			 omp_context *ctx, struct omp_for_data *fd)
 {
-  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
-
   tree c, dtor, copyin_seq, x, ptr;
   bool copyin_by_ref = false;
   bool lastprivate_firstprivate = false;
@@ -3920,8 +3930,6 @@  static void
 lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 			   omp_context *ctx)
 {
-  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
-
   tree x, c, label = NULL, orig_clauses = clauses;
   bool par_clauses = false;
   tree simduid = NULL, lastlane = NULL;
@@ -4057,8 +4065,6 @@  lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 static void
 lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
 {
-  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
-
   gimple_seq sub_seq = NULL;
   gimple stmt;
   tree x, c;
@@ -5849,6 +5855,8 @@  expand_omp_for_generic (struct omp_region *region,
 			enum built_in_function next_fn,
 			gimple inner_stmt)
 {
+  gcc_assert (gimple_omp_for_kind (fd->for_stmt) != GF_OMP_FOR_KIND_OACC_LOOP);
+
   tree type, istart0, iend0, iend;
   tree t, vmain, vback, bias = NULL_TREE;
   basic_block entry_bb, cont_bb, exit_bb, l0_bb, l1_bb, collapse_bb;
@@ -5918,6 +5926,9 @@  expand_omp_for_generic (struct omp_region *region,
   gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
   if (fd->collapse > 1)
     {
+      gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       int first_zero_iter = -1;
       basic_block zero_iter_bb = NULL, l2_dom_bb = NULL;
 
@@ -5946,6 +5957,9 @@  expand_omp_for_generic (struct omp_region *region,
     }
   if (in_combined_parallel)
     {
+      gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       /* In a combined parallel loop, emit a call to
 	 GOMP_loop_foo_next.  */
       t = build_call_expr (builtin_decl_explicit (next_fn), 2,
@@ -5964,6 +5978,9 @@  expand_omp_for_generic (struct omp_region *region,
       t0 = fd->loop.n1;
       if (gimple_omp_for_combined_into_p (fd->for_stmt))
 	{
+	  gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+		      != GF_OMP_FOR_KIND_OACC_LOOP);
+
 	  tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
 					 OMP_CLAUSE__LOOPTEMP_);
 	  gcc_assert (innerc);
@@ -6276,12 +6293,14 @@  expand_omp_for_static_nochunk (struct omp_region *region,
   gimple_stmt_iterator gsi;
   gimple stmt;
   edge ep;
-  enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
-  enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
   bool broken_loop = region->cont == NULL;
   tree *counts = NULL;
   tree n1, n2, step;
 
+  gcc_assert ((gimple_omp_for_kind (fd->for_stmt)
+	       != GF_OMP_FOR_KIND_OACC_LOOP)
+	      || !inner_stmt);
+
   itype = type = TREE_TYPE (fd->loop.v);
   if (POINTER_TYPE_P (type))
     itype = signed_type_for (type);
@@ -6305,14 +6324,11 @@  expand_omp_for_static_nochunk (struct omp_region *region,
   gsi = gsi_last_bb (entry_bb);
   gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
 
-  if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
-    {
-      get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
-      get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
-    }
-
   if (fd->collapse > 1)
     {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       int first_zero_iter = -1;
       basic_block l2_dom_bb = NULL;
 
@@ -6323,7 +6339,12 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       t = NULL_TREE;
     }
   else if (gimple_omp_for_combined_into_p (fd->for_stmt))
+    {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
     t = integer_one_node;
+    }
   else
     t = fold_binary (fd->loop.cond_code, boolean_type_node,
 		     fold_convert (type, fd->loop.n1),
@@ -6357,6 +6378,9 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       ep->probability = REG_BR_PROB_BASE / 2000 - 1;
       if (gimple_in_ssa_p (cfun))
 	{
+	  gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		      != GF_OMP_FOR_KIND_OACC_LOOP);
+
 	  int dest_idx = find_edge (entry_bb, fin_bb)->dest_idx;
 	  for (gsi = gsi_start_phis (fin_bb);
 	       !gsi_end_p (gsi); gsi_next (&gsi))
@@ -6369,14 +6393,32 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       gsi = gsi_last_bb (entry_bb);
     }
 
-  t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
-  t = fold_convert (itype, t);
-  nthreads = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+  switch (gimple_omp_for_kind (fd->for_stmt))
+    {
+    case GF_OMP_FOR_KIND_FOR:
+      nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+      nthreads = build_call_expr (nthreads, 0);
+      threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+      threadid = build_call_expr (threadid, 0);
+      break;
+    case GF_OMP_FOR_KIND_DISTRIBUTE:
+      nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+      nthreads = build_call_expr (nthreads, 0);
+      threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+      threadid = build_call_expr (threadid, 0);
+      break;
+    case GF_OMP_FOR_KIND_OACC_LOOP:
+      nthreads = integer_one_node;
+      threadid = integer_zero_node;
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  nthreads = fold_convert (itype, nthreads);
+  nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
 				       true, GSI_SAME_STMT);
-
-  t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
-  t = fold_convert (itype, t);
-  threadid = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+  threadid = fold_convert (itype, threadid);
+  threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
 				       true, GSI_SAME_STMT);
 
   n1 = fd->loop.n1;
@@ -6384,6 +6426,9 @@  expand_omp_for_static_nochunk (struct omp_region *region,
   step = fd->loop.step;
   if (gimple_omp_for_combined_into_p (fd->for_stmt))
     {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
 				     OMP_CLAUSE__LOOPTEMP_);
       gcc_assert (innerc);
@@ -6462,6 +6507,9 @@  expand_omp_for_static_nochunk (struct omp_region *region,
 
   if (gimple_omp_for_combined_p (fd->for_stmt))
     {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       tree clauses = gimple_code (inner_stmt) == GIMPLE_OMP_PARALLEL
 		     ? gimple_omp_parallel_clauses (inner_stmt)
 		     : gimple_omp_for_clauses (inner_stmt);
@@ -6502,7 +6550,12 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
     }
   if (fd->collapse > 1)
+    {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
     expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar);
+    }
 
   if (!broken_loop)
     {
@@ -6537,13 +6590,21 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       gsi_remove (&gsi, true);
 
       if (fd->collapse > 1 && !gimple_omp_for_combined_p (fd->for_stmt))
+	{
+	  gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		      != GF_OMP_FOR_KIND_OACC_LOOP);
+
 	collapse_bb = extract_omp_for_update_vars (fd, cont_bb, body_bb);
+	}
     }
 
   /* Replace the GIMPLE_OMP_RETURN with a barrier, or nothing.  */
   gsi = gsi_last_bb (exit_bb);
   if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
     {
+      gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		  != GF_OMP_FOR_KIND_OACC_LOOP);
+
       t = gimple_omp_return_lhs (gsi_stmt (gsi));
       gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
     }
@@ -6563,11 +6624,17 @@  expand_omp_for_static_nochunk (struct omp_region *region,
       ep = find_edge (cont_bb, body_bb);
       if (gimple_omp_for_combined_p (fd->for_stmt))
 	{
+	  gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		      != GF_OMP_FOR_KIND_OACC_LOOP);
+
 	  remove_edge (ep);
 	  ep = NULL;
 	}
       else if (fd->collapse > 1)
 	{
+	  gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+		      != GF_OMP_FOR_KIND_OACC_LOOP);
+
 	  remove_edge (ep);
 	  ep = make_edge (cont_bb, collapse_bb, EDGE_TRUE_VALUE);
 	}
@@ -6639,6 +6706,8 @@  static void
 expand_omp_for_static_chunk (struct omp_region *region,
 			     struct omp_for_data *fd, gimple inner_stmt)
 {
+  gcc_assert (gimple_omp_for_kind (fd->for_stmt) != GF_OMP_FOR_KIND_OACC_LOOP);
+
   tree n, s0, e0, e, t;
   tree trip_var, trip_init, trip_main, trip_back, nthreads, threadid;
   tree type, itype, v_main, v_back, v_extra;
@@ -6647,8 +6716,6 @@  expand_omp_for_static_chunk (struct omp_region *region,
   gimple_stmt_iterator si;
   gimple stmt;
   edge se;
-  enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
-  enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
   bool broken_loop = region->cont == NULL;
   tree *counts = NULL;
   tree n1, n2, step;
@@ -6680,12 +6747,6 @@  expand_omp_for_static_chunk (struct omp_region *region,
   si = gsi_last_bb (entry_bb);
   gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_FOR);
 
-  if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
-    {
-      get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
-      get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
-    }
-
   if (fd->collapse > 1)
     {
       int first_zero_iter = -1;
@@ -6744,14 +6805,28 @@  expand_omp_for_static_chunk (struct omp_region *region,
       si = gsi_last_bb (entry_bb);
     }
 
-  t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
-  t = fold_convert (itype, t);
-  nthreads = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+  switch (gimple_omp_for_kind (fd->for_stmt))
+    {
+    case GF_OMP_FOR_KIND_FOR:
+      nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+      nthreads = build_call_expr (nthreads, 0);
+      threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+      threadid = build_call_expr (threadid, 0);
+      break;
+    case GF_OMP_FOR_KIND_DISTRIBUTE:
+      nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+      nthreads = build_call_expr (nthreads, 0);
+      threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+      threadid = build_call_expr (threadid, 0);
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  nthreads = fold_convert (itype, nthreads);
+  nthreads = force_gimple_operand_gsi (&si, nthreads, true, NULL_TREE,
 				       true, GSI_SAME_STMT);
-
-  t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
-  t = fold_convert (itype, t);
-  threadid = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+  threadid = fold_convert (itype, threadid);
+  threadid = force_gimple_operand_gsi (&si, threadid, true, NULL_TREE,
 				       true, GSI_SAME_STMT);
 
   n1 = fd->loop.n1;
@@ -9211,8 +9286,6 @@  lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 static void
 maybe_add_implicit_barrier_cancel (omp_context *ctx, gimple_seq *body)
 {
-  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
-
   gimple omp_return = gimple_seq_last_stmt (*body);
   gcc_assert (gimple_code (omp_return) == GIMPLE_OMP_RETURN);
   if (gimple_omp_return_nowait_p (omp_return))
@@ -9792,6 +9865,8 @@  lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
   if (gimple_omp_for_combined_into_p (stmt))
     {
+      gcc_assert (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_OACC_LOOP);
+
       extract_omp_for_data (stmt, &fd, NULL);
       fdp = &fd;
 
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 13e99d5..78882c0 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,12 @@ 
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* c-c++-common/goacc-gomp/nesting-1.c: New file.
+	* c-c++-common/goacc-gomp/nesting-fail-1.c: Extend.
+	* c-c++-common/goacc/clauses-fail.c: Likewise.
+	* c-c++-common/goacc/nesting-1.c: Likewise.
+	* gcc.dg/goacc/sb-1.c: Likewise.
+	* gcc.dg/goacc/sb-3.c: New file.
+
 	* c-c++-common/goacc/nesting-1.c: New file.
 	* c-c++-common/goacc/nesting-data-1.c: Likewise.
 	* c-c++-common/goacc/nesting-fail-1.c: Update.
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
new file mode 100644
index 0000000..df45bcf
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -0,0 +1,12 @@ 
+void
+f_omp_parallel (void)
+{
+#pragma omp parallel
+  {
+    int i;
+
+#pragma acc loop
+    for (i = 0; i < 2; ++i)
+      ;
+  }
+}
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 14103a6..871fab3 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -24,6 +24,9 @@  f_omp (void)
       ;
 #pragma acc data	/* { dg-error "may not be nested" } */
       ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+      for (i = 0; i < 2; ++i)
+	;
     }
 
 #pragma omp sections
@@ -42,6 +45,12 @@  f_omp (void)
 #pragma acc data	/* { dg-error "may not be nested" } */
       ;
     }
+#pragma omp section
+    {
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+      for (i = 0; i < 2; ++i)
+	;
+    }
   }
 
 #pragma omp single
@@ -52,6 +61,9 @@  f_omp (void)
     ;
 #pragma acc data	/* { dg-error "may not be nested" } */
     ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 
 #pragma omp task
@@ -62,6 +74,9 @@  f_omp (void)
     ;
 #pragma acc data	/* { dg-error "may not be nested" } */
     ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 
 #pragma omp master
@@ -72,6 +87,9 @@  f_omp (void)
     ;
 #pragma acc data	/* { dg-error "may not be nested" } */
     ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 
 #pragma omp critical
@@ -82,6 +100,9 @@  f_omp (void)
     ;
 #pragma acc data	/* { dg-error "may not be nested" } */
     ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 
 #pragma omp ordered
@@ -92,6 +113,9 @@  f_omp (void)
     ;
 #pragma acc data	/* { dg-error "may not be nested" } */
     ;
+#pragma acc loop	/* { dg-error "may not be closely nested" } */
+    for (i = 0; i < 2; ++i)
+      ;
   }
 }
 
@@ -289,3 +313,77 @@  f_acc_data (void)
     ;
   }
 }
+
+/* TODO: Some of these should either be allowed or fail with a more sensible
+   error message.  */
+void
+f_acc_loop (void)
+{
+  int i;
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp parallel	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp for		/* { dg-error "may not be nested" } */
+      for (i = 0; i < 3; i++)
+	;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp sections	/* { dg-error "may not be nested" } */
+      {
+	;
+      }
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp single	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp task	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp master	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp critical	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp atomic write
+      i = 0;		/* { dg-error "may not be nested" } */
+    }
+
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp ordered	/* { dg-error "may not be nested" } */
+      ;
+    }
+}
diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index 133bf81..e8e1278 100644
--- gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -1,6 +1,8 @@ 
 void
 f (void)
 {
+  int i;
+
 #pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
   ;
 
@@ -9,4 +11,8 @@  f (void)
 
 #pragma acc data two /* { dg-error "expected clause before 'two'" } */
   ;
+
+#pragma acc loop deux /* { dg-error "expected clause before 'deux'" } */
+  for (i = 0; i < 2; ++i)
+    ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c
index 3a22292..a489d2d 100644
--- gcc/testsuite/c-c++-common/goacc/nesting-1.c
+++ gcc/testsuite/c-c++-common/goacc/nesting-1.c
@@ -1,13 +1,56 @@ 
 void
+f_acc_parallel (void)
+{
+#pragma acc parallel
+  {
+    int i;
+
+#pragma acc loop
+    for (i = 0; i < 2; ++i)
+      ;
+  }
+}
+
+
+void
+f_acc_kernels (void)
+{
+#pragma acc kernels
+  {
+    int i;
+
+#pragma acc loop
+    for (i = 0; i < 2; ++i)
+      ;
+  }
+}
+
+
+void
 f_acc_data (void)
 {
 #pragma acc data
   {
+    int i;
+
 #pragma acc parallel
     ;
+
+#pragma acc parallel
+    {
+#pragma acc loop
+      for (i = 0; i < 2; ++i)
+	;
+    }
+
 #pragma acc kernels
     ;
+
 #pragma acc data
     ;
+
+#pragma acc loop
+    for (i = 0; i < 2; ++i)
+      ;
   }
 }
diff --git gcc/testsuite/gcc.dg/goacc/sb-1.c gcc/testsuite/gcc.dg/goacc/sb-1.c
index 24c88fe..bcb7272 100644
--- gcc/testsuite/gcc.dg/goacc/sb-1.c
+++ gcc/testsuite/gcc.dg/goacc/sb-1.c
@@ -2,6 +2,8 @@ 
 
 void foo()
 {
+  int l;
+
   bad1:
   #pragma acc parallel
     goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
@@ -9,6 +11,9 @@  void foo()
     goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
   #pragma acc data
     goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
+  #pragma acc loop
+    for (l = 0; l < 2; ++l)
+      goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
 
   goto bad2_parallel; // { dg-error "invalid entry to OpenACC structured block" }
   #pragma acc parallel
@@ -28,6 +33,13 @@  void foo()
       bad2_data: ;
     }
 
+  goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
+  #pragma acc loop
+  for (l = 0; l < 2; ++l)
+    {
+      bad2_loop: ;
+    }
+
   #pragma acc parallel
     {
       int i;
@@ -51,4 +63,13 @@  void foo()
       for (i = 0; i < 10; ++i)
 	{ ok1_data: break; }
     }
+
+  #pragma acc loop
+    for (l = 0; l < 2; ++l)
+      {
+	int i;
+	goto ok1_loop;
+	for (i = 0; i < 10; ++i)
+	  { ok1_loop: break; }
+      }
 }
diff --git gcc/testsuite/gcc.dg/goacc/sb-3.c gcc/testsuite/gcc.dg/goacc/sb-3.c
new file mode 100644
index 0000000..6c2926c
--- /dev/null
+++ gcc/testsuite/gcc.dg/goacc/sb-3.c
@@ -0,0 +1,18 @@ 
+// { dg-do compile }
+
+void f (void)
+{
+  int i, j;
+#pragma acc loop
+  for(i = 1; i < 30; i++)
+    {
+      if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }
+#pragma acc loop // { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+      for(j = 5; j < 10; j++)
+	{
+	  if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }
+	}
+    }
+ out:
+  ;
+}
diff --git gcc/tree-inline.c gcc/tree-inline.c
index cdfe35c..5cfda33 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1342,6 +1342,7 @@  remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	  break;
 
 	case GIMPLE_OMP_FOR:
+	  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id);
 	  copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt),
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 397f851..e8ba1e3 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -622,6 +622,8 @@  walk_gimple_omp_for (gimple for_stmt,
     		     walk_stmt_fn callback_stmt, walk_tree_fn callback_op,
     		     struct nesting_info *info)
 {
+  gcc_assert (!is_gimple_omp_oacc_specifically (for_stmt));
+
   struct walk_stmt_info wi;
   gimple_seq seq;
   tree t;
@@ -1282,6 +1284,7 @@  convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_nonlocal_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
       walk_gimple_omp_for (stmt, convert_nonlocal_reference_stmt,
@@ -1746,6 +1749,7 @@  convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_local_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
       walk_gimple_omp_for (stmt, convert_local_reference_stmt,
@@ -2178,6 +2182,7 @@  convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       walk_body (convert_gimple_call, NULL, info,
 	  	 gimple_omp_for_pre_body_ptr (stmt));
       /* FALLTHRU */
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index 49e5f6c..6c311790 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -2538,14 +2538,14 @@  dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
       pp_string (buffer, "#pragma simd");
       goto dump_omp_loop;
 
-    case OACC_LOOP:
-      pp_string (buffer, "#pragma acc loop");
-      goto dump_omp_loop;
-
     case OMP_DISTRIBUTE:
       pp_string (buffer, "#pragma omp distribute");
       goto dump_omp_loop;
 
+    case OACC_LOOP:
+      pp_string (buffer, "#pragma acc loop");
+      goto dump_omp_loop;
+
     case OMP_TEAMS:
       pp_string (buffer, "#pragma omp teams");
       dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags);
diff --git gcc/tree.def gcc/tree.def
index d9e4eb41..a9916f4 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1065,7 +1065,7 @@  DEFTREECODE (OMP_TASK, "omp_task", tcc_statement, 2)
    private.  N1, N2 and INCR are required to be loop invariant integer
    expressions that are evaluated without any synchronization.
    The evaluation order, frequency of evaluation and side-effects are
-   unspecified by the standard.  */
+   unspecified by the standards.  */
 DEFTREECODE (OMP_FOR, "omp_for", tcc_statement, 6)
 
 /* OpenMP - #pragma omp simd [clause1 ... clauseN]
@@ -1076,14 +1076,14 @@  DEFTREECODE (OMP_SIMD, "omp_simd", tcc_statement, 6)
    Operands like for OMP_FOR.  */
 DEFTREECODE (CILK_SIMD, "cilk_simd", tcc_statement, 6)
 
-/* OpenACC - #pragma acc loop [clause1 ... clauseN]
-   Operands like for OMP_FOR.  */
-DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6)
-
 /* OpenMP - #pragma omp distribute [clause1 ... clauseN]
    Operands like for OMP_FOR.  */
 DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
 
+/* OpenMP - #pragma acc loop [clause1 ... clauseN]
+   Operands like for OMP_FOR.  */
+DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6)
+
 /* OpenMP - #pragma omp teams [clause1 ... clauseN]
    Operand 0: OMP_TEAMS_BODY: Teams body.
    Operand 1: OMP_TEAMS_CLAUSES: List of clauses.  */
diff --git gcc/tree.h gcc/tree.h
index 6668895..196ec3e 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1210,7 +1210,7 @@  extern void protected_set_expr_location (tree, location_t);
 #define OMP_TASKREG_BODY(NODE)    TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 0)
 #define OMP_TASKREG_CLAUSES(NODE) TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 1)
 
-#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OMP_DISTRIBUTE)
+#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OACC_LOOP)
 #define OMP_FOR_BODY(NODE)	   TREE_OPERAND (OMP_LOOP_CHECK (NODE), 0)
 #define OMP_FOR_CLAUSES(NODE)	   TREE_OPERAND (OMP_LOOP_CHECK (NODE), 1)
 #define OMP_FOR_INIT(NODE)	   TREE_OPERAND (OMP_LOOP_CHECK (NODE), 2)