gcc/

	* omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop
	nesting.

	testsuite/
	* 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/loop-1.c: Adjust errors.
	* c-c++-common/goacc/nesting-1.c: Adjust errors.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors.
	* c-c++-common/goacc-gomp/nesting-1.c: Adjust errors.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@229129 138bc75d-0d04-0410-961f-82ee72b054a4

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index d5296ae..21164b1 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2015-10-21  Nathan Sidwell  <nathan@codesourcery.com>
+
+       * omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop
+       nesting.
+
 2015-10-21  Ilya Enkovich  <enkovich.gnu@gmail.com>
 
        * doc/tm.texi: Regenerated.

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index b71609b..ad7c017 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3178,6 +3178,43 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
       /* We split taskloop into task and nested taskloop in it.  */
       if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP)
        return true;
+      if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+       {
+         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:
       if (is_gimple_call (stmt)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 3f325b9..db0f896 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,13 @@
+2015-10-21  Nathan Sidwell  <nathan@codesourcery.com>
+
+       * 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/loop-1.c: Adjust errors.
+       * c-c++-common/goacc/nesting-1.c: Adjust errors.
+       * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors.
+       * c-c++-common/goacc-gomp/nesting-1.c: Adjust errors.
+
 2015-10-21  Ilya Enkovich  <enkovich.gnu@gmail.com>
 
        * g++.dg/ext/vector22.C: Allow VEC_COND_EXPR.

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 df45bcf..1c17818 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -5,7 +5,7 @@ f_omp_parallel (void)
   {
     int i;
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
   }

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 ac614e7..0c8ea54 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -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)
       ;
   }
@@ -379,6 +379,7 @@ f_acc_data (void)
 void
 f_acc_loop (void)
 {
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -386,6 +387,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -394,6 +396,7 @@ f_acc_loop (void)
        ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -403,6 +406,7 @@ f_acc_loop (void)
       }
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -410,6 +414,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -417,6 +422,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -424,6 +430,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -431,6 +438,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -438,6 +446,7 @@ f_acc_loop (void)
       i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
@@ -445,6 +454,7 @@ f_acc_loop (void)
       ;
     }
 
+#pragma acc parallel
 #pragma acc loop
   for (i = 0; i < 2; ++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..661d364 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -12,6 +12,7 @@ f (void)
 #pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */
   ;
 
+#pragma acc parallel
 #pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */
   for (i = 0; i < 2; ++i)
     ;

diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
index fea40e0..2d8f4d5 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -38,16 +38,16 @@ int test1()
       i = d;
       a[i] = 1;
     }
-  #pragma acc loop
+  #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
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
   for (i = 1; i < 10; i++)
     {
     }
-  #pragma acc loop
+  #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;

diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c
index b4b863f..3a8f838 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c
@@ -58,7 +58,7 @@ f_acc_data (void)
 
 #pragma acc exit data delete(i)
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
     for (i = 0; i < 2; ++i)
       ;
 
@@ -93,7 +93,7 @@ f_acc_data (void)
 
 #pragma acc exit data delete(i)
 
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
       for (i = 0; i < 2; ++i)
        ;
     }

diff --git a/gcc/testsuite/c-c++-common/goacc/sb-1.c b/gcc/testsuite/c-c++-common/goacc/sb-1.c
index 5e55c95..1ce41df 100644 (file)
--- 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 "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
+  #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
+  #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC 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 147b7b0..620498e 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/sb-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/sb-3.c
@@ -3,11 +3,11 @@
 void f (void)
 {
   int i, j;
-#pragma acc loop
+#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" }
-#pragma acc loop // { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+#pragma acc loop
       for(j = 5; j < 10; j++)
        {
          if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }