omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop nesting.

gcc/
	* omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop
	nesting.

	testsuite/
	* c-c++-common/goacc/clauses-fail.c: Adjust errors.
	* c-c++-common/goacc/sb-1.c: Adjust errors.
	* c-c++-common/goacc/sb-3.c: Adjust errors.
	* c-c++-common/goacc/loop-1.c: Adjust errors.
	* c-c++-common/goacc/nesting-1.c: Adjust errors.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors.
	* c-c++-common/goacc-gomp/nesting-1.c: Adjust errors.

From-SVN: r229129

gcc/ChangeLog

@@ -1,3 +1,8 @@
+2015-10-21  Nathan Sidwell  <nathan@codesourcery.com>
+
+	* omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop
+	nesting.
+
 2015-10-21  Ilya Enkovich  <enkovich.gnu@gmail.com>
 
 	* doc/tm.texi: Regenerated.

gcc/omp-low.c

@@ -3178,6 +3178,43 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
       /* We split taskloop into task and nested taskloop in it.  */
       if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP)
 	return true;
+      if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+	{
+	  bool ok = false;
+	  
+	  if (ctx)
+	    switch (gimple_code (ctx->stmt))
+	      {
+	      case GIMPLE_OMP_FOR:
+		ok = (gimple_omp_for_kind (ctx->stmt)
+		      == GF_OMP_FOR_KIND_OACC_LOOP);
+		break;
+
+	      case GIMPLE_OMP_TARGET:
+		switch (gimple_omp_target_kind (ctx->stmt))
+		  {
+		  case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+		  case GF_OMP_TARGET_KIND_OACC_KERNELS:
+		    ok = true;
+		    break;
+
+		  default:
+		    break;
+		  }
+
+	      default:
+		break;
+	      }
+	  else if (get_oacc_fn_attrib (current_function_decl))
+	    ok = true;
+	  if (!ok)
+	    {
+	      error_at (gimple_location (stmt),
+			"OpenACC loop directive must be associated with"
+			" an OpenACC compute region");
+	      return false;
+	    }
+	}
       /* FALLTHRU */
     case GIMPLE_CALL:
       if (is_gimple_call (stmt)

gcc/testsuite/ChangeLog

@@ -1,3 +1,13 @@
+2015-10-21  Nathan Sidwell  <nathan@codesourcery.com>
+
+	* c-c++-common/goacc/clauses-fail.c: Adjust errors.
+	* c-c++-common/goacc/sb-1.c: Adjust errors.
+	* c-c++-common/goacc/sb-3.c: Adjust errors.
+	* c-c++-common/goacc/loop-1.c: Adjust errors.
+	* c-c++-common/goacc/nesting-1.c: Adjust errors.
+	* c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors.
+	* c-c++-common/goacc-gomp/nesting-1.c: Adjust errors.
+
 2015-10-21  Ilya Enkovich  <enkovich.gnu@gmail.com>
 
 	* g++.dg/ext/vector22.C: Allow VEC_COND_EXPR.

gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c

@@ -5,7 +5,7 @@ f_omp_parallel (void)
   {
     int i;
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }

gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c

@@ -28,7 +28,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
       for (i = 0; i < 2; ++i)
 	;
     }
@@ -63,7 +63,7 @@ f_omp (void)
     }
 #pragma omp section
     {
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
       for (i = 0; i < 2; ++i)
 	;
     }
@@ -80,7 +80,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -96,7 +96,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -112,7 +112,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -128,7 +128,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -144,7 +144,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -160,7 +160,7 @@ f_omp (void)
 #pragma acc update host(i) /* { dg-error "OpenACC update construct inside of OpenMP target region" } */
 #pragma acc enter data copyin(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */
 #pragma acc exit data delete(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }
@@ -379,6 +379,7 @@ f_acc_data (void)
 void
 f_acc_loop (void)
 {
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -386,6 +387,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -394,6 +396,7 @@ f_acc_loop (void)
 	;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -403,6 +406,7 @@ f_acc_loop (void)
       }
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -410,6 +414,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -417,6 +422,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -424,6 +430,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -431,6 +438,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -438,6 +446,7 @@ f_acc_loop (void)
       i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -445,6 +454,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {

gcc/testsuite/c-c++-common/goacc/clauses-fail.c

@@ -12,6 +12,7 @@ f (void)
 #pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */
   ;
 
+#pragma acc parallel
 #pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */
   for (i = 0; i < 2; ++i)
     ;

gcc/testsuite/c-c++-common/goacc/loop-1.c

@@ -38,16 +38,16 @@ int test1()
       i = d;
       a[i] = 1;
     }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 1; i < 30; i++ )
     if (i == 16) break; /* { dg-error "break statement used" } */
 
 /* different types of for loop are allowed */
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 1; i < 10; i++)
     {
     }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 1; i < 10; i+=2)
     {
       a[i] = i;

gcc/testsuite/c-c++-common/goacc/nesting-1.c

@@ -58,7 +58,7 @@ f_acc_data (void)
 
 #pragma acc exit data delete(i)
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
 
@@ -93,7 +93,7 @@ f_acc_data (void)
 
 #pragma acc exit data delete(i)
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
       for (i = 0; i < 2; ++i)
 	;
     }

gcc/testsuite/c-c++-common/goacc/sb-1.c

@@ -11,7 +11,7 @@ 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
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (l = 0; l < 2; ++l)
       goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
 
@@ -34,7 +34,7 @@ void foo()
     }
 
   goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (l = 0; l < 2; ++l)
     {
       bad2_loop: ;
@@ -64,7 +64,7 @@ void foo()
 	{ ok1_data: break; }
     }
 
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (l = 0; l < 2; ++l)
       {
 	int i;

gcc/testsuite/c-c++-common/goacc/sb-3.c

@@ -3,11 +3,11 @@
 void f (void)
 {
   int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   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" }
+#pragma acc loop
       for(j = 5; j < 10; j++)
 	{
 	  if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }