@@ -1,4 +1,11 @@
2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (scan_omp_target): Remove taskreg_nesting_level and
+ target_nesting_level assertions.
+ (check_omp_nesting_restrictions): Rework OpenACC constructs
+ handling. Update and extend the relevant test cases.
+
+2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
Bernd Schmidt <bernds@codesourcery.com>
* gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge
@@ -2647,12 +2647,6 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
tree name;
bool offloaded = is_gimple_omp_offloaded (stmt);
- if (is_gimple_omp_oacc_specifically (stmt))
- {
- gcc_assert (taskreg_nesting_level == 0);
- gcc_assert (target_nesting_level == 0);
- }
-
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
@@ -2706,46 +2700,26 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx)
scan_omp (gimple_omp_body_ptr (stmt), ctx);
}
-/* Check OpenMP nesting restrictions. */
+/* Check nesting restrictions. */
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
- /* TODO: While the OpenACC specification does allow for certain kinds of
- nesting, we don't support many of these yet. */
- if (is_gimple_omp (stmt)
- && is_gimple_omp_oacc_specifically (stmt))
+ /* TODO: Some OpenACC/OpenMP nesting should be allowed. */
+
+ /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
+ inside an OpenACC CTX. */
+ if (!(is_gimple_omp (stmt)
+ && is_gimple_omp_oacc_specifically (stmt)))
{
- /* Regular handling of OpenACC loop constructs. */
- if (gimple_code (stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
- goto cont;
- /* 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");
+ "non-OpenACC construct inside of OpenACC region");
return false;
}
}
- cont:
if (ctx != NULL)
{
@@ -3003,20 +2977,74 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
break;
case GIMPLE_OMP_TARGET:
for (; ctx != NULL; ctx = ctx->outer)
- if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
- {
- const char *name;
- switch (gimple_omp_target_kind (stmt))
- {
- case GF_OMP_TARGET_KIND_REGION: name = "target"; break;
- case GF_OMP_TARGET_KIND_DATA: name = "target data"; break;
- case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break;
- default: gcc_unreachable ();
- }
- warning_at (gimple_location (stmt), 0,
- "%s construct inside of target region", name);
- }
+ {
+ if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
+ {
+ if (is_gimple_omp (stmt)
+ && is_gimple_omp_oacc_specifically (stmt)
+ && is_gimple_omp (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "OpenACC construct inside of non-OpenACC region");
+ return false;
+ }
+ continue;
+ }
+
+ const char *stmt_name, *ctx_stmt_name;
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION: stmt_name = "target"; break;
+ case GF_OMP_TARGET_KIND_DATA: stmt_name = "target data"; break;
+ case GF_OMP_TARGET_KIND_UPDATE: stmt_name = "target update"; break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break;
+ case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: stmt_name = "enter/exit data"; break;
+ default: gcc_unreachable ();
+ }
+ switch (gimple_omp_target_kind (ctx->stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION: ctx_stmt_name = "target"; break;
+ case GF_OMP_TARGET_KIND_DATA: ctx_stmt_name = "target data"; break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL: ctx_stmt_name = "parallel"; break;
+ case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
+ default: gcc_unreachable ();
+ }
+
+ /* OpenACC/OpenMP mismatch? */
+ if (is_gimple_omp_oacc_specifically (stmt)
+ != is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "%s %s construct inside of %s %s region",
+ (is_gimple_omp_oacc_specifically (stmt)
+ ? "OpenACC" : "OpenMP"), stmt_name,
+ (is_gimple_omp_oacc_specifically (ctx->stmt)
+ ? "OpenACC" : "OpenMP"), ctx_stmt_name);
+ return false;
+ }
+ if (is_gimple_omp_offloaded (ctx->stmt))
+ {
+ /* No GIMPLE_OMP_TARGET inside offloaded OpenACC CTX. */
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "%s construct inside of %s region",
+ stmt_name, ctx_stmt_name);
+ return false;
+ }
+ else
+ {
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
+ warning_at (gimple_location (stmt), 0,
+ "%s construct inside of %s region",
+ stmt_name, ctx_stmt_name);
+ }
+ }
+ }
break;
default:
break;
@@ -1,5 +1,3 @@
-/* TODO: Some of these should either be allowed or fail with a more sensible
- error message. */
void
f_omp (void)
{
@@ -7,24 +5,30 @@ f_omp (void)
#pragma omp parallel
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
+#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 omp for
for (i = 0; i < 3; i++)
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
@@ -32,22 +36,34 @@ f_omp (void)
#pragma omp sections
{
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
}
#pragma omp section
{
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
}
#pragma omp section
{
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
}
#pragma omp section
{
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
+ }
+#pragma omp section
+ {
+#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
+ }
+#pragma omp section
+ {
+#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
+ }
+#pragma omp section
+ {
+#pragma acc loop /* { dg-error "may not be closely nested" } */
for (i = 0; i < 2; ++i)
;
}
@@ -55,105 +71,121 @@ f_omp (void)
#pragma omp single
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
#pragma omp task
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
#pragma omp master
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
#pragma omp critical
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
#pragma omp ordered
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
;
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#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" } */
for (i = 0; i < 2; ++i)
;
}
#pragma omp target
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+#pragma acc parallel /* { dg-error "OpenACC parallel construct inside of OpenMP target region" } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-error "OpenACC kernels construct inside of OpenMP target region" } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "OpenACC data construct inside of OpenMP target region" } */
;
+#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
for (i = 0; i < 2; ++i)
;
}
}
-/* TODO: Some of these should either be allowed or fail with a more sensible
- error message. */
void
f_acc_parallel (void)
{
#pragma acc parallel
{
-#pragma omp parallel /* { dg-error "may not be nested" } */
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc parallel
{
int i;
-#pragma omp for /* { dg-error "may not be nested" } */
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
for (i = 0; i < 3; i++)
;
}
#pragma acc parallel
{
-#pragma omp sections /* { dg-error "may not be nested" } */
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
{
;
}
@@ -161,25 +193,25 @@ f_acc_parallel (void)
#pragma acc parallel
{
-#pragma omp single /* { dg-error "may not be nested" } */
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc parallel
{
-#pragma omp task /* { dg-error "may not be nested" } */
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc parallel
{
-#pragma omp master /* { dg-error "may not be nested" } */
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc parallel
{
-#pragma omp critical /* { dg-error "may not be nested" } */
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
@@ -187,44 +219,47 @@ f_acc_parallel (void)
{
int i;
#pragma omp atomic write
- i = 0; /* { dg-error "may not be nested" } */
+ i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
#pragma acc parallel
{
-#pragma omp ordered /* { dg-error "may not be nested" } */
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc parallel
{
-#pragma omp target /* { dg-error "may not be nested" } */
+ int i;
+
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+ ;
+#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
}
-/* TODO: Some of these should either be allowed or fail with a more sensible
- error message. */
void
f_acc_kernels (void)
{
#pragma acc kernels
{
-#pragma omp parallel /* { dg-error "may not be nested" } */
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc kernels
{
int i;
-#pragma omp for /* { dg-error "may not be nested" } */
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
for (i = 0; i < 3; i++)
;
}
#pragma acc kernels
{
-#pragma omp sections /* { dg-error "may not be nested" } */
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
{
;
}
@@ -232,25 +267,25 @@ f_acc_kernels (void)
#pragma acc kernels
{
-#pragma omp single /* { dg-error "may not be nested" } */
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc kernels
{
-#pragma omp task /* { dg-error "may not be nested" } */
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc kernels
{
-#pragma omp master /* { dg-error "may not be nested" } */
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc kernels
{
-#pragma omp critical /* { dg-error "may not be nested" } */
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
@@ -258,44 +293,47 @@ f_acc_kernels (void)
{
int i;
#pragma omp atomic write
- i = 0; /* { dg-error "may not be nested" } */
+ i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
#pragma acc kernels
{
-#pragma omp ordered /* { dg-error "may not be nested" } */
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc kernels
{
-#pragma omp target /* { dg-error "may not be nested" } */
+ int i;
+
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+ ;
+#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
}
-/* TODO: Some of these should either be allowed or fail with a more sensible
- error message. */
void
f_acc_data (void)
{
#pragma acc data
{
-#pragma omp parallel /* { dg-error "may not be nested" } */
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc data
{
int i;
-#pragma omp for /* { dg-error "may not be nested" } */
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
for (i = 0; i < 3; i++)
;
}
#pragma acc data
{
-#pragma omp sections /* { dg-error "may not be nested" } */
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
{
;
}
@@ -303,25 +341,25 @@ f_acc_data (void)
#pragma acc data
{
-#pragma omp single /* { dg-error "may not be nested" } */
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc data
{
-#pragma omp task /* { dg-error "may not be nested" } */
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc data
{
-#pragma omp master /* { dg-error "may not be nested" } */
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc data
{
-#pragma omp critical /* { dg-error "may not be nested" } */
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
@@ -329,24 +367,27 @@ f_acc_data (void)
{
int i;
#pragma omp atomic write
- i = 0; /* { dg-error "may not be nested" } */
+ i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
#pragma acc data
{
-#pragma omp ordered /* { dg-error "may not be nested" } */
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc data
{
-#pragma omp target /* { dg-error "may not be nested" } */
+ int i;
+
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+ ;
+#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
}
-/* TODO: Some of these should either be allowed or fail with a more sensible
- error message. */
void
f_acc_loop (void)
{
@@ -355,14 +396,14 @@ f_acc_loop (void)
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp parallel /* { dg-error "may not be nested" } */
+#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp for /* { dg-error "may not be nested" } */
+#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
for (i = 0; i < 3; i++)
;
}
@@ -370,7 +411,7 @@ f_acc_loop (void)
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp sections /* { dg-error "may not be nested" } */
+#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
{
;
}
@@ -379,28 +420,28 @@ f_acc_loop (void)
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp single /* { dg-error "may not be nested" } */
+#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp task /* { dg-error "may not be nested" } */
+#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp master /* { dg-error "may not be nested" } */
+#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp critical /* { dg-error "may not be nested" } */
+#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
@@ -408,20 +449,23 @@ f_acc_loop (void)
for (i = 0; i < 2; ++i)
{
#pragma omp atomic write
- i = 0; /* { dg-error "may not be nested" } */
+ i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp ordered /* { dg-error "may not be nested" } */
+#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
}
#pragma acc loop
for (i = 0; i < 2; ++i)
{
-#pragma omp target /* { dg-error "may not be nested" } */
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
;
+#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+ ;
+#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
}
@@ -46,11 +46,60 @@ f_acc_data (void)
#pragma acc kernels
;
+#pragma acc kernels
+ {
+#pragma acc loop
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+
#pragma acc data
;
+#pragma acc update host(i)
+
+#pragma acc enter data copyin(i)
+
+#pragma acc exit data delete(i)
+
#pragma acc loop
for (i = 0; i < 2; ++i)
;
+
+#pragma acc data
+ {
+#pragma acc parallel
+ ;
+
+#pragma acc parallel
+ {
+#pragma acc loop
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+
+#pragma acc kernels
+ ;
+
+#pragma acc kernels
+ {
+#pragma acc loop
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+
+#pragma acc data
+ ;
+
+#pragma acc update host(i)
+
+#pragma acc enter data copyin(i)
+
+#pragma acc exit data delete(i)
+
+#pragma acc loop
+ for (i = 0; i < 2; ++i)
+ ;
+ }
}
}
deleted file mode 100644
@@ -1,11 +0,0 @@
-int i;
-
-void
-f_acc_data (void)
-{
-#pragma acc data
- {
-#pragma acc update host(i)
-#pragma acc enter data copyin(i)
- }
-}
@@ -5,12 +5,17 @@ f_acc_parallel (void)
{
#pragma acc parallel
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+ int i;
+
+#pragma acc parallel /* { dg-bogus "parallel construct inside of parallel region" "not implemented" { xfail *-*-* } } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-bogus "kernels construct inside of parallel region" "not implemented" { xfail *-*-* } } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "data construct inside of parallel region" } */
;
+#pragma acc update host(i) /* { dg-error "update construct inside of parallel region" } */
+#pragma acc enter data copyin(i) /* { dg-error "enter/exit data construct inside of parallel region" } */
+#pragma acc exit data delete(i) /* { dg-error "enter/exit data construct inside of parallel region" } */
}
}
@@ -21,11 +26,16 @@ f_acc_kernels (void)
{
#pragma acc kernels
{
-#pragma acc parallel /* { dg-error "may not be nested" } */
+ int i;
+
+#pragma acc parallel /* { dg-bogus "parallel construct inside of kernels region" "not implemented" { xfail *-*-* } } */
;
-#pragma acc kernels /* { dg-error "may not be nested" } */
+#pragma acc kernels /* { dg-bogus "kernels construct inside of kernels region" "not implemented" { xfail *-*-* } } */
;
-#pragma acc data /* { dg-error "may not be nested" } */
+#pragma acc data /* { dg-error "data construct inside of kernels region" } */
;
+#pragma acc update host(i) /* { dg-error "update construct inside of kernels region" } */
+#pragma acc enter data copyin(i) /* { dg-error "enter/exit data construct inside of kernels region" } */
+#pragma acc exit data delete(i) /* { dg-error "enter/exit data construct inside of kernels region" } */
}
}
@@ -1,7 +1,7 @@
! { dg-do compile }
-! OpenACC 2.0 allows nested parallel/kernels regions
-! However, in middle-end there is check for nested parallel
+! OpenACC 2.0 allows nested parallel/kernels regions, but this is not yet
+! supported.
program test
implicit none
@@ -9,48 +9,47 @@ program test
integer :: i
!$acc parallel
- !$acc kernels
+ !$acc kernels ! { dg-bogus "kernels construct inside of parallel region" "not implemented" { xfail *-*-* } }
!$acc end kernels
!$acc end parallel
!$acc parallel
- !$acc parallel ! { dg-error "may not be nested" }
+ !$acc parallel ! { dg-bogus "parallel construct inside of parallel region" "not implemented" { xfail *-*-* } }
!$acc end parallel
!$acc end parallel
!$acc parallel
- !$acc parallel ! { dg-error "may not be nested" }
+ !$acc parallel ! { dg-bogus "parallel construct inside of parallel region" "not implemented" { xfail *-*-* } }
!$acc end parallel
- !$acc kernels
+ !$acc kernels ! { dg-bogus "kernels construct inside of parallel region" "not implemented" { xfail *-*-* } }
!$acc end kernels
!$acc end parallel
!$acc kernels
- !$acc kernels
+ !$acc kernels ! { dg-bogus "kernels construct inside of kernels region" "not implemented" { xfail *-*-* } }
!$acc end kernels
!$acc end kernels
!$acc kernels
- !$acc parallel
+ !$acc parallel ! { dg-bogus "parallel construct inside of kernels region" "not implemented" { xfail *-*-* } }
!$acc end parallel
!$acc end kernels
!$acc kernels
- !$acc parallel
+ !$acc parallel ! { dg-bogus "parallel construct inside of kernels region" "not implemented" { xfail *-*-* } }
!$acc end parallel
- !$acc kernels
+ !$acc kernels ! { dg-bogus "kernels construct inside of kernels region" "not implemented" { xfail *-*-* } }
!$acc end kernels
!$acc end kernels
!$acc parallel
- !$acc data ! { dg-error "may not be nested" }
+ !$acc data ! { dg-error "data construct inside of parallel region" }
!$acc end data
!$acc end parallel
!$acc kernels
- !$acc data
+ !$acc data ! { dg-error "data construct inside of kernels region" }
!$acc end data
!$acc end kernels
end program test
-! { dg-prune-output "Error: may not be nested" }