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

	gcc/
	* omp-low.c (check_omp_nesting_restrictions): Check for acc loops not
	associated with acc regions or routines.

	gcc/testsuite/
	* c-c++-common/goacc/non-routine.c: New test.
	* c-c++-common/goacc-gomp/nesting-1.c: Add checks for invalid loop
	nesting.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/clauses-fail.c: Likewise.
	* c-c++-common/goacc/sb-1.c: Likewise.
	* c-c++-common/goacc/sb-3.c: Likewise.
	* gcc.dg/goacc/sb-1.c: Likewise.
	* gcc.dg/goacc/sb-3.c: Likewise.


diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 99b3939..2329a71 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2901,6 +2901,14 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	    }
 	  return true;
 	}
+      if (is_gimple_omp_oacc (stmt) && ctx == NULL
+	  && get_oacc_fn_attrib (current_function_decl) == NULL)
+	{
+	  error_at (gimple_location (stmt),
+		    "acc loops must be associated with an acc region or "
+		    "routine");
+	  return false;
+	}
       /* FALLTHRU */
     case GIMPLE_CALL:
       if (is_gimple_call (stmt)
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index b38e181..75d6a1d 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -20,6 +20,7 @@ f_acc_kernels (void)
   }
 }
 
+#pragma acc routine
 void
 f_acc_loop (void)
 {
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 14c6aa6..6d91484 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -361,72 +361,72 @@ f_acc_data (void)
 void
 f_acc_loop (void)
 {
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp parallel
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp for
       for (i = 0; i < 3; i++)
 	;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp sections
       {
 	;
       }
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp single
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp task
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp master
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp critical
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp ordered
       ;
     }
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 0; i < 2; ++i)
     {
-#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target
       ;
-#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target data
       ;
-#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target update to(i)
     }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index 8990180..5a572f6 100644
--- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -16,3 +16,5 @@ f (void)
   for (i = 0; i < 2; ++i)
     ;
 }
+
+/* { dg-error "acc loops must be associated with an acc region or routine" "" { target *-*-* } 15 } */
\ No newline at end of file
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
index 5e1a248..bb8b9f3 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -36,16 +36,16 @@ int test1()
       i = d;
       a[i] = 1;
     }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   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 "acc loops must be associated with an acc region or routine" } */
   for (i = 1; i < 10; i++)
     {
     }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (i = 1; i < 10; i+=2)
     {
       a[i] = i;
diff --git a/gcc/testsuite/c-c++-common/goacc/non-routine.c b/gcc/testsuite/c-c++-common/goacc/non-routine.c
new file mode 100644
index 0000000..0af69cf
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/non-routine.c
@@ -0,0 +1,16 @@
+/* This program validates the behavior of acc loops which are
+   not associated with a parallel or kernles region or routine.  */
+
+/* { dg-do compile } */
+
+int
+main ()
+{
+  int i, v = 0;
+
+#pragma acc loop gang reduction (+:v) /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  for (i = 0; i < 10; i++)
+    v++;
+
+  return v;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/sb-1.c b/gcc/testsuite/c-c++-common/goacc/sb-1.c
index 5e55c95..968ce5f 100644
--- a/gcc/testsuite/c-c++-common/goacc/sb-1.c
+++ b/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 "acc loops must be associated with an acc region or routine" } */
     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 "acc loops must be associated with an acc region or routine" } */
   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 "acc loops must be associated with an acc region or routine" } */
     for (l = 0; l < 2; ++l)
       {
 	int i;
diff --git a/gcc/testsuite/c-c++-common/goacc/sb-3.c b/gcc/testsuite/c-c++-common/goacc/sb-3.c
index 1567a10..dd57a8f 100644
--- a/gcc/testsuite/c-c++-common/goacc/sb-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/sb-3.c
@@ -3,7 +3,7 @@
 void f (void)
 {
   int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for(i = 1; i < 30; i++)
     {
       if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }
diff --git a/gcc/testsuite/gcc.dg/goacc/sb-1.c b/gcc/testsuite/gcc.dg/goacc/sb-1.c
index 0a3316c..cf61518 100644
--- a/gcc/testsuite/gcc.dg/goacc/sb-1.c
+++ b/gcc/testsuite/gcc.dg/goacc/sb-1.c
@@ -9,7 +9,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 "acc loops must be associated with an acc region or routine" } */
     for (l = 0; l < 2; ++l)
       goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
 
@@ -32,7 +32,7 @@ void foo()
     }
 
   goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for (l = 0; l < 2; ++l)
     {
       bad2_loop: ;
@@ -62,7 +62,7 @@ void foo()
 	{ ok1_data: break; }
     }
 
-  #pragma acc loop
+  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
     for (l = 0; l < 2; ++l)
       {
 	int i;
diff --git a/gcc/testsuite/gcc.dg/goacc/sb-3.c b/gcc/testsuite/gcc.dg/goacc/sb-3.c
index acbe9d5..7999ebe 100644
--- a/gcc/testsuite/gcc.dg/goacc/sb-3.c
+++ b/gcc/testsuite/gcc.dg/goacc/sb-3.c
@@ -1,7 +1,7 @@
 void f (void)
 {
   int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
   for(i = 1; i < 30; i++)
     {
       if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }