diff mbox

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

Message ID 56097412.4080600@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Sept. 28, 2015, 5:08 p.m. UTC 
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

Comments

Thomas Schwinge Sept. 29, 2015, 9:48 a.m. UTC | #1 
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
diff mbox

Patch

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