diff mbox

[gomp4] loop nesting check

Message ID 5627C4C8.1060308@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Oct. 21, 2015, 5 p.m. UTC 
This is the gomp4-branch variant of the loop nesting patch I just committed to 
trunk. The gomp4 branch had some checking, but
a) it didn't catch all erroreous cases
b) gave an ambiguous error, by not mentioning 'OpenACC'

committed to gomp4

nathan
diff mbox

Patch

2015-10-21  Nathan Sidwell  <nathan@codesourcery.com>

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

	gcc/testsuite/g
	* c-c++-common/goacc/non-routine.c: Adjust errors.
	* c-c++-common/goacc/loop-1.c: Adjust errors.
	* 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-gomp/nesting-fail-1.c: Adjust errors.
	* gfortran.dg/goacc/loop-4.f95 : Adjust errors.
	* gcc.dg/goacc/sb-3.c: Adjust errors.
	* gcc.dg/goacc/sb-1.c: Adjust errors.

Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 229094)
+++ gcc/omp-low.c	(working copy)
@@ -3106,13 +3106,42 @@  check_omp_nesting_restrictions (gimple *
       /* We split taskloop into task and nested taskloop in it.  */
       if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP)
 	return true;
-      if (is_gimple_omp_oacc (stmt) && ctx == NULL
-	  && get_oacc_fn_attrib (current_function_decl) == NULL)
+      if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
 	{
-	  error_at (gimple_location (stmt),
-		    "loop directive must be associated with a compute "
-		    "region");
-	  return false;
+	  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:
Index: gcc/testsuite/c-c++-common/goacc/non-routine.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/non-routine.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc/non-routine.c	(working copy)
@@ -8,7 +8,7 @@  main ()
 {
   int i, v = 0;
 
-#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with a compute region" } */
+#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 0; i < 10; i++)
     v++;
 
Index: gcc/testsuite/c-c++-common/goacc/loop-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/loop-1.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc/loop-1.c	(working copy)
@@ -36,16 +36,16 @@  int test1()
       i = d;
       a[i] = 1;
     }
-  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+  #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 /* { dg-error "loop directive must be associated with a compute region" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 1; i < 10; i++)
     {
     }
-  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+  #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;
Index: gcc/testsuite/c-c++-common/goacc/clauses-fail.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/clauses-fail.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c	(working copy)
@@ -17,4 +17,4 @@  f (void)
     ;
 }
 
-/* { dg-error "loop directive must be associated with a compute region" "" { target *-*-* } 15 } */
+/* { dg-error "loop directive must be associated with an OpenACC compute region" "" { target *-*-* } 15 } */
Index: gcc/testsuite/c-c++-common/goacc/sb-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/sb-1.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc/sb-1.c	(working copy)
@@ -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 "loop directive must be associated with a compute region" } */
+  #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 /* { dg-error "loop directive must be associated with a compute region" } */
+  #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 /* { dg-error "loop directive must be associated with a compute region" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (l = 0; l < 2; ++l)
       {
 	int i;
Index: gcc/testsuite/c-c++-common/goacc/sb-3.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/sb-3.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc/sb-3.c	(working copy)
@@ -3,7 +3,7 @@ 
 void f (void)
 {
   int i, j;
-#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+#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" }
Index: gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c	(revision 229094)
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c	(working copy)
@@ -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)
       ;
   }
Index: gcc/testsuite/gfortran.dg/goacc/loop-4.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-4.f95	(revision 229094)
+++ gcc/testsuite/gfortran.dg/goacc/loop-4.f95	(working copy)
@@ -1,7 +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" }
+    !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" }
   DO i = 1,10
   ENDDO
 end subroutine test1
Index: gcc/testsuite/gcc.dg/goacc/sb-3.c
===================================================================
--- gcc/testsuite/gcc.dg/goacc/sb-3.c	(revision 229094)
+++ gcc/testsuite/gcc.dg/goacc/sb-3.c	(working copy)
@@ -1,7 +1,7 @@ 
 void f (void)
 {
   int i, j;
-#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+#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" }
Index: gcc/testsuite/gcc.dg/goacc/sb-1.c
===================================================================
--- gcc/testsuite/gcc.dg/goacc/sb-1.c	(revision 229094)
+++ gcc/testsuite/gcc.dg/goacc/sb-1.c	(working copy)
@@ -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 "loop directive must be associated with a compute region" } */
+  #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" }
 
@@ -32,7 +32,7 @@  void foo()
     }
 
   goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
-  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (l = 0; l < 2; ++l)
     {
       bad2_loop: ;
@@ -62,7 +62,7 @@  void foo()
 	{ ok1_data: break; }
     }
 
-  #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (l = 0; l < 2; ++l)
       {
 	int i;