[gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 1aebc4d..f43452c 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@ 
+2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* omp-low.c (check_omp_nesting_restrictions): Allow nesting of
+	OpenACC constructs inside of OpenACC data constructs.
+
 2014-03-18  Ilmir Usmanov  <i.usmanov@samsung.com>
 
 	* tree.def (OACC_LOOP): New tree code.
diff --git gcc/omp-low.c gcc/omp-low.c
index f1b0fa5..23a0dda 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2416,26 +2416,31 @@  scan_omp_teams (gimple stmt, omp_context *outer_ctx)
 static bool
 check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 {
-  omp_context *ctx_;
-
   /* TODO: While the OpenACC specification does allow for certain kinds of
-     nesting, we don't support that yet.  */
-  /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin)
-     inside any OpenACC CTX.  */
-  for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
-    if (is_gimple_omp (ctx_->stmt)
-	&& is_gimple_omp_oacc_specifically (ctx_->stmt))
-      {
-	error_at (gimple_location (stmt),
-		  "may not be nested");
-	return false;
-      }
-  /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX.  */
+     nesting, we don't support many of these yet.  */
   if (is_gimple_omp (stmt)
       && is_gimple_omp_oacc_specifically (stmt))
     {
-      for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
-	if (is_gimple_omp (ctx_->stmt))
+      /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different
+	 from an OpenACC data construct.  */
+      for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+	if (is_gimple_omp (ctx_->stmt)
+	    && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET
+		 && (gimple_omp_target_kind (ctx_->stmt)
+		     == GF_OMP_TARGET_KIND_OACC_DATA)))
+	  {
+	    error_at (gimple_location (stmt),
+		      "may not be nested");
+	    return false;
+	  }
+    }
+  else
+    {
+      /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP
+	 builtin) inside any OpenACC CTX.  */
+      for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+	if (is_gimple_omp (ctx_->stmt)
+	    && is_gimple_omp_oacc_specifically (ctx_->stmt))
 	  {
 	    error_at (gimple_location (stmt),
 		      "may not be nested");
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index fd38d80..13e99d5 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,9 @@ 
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* c-c++-common/goacc/nesting-1.c: New file.
+	* c-c++-common/goacc/nesting-data-1.c: Likewise.
+	* c-c++-common/goacc/nesting-fail-1.c: Update.
+
 	* c-c++-common/goacc/nesting-fail-1.c (f_acc_kernels): Replace
 	OpenACC parallel with kernels directive.
 
diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c
new file mode 100644
index 0000000..3a22292
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-1.c
@@ -0,0 +1,13 @@ 
+void
+f_acc_data (void)
+{
+#pragma acc data
+  {
+#pragma acc parallel
+    ;
+#pragma acc kernels
+    ;
+#pragma acc data
+    ;
+  }
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-data-1.c gcc/testsuite/c-c++-common/goacc/nesting-data-1.c
new file mode 100644
index 0000000..fefe6cd
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-data-1.c
@@ -0,0 +1,61 @@ 
+void
+f (void)
+{
+  unsigned char c, ca[15], caa[20][30];
+
+#pragma acc data copyin(c)
+  {
+    c = 5;
+    ca[3] = c;
+    caa[3][12] = ca[3] + caa[3][12];
+
+#pragma acc data copyin(ca[2:4])
+    {
+      c = 6;
+      ca[4] = c;
+      caa[3][12] = ca[3] + caa[3][12];
+    }
+
+#pragma acc parallel copyout(ca[3:4])
+    {
+      c = 7;
+      ca[5] = c;
+      caa[3][12] = ca[3] + caa[3][12];
+    }
+
+#pragma acc kernels copy(ca[4:4])
+    {
+      c = 8;
+      ca[6] = c;
+      caa[3][12] = ca[3] + caa[3][12];
+    }
+
+#pragma acc data pcopy(ca[5:7])
+    {
+      c = 15;
+      ca[7] = c;
+      caa[3][12] = ca[3] + caa[3][12];
+
+#pragma acc data pcopyin(caa[3:7][0:30])
+      {
+	c = 16;
+	ca[8] = c;
+	caa[3][12] = ca[3] + caa[3][12];
+      }
+
+#pragma acc parallel pcopyout(caa[3:7][0:30])
+      {
+	c = 17;
+	ca[9] = c;
+	caa[3][12] = ca[3] + caa[3][12];
+      }
+
+#pragma acc kernels pcopy(caa[3:7][0:30])
+      {
+	c = 18;
+	ca[10] = c;
+	caa[3][12] = ca[3] + caa[3][12];
+      }
+    }
+  }
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index ca8921f..00dc602 100644
--- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -1,5 +1,5 @@ 
 /* TODO: While the OpenACC specification does allow for certain kinds of
-   nesting, we don't support that yet.  */
+   nesting, we don't support many of these yet.  */
 void
 f_acc_parallel (void)
 {
@@ -15,7 +15,7 @@  f_acc_parallel (void)
 }
 
 /* TODO: While the OpenACC specification does allow for certain kinds of
-   nesting, we don't support that yet.  */
+   nesting, we don't support many of these yet.  */
 void
 f_acc_kernels (void)
 {
@@ -29,19 +29,3 @@  f_acc_kernels (void)
     ;
   }
 }
-
-/* TODO: While the OpenACC specification does allow for certain kinds of
-   nesting, we don't support that yet.  */
-void
-f_acc_data (void)
-{
-#pragma acc data
-  {
-#pragma acc parallel	/* { dg-error "may not be nested" } */
-    ;
-#pragma acc kernels	/* { dg-error "may not be nested" } */
-    ;
-#pragma acc data	/* { dg-error "may not be nested" } */
-    ;
-  }
-}