Hi! Applied in r208701 to gomp-4_0-branch:
commit 22dd36a31c433dcd8bcc890d245a9e4ac6ed9c7f
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Mar 20 14:33:28 2014 +0000
Nesting of OpenACC constructs inside of OpenACC data constructs.
gcc/
* omp-low.c (check_omp_nesting_restrictions): Allow nesting of
OpenACC constructs inside of OpenACC data constructs.
gcc/testsuite/
* 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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208701
138bc75d-0d04-0410-961f-82ee72b054a4
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 <tho...@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.usma...@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 <tho...@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" } */
- ;
- }
-}
Grüße,
Thomas
pgp9sHUE0nQyV.pgp
Description: PGP signature
