Hi Jakub! On Wed, 10 Dec 2014 11:02:24 +0100, I wrote: > OpenACC: Rework nested constructs checking. > > gcc/ > * 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.
Regarding the check_omp_nesting_restrictions rework:
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -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;
For guarding against OpenMP target regressions, I used the following
tests -- OK to commit to trunk?
gcc/testsuite/c-c++-common/gomp/nesting-1.c | 77 ++++++++++++++++++++++++
gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c | 23 +++++++
2 files changed, 100 insertions(+)
diff --git gcc/testsuite/c-c++-common/gomp/nesting-1.c
gcc/testsuite/c-c++-common/gomp/nesting-1.c
new file mode 100644
index 0000000..70b3e8d
--- /dev/null
+++ gcc/testsuite/c-c++-common/gomp/nesting-1.c
@@ -0,0 +1,77 @@
+void
+f_omp_parallel (void)
+{
+#pragma omp parallel
+ {
+ int i;
+
+#pragma omp parallel
+ ;
+
+#pragma omp target
+ ;
+
+#pragma omp target data
+ ;
+
+#pragma omp target update to(i)
+
+#pragma omp target data
+ {
+#pragma omp parallel
+ ;
+
+#pragma omp target
+ ;
+
+#pragma omp target data
+ ;
+
+#pragma omp target update to(i)
+ }
+ }
+}
+
+void
+f_omp_target (void)
+{
+#pragma omp target
+ {
+#pragma omp parallel
+ ;
+ }
+}
+
+void
+f_omp_target_data (void)
+{
+#pragma omp target data
+ {
+ int i;
+
+#pragma omp parallel
+ ;
+
+#pragma omp target
+ ;
+
+#pragma omp target data
+ ;
+
+#pragma omp target update to(i)
+
+#pragma omp target data
+ {
+#pragma omp parallel
+ ;
+
+#pragma omp target
+ ;
+
+#pragma omp target data
+ ;
+
+#pragma omp target update to(i)
+ }
+ }
+}
diff --git gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
new file mode 100644
index 0000000..6a65699
--- /dev/null
+++ gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
@@ -0,0 +1,23 @@
+void
+f_omp_target (void)
+{
+#pragma omp target
+ {
+ int i;
+
+#pragma omp target /* { dg-warning "target construct inside of target region"
} */
+ ;
+#pragma omp target data /* { dg-warning "target data construct inside of
target region" } */
+ ;
+#pragma omp target update to(i) /* { dg-warning "target update construct
inside of target region" } */
+
+#pragma omp parallel
+ {
+#pragma omp target /* { dg-warning "target construct inside of target region"
} */
+ ;
+#pragma omp target data /* { dg-warning "target data construct inside of
target region" } */
+ ;
+#pragma omp target update to(i) /* { dg-warning "target update construct
inside of target region" } */
+ }
+ }
+}
Grüße,
Thomas
signature.asc
Description: PGP signature
