[gomp4] error on acc loops not associated with offloaded acc regions

[gomp4] error on acc loops not associated with offloaded acc regions

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 346 bytes --]

I've applied this patch to gomp-4_0-branch which teaches omplower how to
error when it detects acc loops which aren't nested inside an acc
parallel or kernels region or located within a function marked as an acc
routine. A couple of test cases needed to be updated.

The error message is kind of long. Let me know if it should be revised.

Cesar

[-- Attachment #2: lonely-loops.diff --]
[-- Type: text/x-patch, Size: 9909 bytes --]

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" }

Re: [gomp4] error on acc loops not associated with offloaded acc regions

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 3139 bytes --]

Hi Cesar!

On Mon, 28 Sep 2015 10:08:34 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> I've applied this patch to gomp-4_0-branch which teaches omplower how to
> error when it detects acc loops which aren't nested inside an acc
> parallel or kernels region or located within a function marked as an acc
> routine. A couple of test cases needed to be updated.
> 
> The error message is kind of long. Let me know if it should be revised.

> 	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.

What about any Fortran test cases?

> --- 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)

I see that the error reporting doesn't really use a consistent style
currently, but what about something like "loop directive must be
associated with compute region" (where "compute region" is the language
used by OpenACC 2.0a to mean the structured block associated with a
compute construct as well as routine directive)?

> --- 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)
>  {

OK, but...

> --- 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
>        ;
>      }

... here you're changing what this is meant to be testing, so please
restore the original meaning (by adding "#pragma acc routine" to this
function, I suppose), and then perhaps add whichever additional test
cases you deem necessary.

> --- /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.  */

:-) Thanks for adding such a comment -- this is missing in too many test
cases.


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [gomp4] error on acc loops not associated with offloaded acc regions

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 3980 bytes --]

On 09/29/2015 02:48 AM, Thomas Schwinge wrote:

> On Mon, 28 Sep 2015 10:08:34 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
>> I've applied this patch to gomp-4_0-branch which teaches omplower how to
>> error when it detects acc loops which aren't nested inside an acc
>> parallel or kernels region or located within a function marked as an acc
>> routine. A couple of test cases needed to be updated.
>>
>> The error message is kind of long. Let me know if it should be revised.
> 
>> 	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.
> 
> What about any Fortran test cases?

My first thought was that we didn't need one because this is generic
error handling in omplow, and there are already a lot of c tests cases
exercising it. However a fortran test can't hurt, so I added one in this
new patch. Note that I had to create a new test instead of hijacking an
existing test, because the fortran front end bails out when it detects
errors before it hands anything over to omplow. And the existing tests
had a bunch of expected front end errors.

>> --- 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)
> 
> I see that the error reporting doesn't really use a consistent style
> currently, but what about something like "loop directive must be
> associated with compute region" (where "compute region" is the language
> used by OpenACC 2.0a to mean the structured block associated with a
> compute construct as well as routine directive)?

That sounds reasonable, but it's not much shorter.

>> --- 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)
>>  {
> 
> OK, but...
> 
>> --- 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
>>        ;
>>      }
> 
> ... here you're changing what this is meant to be testing, so please
> restore the original meaning (by adding "#pragma acc routine" to this
> function, I suppose), and then perhaps add whichever additional test
> cases you deem necessary.

I was wondering about that too. After thinking about it some more, I did
as you suggested -- revert those changes and used a routine pragma.

>> --- /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.  */
> 
> :-) Thanks for adding such a comment -- this is missing in too many test
> cases.

We definitely need more of them. I'm not starting to forget what I was
trying to test several months ago.

I'll apply this patch to gomp4.

Cesar


[-- Attachment #2: lonely-loops-update.diff --]
[-- Type: text/x-patch, Size: 10935 bytes --]

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

	gcc/
	* omp-low.c (check_omp_nesting_restrictions): Update the error
	message for loops not affliated with acc compute regions.

	gcc/testsuite/
	* c-c++-common/goacc-gomp/nesting-fail-1.c (f_omp): Revert changes and
	mark the function as an acc routine.
	* c-c++-common/goacc/clauses-fail.c: Likewise.
	* c-c++-common/goacc/loop-1.c: Likewise.
	* c-c++-common/goacc/non-routine.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.
	* gfortran.dg/goacc/loop-4.f95: New test.


diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ba8cdf4..dff013d 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2923,8 +2923,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	  && get_oacc_fn_attrib (current_function_decl) == NULL)
 	{
 	  error_at (gimple_location (stmt),
-		    "acc loops must be associated with an acc region or "
-		    "routine");
+		    "loop directive must be associated with a compute "
+		    "region");
 	  return false;
 	}
       /* FALLTHRU */
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 6d91484..fa78a66 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
@@ -358,75 +358,76 @@ f_acc_data (void)
   }
 }
 
+#pragma acc routine
 void
 f_acc_loop (void)
 {
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp parallel
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp for
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       for (i = 0; i < 3; i++)
 	;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp sections
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       {
 	;
       }
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp single
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp task
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp master
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp critical
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp ordered
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
 
-#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp target
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
-#pragma omp target data
+#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
-#pragma omp target update to(i)
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index 5a572f6..529ee54 100644
--- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -17,4 +17,4 @@ f (void)
     ;
 }
 
-/* { dg-error "acc loops must be associated with an acc region or routine" "" { target *-*-* } 15 } */
\ No newline at end of file
+/* { dg-error "loop directive must be associated with a compute region" "" { target *-*-* } 15 } */
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
index bb8b9f3..501b350 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   for (i = 1; i < 10; i++)
     {
     }
-  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   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
index 0af69cf..9caafd3 100644
--- a/gcc/testsuite/c-c++-common/goacc/non-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/non-routine.c
@@ -8,7 +8,7 @@ main ()
 {
   int i, v = 0;
 
-#pragma acc loop gang reduction (+:v) /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with a compute region" } */
   for (i = 0; i < 10; i++)
     v++;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/sb-1.c b/gcc/testsuite/c-c++-common/goacc/sb-1.c
index 968ce5f..61c44e1 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   for (l = 0; l < 2; ++l)
     {
       bad2_loop: ;
@@ -64,7 +64,7 @@ void foo()
 	{ ok1_data: break; }
     }
 
-  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
     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 dd57a8f..f48d52e 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   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 cf61518..23bcd6b 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
     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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   for (l = 0; l < 2; ++l)
     {
       bad2_loop: ;
@@ -62,7 +62,7 @@ void foo()
 	{ ok1_data: break; }
     }
 
-  #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
     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 7999ebe..acce18a 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 /* { dg-error "acc loops must be associated with an acc region or routine" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
   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/gfortran.dg/goacc/loop-4.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
new file mode 100644
index 0000000..76ffe74
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
@@ -0,0 +1,7 @@
+! Ensure that loops not affiliated with acc compute regions cause an error.
+
+subroutine test1
+    !$acc loop gang ! { dg-error "loop directive must be associated with a compute region" }
+  DO i = 1,10
+  ENDDO
+end subroutine test1