diff mbox

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

Message ID 560AD984.5070601@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Sept. 29, 2015, 6:33 p.m. UTC 
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
diff mbox

Patch

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