@@ -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.
@@ -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");
@@ -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.
new file mode 100644
@@ -0,0 +1,13 @@
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma acc parallel
+ ;
+#pragma acc kernels
+ ;
+#pragma acc data
+ ;
+ }
+}
new file mode 100644
@@ -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];
+ }
+ }
+ }
+}
@@ -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" } */
- ;
- }
-}