From 30b2cd7ac340764d4f7eb14730b16a49e8799e32 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 10 Dec 2014 09:52:42 +0000 Subject: [PATCH] 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. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218569 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 7 + gcc/omp-low.c | 124 +-- .../c-c++-common/goacc-gomp/nesting-fail-1.c | 898 +++++++++++---------- gcc/testsuite/c-c++-common/goacc/nesting-1.c | 49 ++ gcc/testsuite/c-c++-common/goacc/nesting-2.c | 11 - gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 22 +- .../gfortran.dg/goacc/parallel-kernels-regions.f95 | 25 +- 7 files changed, 631 insertions(+), 505 deletions(-) rewrite gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (61%) delete mode 100644 gcc/testsuite/c-c++-common/goacc/nesting-2.c diff --git a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp index 06e858363d5..970e744fa9b 100644 --- a/gcc/ChangeLog.gomp +++ b/gcc/ChangeLog.gomp @@ -1,4 +1,11 @@ 2014-12-10 Thomas Schwinge + + * 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 Bernd Schmidt * gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 39e2f225e84..d16e2de0834 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -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)) - { - /* 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 + /* 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))) { - /* 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; diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c dissimilarity index 61% index 59f41a8af02..d52c7c08a6b 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -1,427 +1,471 @@ -/* TODO: Some of these should either be allowed or fail with a more sensible - error message. */ -void -f_omp (void) -{ - int i; - -#pragma omp parallel - { -#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" } */ - ; - } - -#pragma omp for - for (i = 0; i < 3; i++) - { -#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" } */ - ; -#pragma acc loop /* { dg-error "may not be closely nested" } */ - for (i = 0; i < 2; ++i) - ; - } - -#pragma omp sections - { - { -#pragma acc parallel /* { dg-error "may not be nested" } */ - ; - } -#pragma omp section - { -#pragma acc kernels /* { dg-error "may not be nested" } */ - ; - } -#pragma omp section - { -#pragma acc data /* { dg-error "may not be nested" } */ - ; - } -#pragma omp section - { -#pragma acc loop /* { dg-error "may not be closely nested" } */ - for (i = 0; i < 2; ++i) - ; - } - } - -#pragma omp single - { -#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" } */ - ; -#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 kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; -#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 kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; -#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 kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; -#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 kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; -#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 kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; -#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 acc parallel - { - int i; -#pragma omp for /* { dg-error "may not be nested" } */ - for (i = 0; i < 3; i++) - ; - } - -#pragma acc parallel - { -#pragma omp sections /* { dg-error "may not be nested" } */ - { - ; - } - } - -#pragma acc parallel - { -#pragma omp single /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc parallel - { -#pragma omp task /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc parallel - { -#pragma omp master /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc parallel - { -#pragma omp critical /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc parallel - { - int i; -#pragma omp atomic write - i = 0; /* { dg-error "may not be nested" } */ - } - -#pragma acc parallel - { -#pragma omp ordered /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc parallel - { -#pragma omp target /* { dg-error "may not be nested" } */ - ; - } -} - -/* 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 acc kernels - { - int i; -#pragma omp for /* { dg-error "may not be nested" } */ - for (i = 0; i < 3; i++) - ; - } - -#pragma acc kernels - { -#pragma omp sections /* { dg-error "may not be nested" } */ - { - ; - } - } - -#pragma acc kernels - { -#pragma omp single /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc kernels - { -#pragma omp task /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc kernels - { -#pragma omp master /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc kernels - { -#pragma omp critical /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc kernels - { - int i; -#pragma omp atomic write - i = 0; /* { dg-error "may not be nested" } */ - } - -#pragma acc kernels - { -#pragma omp ordered /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc kernels - { -#pragma omp target /* { dg-error "may not be nested" } */ - ; - } -} - -/* 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 acc data - { - int i; -#pragma omp for /* { dg-error "may not be nested" } */ - for (i = 0; i < 3; i++) - ; - } - -#pragma acc data - { -#pragma omp sections /* { dg-error "may not be nested" } */ - { - ; - } - } - -#pragma acc data - { -#pragma omp single /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc data - { -#pragma omp task /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc data - { -#pragma omp master /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc data - { -#pragma omp critical /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc data - { - int i; -#pragma omp atomic write - i = 0; /* { dg-error "may not be nested" } */ - } - -#pragma acc data - { -#pragma omp ordered /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc data - { -#pragma omp target /* { dg-error "may not be nested" } */ - ; - } -} - -/* TODO: Some of these should either be allowed or fail with a more sensible - error message. */ -void -f_acc_loop (void) -{ - int i; - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp parallel /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp for /* { dg-error "may not be nested" } */ - for (i = 0; i < 3; i++) - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp sections /* { dg-error "may not be nested" } */ - { - ; - } - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp single /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp task /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp master /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp critical /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp atomic write - i = 0; /* { dg-error "may not be nested" } */ - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp ordered /* { dg-error "may not be nested" } */ - ; - } - -#pragma acc loop - for (i = 0; i < 2; ++i) - { -#pragma omp target /* { dg-error "may not be nested" } */ - ; - } -} +void +f_omp (void) +{ + int i; + +#pragma omp parallel + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp sections + { + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; + } +#pragma omp section + { +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; + } +#pragma omp section + { +#pragma acc data /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; + } +#pragma omp section + { +#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) + ; + } + } + +#pragma omp single + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp task + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp master + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp critical + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp ordered + { +#pragma acc parallel /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ + ; +#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 acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } + +#pragma omp target + { +#pragma acc parallel /* { dg-error "OpenACC parallel construct inside of OpenMP target region" } */ + ; +#pragma acc kernels /* { dg-error "OpenACC kernels construct inside of OpenMP target region" } */ + ; +#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) + ; + } +} + +void +f_acc_parallel (void) +{ +#pragma acc parallel + { +#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { + int i; +#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 "non-OpenACC construct inside of OpenACC region" } */ + { + ; + } + } + +#pragma acc parallel + { +#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { +#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { +#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { +#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + } + +#pragma acc parallel + { +#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc parallel + { + 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" } */ + } +} + +void +f_acc_kernels (void) +{ +#pragma acc kernels + { +#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { + int i; +#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 "non-OpenACC construct inside of OpenACC region" } */ + { + ; + } + } + +#pragma acc kernels + { +#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { +#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { +#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { +#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + } + +#pragma acc kernels + { +#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc kernels + { + 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" } */ + } +} + +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { + int i; +#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 "non-OpenACC construct inside of OpenACC region" } */ + { + ; + } + } + +#pragma acc data + { +#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { +#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { +#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { +#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + } + +#pragma acc data + { +#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc data + { + 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" } */ + } +} + +void +f_acc_loop (void) +{ + int i; + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#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 "non-OpenACC construct inside of OpenACC region" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ + { + ; + } + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#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 "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#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 "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp atomic write + 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 "non-OpenACC construct inside of OpenACC region" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++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" } */ + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c index a489d2d6394..4fbf01802ca 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -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) + ; + } } } diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-2.c b/gcc/testsuite/c-c++-common/goacc/nesting-2.c deleted file mode 100644 index 0d350c6e42a..00000000000 --- a/gcc/testsuite/c-c++-common/goacc/nesting-2.c +++ /dev/null @@ -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) - } -} diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 00dc6022603..a83380658b3 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -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" } */ } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 index 33cb9cb621f..8b8e9893995 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 @@ -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" } -- 2.11.4.GIT