From patchwork Wed Dec 10 10:02:24 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 419506 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 462E41400B7 for ; Wed, 10 Dec 2014 21:03:01 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=L7BbMwIYHvyrKyzXCH38+Jo3m28Gs rlOast4jnxmAR0Xf808INR/T4QQ+5q0pcBR1MclzK93DG+wNIMi9xf8S4kEgcYfR bw757dlEh7b0zZ+pHY4r0XPDY21/h+dCJ47b5W5GtFqTsOlYlbx63Nq8+pcqRU7A VaS0caRlZZouTo= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=x/iB7bGX/BFW4fEz2F9o9yT97FQ=; b=yZi QPP57vKTeMNuVZXhB3CDiVqELvXlndJKzKCOrYSi+50h4NF/sYvSuTrmjAKifYIZ 7g6hlOVePuh22a9FBxJPf+p66cdMILcH3mjo2XTH7YYQrx0sqPLdDffDiO2z1q4g YHWag+vGWKyNVHIUDQULOu2A+GY9n5QYJZSgNbpM= Received: (qmail 14588 invoked by alias); 10 Dec 2014 10:02:52 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 14575 invoked by uid 89); 10 Dec 2014 10:02:51 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 10 Dec 2014 10:02:48 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Xye6V-0004Br-PE from Thomas_Schwinge@mentor.com ; Wed, 10 Dec 2014 02:02:44 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Wed, 10 Dec 2014 10:02:40 +0000 From: Thomas Schwinge To: Jakub Jelinek , Subject: Nested OpenACC/OpenMP constructs (was: OpenACC GIMPLE_OACC_* -- or not?) In-Reply-To: <20141112134502.GC5026@tucnak.redhat.com> References: <1383766943-8863-2-git-send-email-thomas@codesourcery.com> <1383766943-8863-3-git-send-email-thomas@codesourcery.com> <1383766943-8863-4-git-send-email-thomas@codesourcery.com> <1383766943-8863-5-git-send-email-thomas@codesourcery.com> <1383766943-8863-6-git-send-email-thomas@codesourcery.com> <1383766943-8863-7-git-send-email-thomas@codesourcery.com> <1383766943-8863-8-git-send-email-thomas@codesourcery.com> <1383766943-8863-9-git-send-email-thomas@codesourcery.com> <8761s5joir.fsf@schwinge.name> <87sihoczm0.fsf@kepler.schwinge.homeip.net> <20141112134502.GC5026@tucnak.redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Wed, 10 Dec 2014 11:02:24 +0100 Message-ID: <87sign7rgv.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Wed, 12 Nov 2014 14:45:02 +0100, Jakub Jelinek wrote: > please make sure we'll have > some tests on mixing OpenMP and OpenACC directives in the same functions > (it is fine if we error out on combinations that don't make sense or are > too hard to support). > E.g. supporting OpenACC #pragma omp target counterpart inside > of #pragma omp parallel or #pragma omp task should be presumably fine, > supporting OpenACC inside of #pragma omp target should be IMHO just > diagnosed, mixing target data and openacc is generically hard to diagnose, > perhaps at runtime, supporting #pragma omp directives inside of OpenACC > regions not needed (perhaps there are exceptions you want to support?). We have not yet tested such nested OpenACC/OpenMP constructs (one thing after the other), so I'm not confident to claim support for that, and earlier on had already enable checking for such nesting. In r218569, I have now committed to gomp-4_0-branch the following patch to rework this: commit 30b2cd7ac340764d4f7eb14730b16a49e8799e32 Author: tschwinge Date: Wed Dec 10 09:52:42 2014 +0000 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 | 212 +++++++++++++-------- 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, 288 insertions(+), 162 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 06e8583..970e744 100644 --- gcc/ChangeLog.gomp +++ 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 gcc/omp-low.c gcc/omp-low.c index 39e2f22..d16e2de 100644 --- gcc/omp-low.c +++ 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)) + /* 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; diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 59f41a8..d52c7c0 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -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" } */ } } diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c index a489d2d..4fbf018 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ 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 gcc/testsuite/c-c++-common/goacc/nesting-2.c gcc/testsuite/c-c++-common/goacc/nesting-2.c deleted file mode 100644 index 0d350c6..0000000 --- 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 gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 00dc602..a833806 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ 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 gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 index 33cb9cb..8b8e989 100644 --- gcc/testsuite/gfortran.dg/goacc/parallel-kernels-regions.f95 +++ 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" }