From patchwork Tue Aug 15 14:35:54 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 135667 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:a888:0:b0:3f2:4152:657d with SMTP id x8csp475610vqo; Tue, 15 Aug 2023 07:38:53 -0700 (PDT) X-Google-Smtp-Source: AGHT+IEidkzfvAAHVSFeTT721ZoNA0j7KbdcxQ3mo/Ii5w34GIc9gZeoiSwc3E5wyG8aCueXtAqe X-Received: by 2002:a50:fc14:0:b0:525:6d9e:67c0 with SMTP id i20-20020a50fc14000000b005256d9e67c0mr3445934edr.23.1692110333511; Tue, 15 Aug 2023 07:38:53 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1692110333; cv=none; d=google.com; s=arc-20160816; b=avxxpq9SZoynnhKR/zLDnTq178Ouy9bEAsnQ1MhDnmo2xxgZTt9oKuoMP7gNJKWxnt GmjiqDnD4ji1WXuAb3ja9Jw1FViyHcADDnI2LC8Ajdp7SakFs1uggis/Q/fMViUM4Bt4 4MdG9QFRaQHEUzhT4RT+Lz7AH9BUs49jYo7nGyIkarBIN2HSVdI3kmbtxgb2M8cDQed9 I6dJ1zeHBhL3igKmKnQYHFu/RJpccwo6aL0E9HivkNZf8ZAVIfaf3nM2ZzLMVYnxN8xp 8A428KI5tP79Wa8v22OUfgO9kWQyTn7BDft1ME9AGTQ9ha8S9F0O/4Gft+7liZT4CJcF 6OYA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:mime-version:message-id:date :user-agent:references:in-reply-to:subject:cc:to:from:ironport-sdr :dmarc-filter:delivered-to; bh=1ypYK0vOwNzRtqZZhVob7qCrB5LA1paZzwu+gnN0pZU=; fh=2diLeGlUOqBfoL+uKdAAgpzRS33i5FnH83HB6UAXwlI=; b=BGqHP9Z2cYw8y82TwWTGo7dyCMd8yC0EYXVQnRWrp1Eb+58UEm0/wlqYwfppBK+HrM Q3w16J41pFN0b6hhWkZV2qFCdIDC+OpkDGR3RHPwtjjnA59JvDyhw2hoMw5o5y3JDz42 Rw11PqnxAnZEM7IWR7ssyQ9aBxX7ElEeFhMjSNGNUCAYcttej40leCJzTXBHbt+tVwy5 Jrol9mDiy42EOW8Awo3yEx/pmr1rc5eNWmv7GB+vyhVCu7koiA0gFLFbPHX+UxU+787z uF2HzashpEwrsnuNkqzC4sh3I1LVcHqv7+pV6XUj6Ppa8tyemx6WY4euMR5zABXWngJ0 rYVw== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id t26-20020aa7d71a000000b0052220448d66si9362935edq.283.2023.08.15.07.38.53 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Tue, 15 Aug 2023 07:38:53 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 0F5FE3855581 for ; Tue, 15 Aug 2023 14:38:28 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id CD61A3857357 for ; Tue, 15 Aug 2023 14:37:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CD61A3857357 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,174,1684828800"; d="scan'208,223";a="14479567" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 15 Aug 2023 06:36:07 -0800 IronPort-SDR: 0qjWSDz9hg+jpeszkHh//JwQRdBHX4x3a1kwnpNP6ZbAGjYZiHcg88O3x0NUHIiuvH65Da4esd 1OCXOVdAkCijCbsd4qDW93lXy29rPyECBIo1uZUGEn9NXQNKWOjBcW6iRV1lFraMmAwamK/xHO FWaZyvusU/kicoJjHVcAHsbbKwRk3vB+fJ+pEDdaZaRmVmIkJBVnIYGPQde9fKSHhzxYqkY5Fv 9XAO+YUoPfu3RUhAgTv+8ZMNQ49rBbsnujnC1bpf6+lBUcq323SQMntKWozM/0CMyYnK6dWkIb bm0= From: Thomas Schwinge To: Chung-Lin Tang , CC: Catherine Moore Subject: [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs) In-Reply-To: <3f274de5-cc52-39c7-c399-f85a1a1a4640@siemens.com> References: <3f274de5-cc52-39c7-c399-f85a1a1a4640@siemens.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Tue, 15 Aug 2023 16:35:54 +0200 Message-ID: <87jztwcr05.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-8.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPAM_BODY, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1774306285326999850 X-GMAIL-MSGID: 1774306285326999850 Hi! On 2023-08-01T23:35:16+0800, Chung-Lin Tang wrote: > this is v2 of the patch for implementing the OpenACC 2.7 addition of > default(none|present) support for data constructs. Thanks! > Instead of propagating an additional 'oacc_default_kind' for OpenACC, > this patch does it in a more complete way: it directly propagates the > gimplify_omp_ctx* pointer of the inner most context where we found > a default-clause. Right -- but reviewing this, it came upon me that we don't need any such new code at all, and instead may in 'gcc/gimplify.cc:oacc_default_clause' simply look through the 'ctx's to find the 'default' clause information. This centralizes the logic in the one place where it's relevant. > This supports displaying the location/type of OpenACC > construct where the default-clause is in the error messages. This is preserved... > The testcases also have the multiple nested data construct testing added, > where we can now have messages referring precisely to the exact innermost > default clause that was active at that program point. ..., but we should also still 'inform' about the compute construct, where the user is expected to add explicit data clauses (if not adding to the 'data' construct where the 'default(none)' clause appears): > --- a/gcc/gimplify.cc > +++ b/gcc/gimplify.cc > @@ -7785,16 +7809,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) > - else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE) > + else if (default_kind == OMP_CLAUSE_DEFAULT_NONE) > { > error ("%qE not specified in enclosing OpenACC %qs construct", > - DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind); > - inform (ctx->location, "enclosing OpenACC %qs construct", rkind); > - } > - else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT) > + DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), > + oacc_region_type_name (ctx->region_type)); > + inform (ctx->oacc_default_clause_ctx->location, > + "enclosing OpenACC %qs construct", > + oacc_region_type_name > + (ctx->oacc_default_clause_ctx->region_type)); > + } That is, we should keep here the original 'inform' for 'ctx->location', and *add another* 'inform' for 'ctx->oacc_default_clause_ctx->location'. Otherwise that's confusing to users. Instead of requiring another iteration through you, I've now implemented that, and with test cases enhanced some more, pushed to master branch commit bed993884b149851fe930b43cf11cbcdf05f1578 "OpenACC 2.7: default clause support for data constructs", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From bed993884b149851fe930b43cf11cbcdf05f1578 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Tue, 6 Jun 2023 03:46:29 -0700 Subject: [PATCH] OpenACC 2.7: default clause support for data constructs This patch implements the OpenACC 2.7 addition of default(none|present) support for data constructs. Now, specifying "default(none|present)" on a data construct turns on same default clause behavior for all lexically enclosed compute constructs (which don't already themselves have a default clause). gcc/c/ChangeLog: * c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/cp/ChangeLog: * parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/fortran/ChangeLog: * openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT. gcc/ChangeLog: * gimplify.cc (oacc_region_type_name): New function. (oacc_default_clause): If no 'default' clause appears on this compute construct, see if one appears on a lexically containing 'data' construct. (gimplify_scan_omp_clauses): Upon OMP_CLAUSE_DEFAULT case, set ctx->oacc_default_clause_ctx to current context. gcc/testsuite/ChangeLog: * c-c++-common/goacc/default-3.c: Adjust testcase. * c-c++-common/goacc/default-4.c: Adjust testcase. * c-c++-common/goacc/default-5.c: Adjust testcase. * gfortran.dg/goacc/default-3.f95: Adjust testcase. * gfortran.dg/goacc/default-4.f: Adjust testcase. * gfortran.dg/goacc/default-5.f: Adjust testcase. Co-authored-by: Thomas Schwinge --- gcc/c/c-parser.cc | 1 + gcc/cp/parser.cc | 1 + gcc/fortran/openmp.cc | 3 +- gcc/gimplify.cc | 64 +++++++++++---- gcc/testsuite/c-c++-common/goacc/default-3.c | 59 +++++++++++++- gcc/testsuite/c-c++-common/goacc/default-4.c | 42 ++++++++++ gcc/testsuite/c-c++-common/goacc/default-5.c | 19 ++++- gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 77 ++++++++++++++++++- gcc/testsuite/gfortran.dg/goacc/default-4.f | 36 +++++++++ gcc/testsuite/gfortran.dg/goacc/default-5.f | 19 ++++- 10 files changed, 298 insertions(+), 23 deletions(-) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index cabb18d04f7..33fe7b115ff 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18284,6 +18284,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 7f646704d3f..774706ac607 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45854,6 +45854,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 234d896b2ce..bee3015e484 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -3802,7 +3802,8 @@ error: #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ - | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH) + | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH \ + | OMP_CLAUSE_DEFAULT) #define OACC_LOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \ diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 320920ed74c..7549436944c 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7699,6 +7699,25 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, return flags; } +/* Return string name for types of OpenACC constructs from ORT_* values. */ + +static const char * +oacc_region_type_name (enum omp_region_type region_type) +{ + switch (region_type) + { + case ORT_ACC_DATA: + return "data"; + case ORT_ACC_PARALLEL: + return "parallel"; + case ORT_ACC_KERNELS: + return "kernels"; + case ORT_ACC_SERIAL: + return "serial"; + default: + gcc_unreachable (); + } +} /* Determine outer default flags for DECL mentioned in an OACC region but not declared in an enclosing clause. */ @@ -7706,7 +7725,23 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, static unsigned oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) { - const char *rkind; + struct gimplify_omp_ctx *ctx_default = ctx; + /* If no 'default' clause appears on this compute construct... */ + if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED) + { + /* ..., see if one appears on a lexically containing 'data' + construct. */ + while ((ctx_default = ctx_default->outer_context)) + { + if (ctx_default->region_type == ORT_ACC_DATA + && ctx_default->default_kind != OMP_CLAUSE_DEFAULT_SHARED) + break; + } + /* If not, reset. */ + if (!ctx_default) + ctx_default = ctx; + } + bool on_device = false; bool is_private = false; bool declared = is_oacc_declared (decl); @@ -7738,14 +7773,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) switch (ctx->region_type) { case ORT_ACC_KERNELS: - rkind = "kernels"; - if (is_private) flags |= GOVD_FIRSTPRIVATE; else if (AGGREGATE_TYPE_P (type)) { /* Aggregates default to 'present_or_copy', or 'present'. */ - if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) + if (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7758,8 +7791,6 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) case ORT_ACC_PARALLEL: case ORT_ACC_SERIAL: - rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial"; - if (is_private) flags |= GOVD_FIRSTPRIVATE; else if (on_device || declared) @@ -7767,7 +7798,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) else if (AGGREGATE_TYPE_P (type)) { /* Aggregates default to 'present_or_copy', or 'present'. */ - if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) + if (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7785,16 +7816,23 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) if (DECL_ARTIFICIAL (decl)) ; /* We can get compiler-generated decls, and should not complain about them. */ - else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE) + else if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_NONE) { error ("%qE not specified in enclosing OpenACC %qs construct", - DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind); - inform (ctx->location, "enclosing OpenACC %qs construct", rkind); - } - else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT) + DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), + oacc_region_type_name (ctx->region_type)); + if (ctx_default != ctx) + inform (ctx->location, "enclosing OpenACC %qs construct and", + oacc_region_type_name (ctx->region_type)); + inform (ctx_default->location, + "enclosing OpenACC %qs construct with %qs clause", + oacc_region_type_name (ctx_default->region_type), + "default(none)"); + } + else if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_PRESENT) ; /* Handled above. */ else - gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED); + gcc_checking_assert (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED); return flags; } diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c index ac169a903e9..73dbc908475 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-3.c +++ b/gcc/testsuite/c-c++-common/goacc/default-3.c @@ -4,13 +4,66 @@ void f1 () { int f1_a = 2; float f1_b[2]; - -#pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */ + +#pragma acc kernels default (none) /* { dg-note "enclosing OpenACC 'kernels' construct with 'default\\\(none\\\)' clause" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */ + } +#pragma acc parallel default (none) /* { dg-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } + +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc kernels /* { dg-note "enclosing OpenACC 'kernels' construct and" } */ { f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */ } -#pragma acc parallel default (none) /* { dg-message "enclosing OpenACC .parallel. construct" } */ +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } + +#pragma acc data default (none) +#pragma acc parallel default (none) /* { dg-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } + +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc data +#pragma acc data +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } +#pragma acc data +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc data +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } +#pragma acc data +#pragma acc data +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ + = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ + } +#pragma acc data +#pragma acc data default (none) +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ { f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c index 867175d4847..e12cb86d097 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-4.c +++ b/gcc/testsuite/c-c++-common/goacc/default-4.c @@ -44,6 +44,27 @@ void f2 () } } +void f2_ () +{ + int f2__a = 2; + float f2__b[2]; + +#pragma acc data default (none) copyin (f2__a) copyout (f2__b) + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2__b \[^\\)\]+\\) map\\(to:f2__a \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } } */ + { +#pragma acc kernels + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */ + { + f2__b[0] = f2__a; + } +#pragma acc parallel + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */ + { + f2__b[0] = f2__a; + } + } +} + void f3 () { int f3_a = 2; @@ -64,3 +85,24 @@ void f3 () } } } + +void f3_ () +{ + int f3__a = 2; + float f3__b[2]; + +#pragma acc data default (present) copyin (f3__a) copyout (f3__b) + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3__b \[^\\)\]+\\) map\\(to:f3__a \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } } */ + { +#pragma acc kernels + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } */ + { + f3__b[0] = f3__a; + } +#pragma acc parallel + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } */ + { + f3__b[0] = f3__a; + } + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c index 37e3c3555cd..59ac1d79668 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-5.c +++ b/gcc/testsuite/c-c++-common/goacc/default-5.c @@ -4,8 +4,8 @@ void f1 () { - int f1_a = 2; - float f1_b[2]; + int f1_a = 2, f1_c = 3; + float f1_b[2], f1_d[2]; #pragma acc kernels default (present) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } */ @@ -17,4 +17,19 @@ void f1 () { f1_b[0] = f1_a; } + + /* { dg-final { scan-tree-dump-times "omp target oacc_data default\\(present\\)" 2 "gimple" } } */ +#pragma acc data default (present) +#pragma acc kernels + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } */ + { + f1_d[0] = f1_c; + } +#pragma acc data default (none) +#pragma acc data default (present) +#pragma acc parallel + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } } */ + { + f1_d[0] = f1_c; + } } diff --git a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 index 98ed34200c6..c1edf4c8137 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 @@ -5,14 +5,87 @@ subroutine f1 integer :: f1_a = 2 real, dimension (2) :: f1_b - !$acc kernels default (none) ! { dg-message "enclosing OpenACC .kernels. construct" } + !$acc kernels default (none) ! { dg-note "enclosing OpenACC .kernels. construct with 'default\\\(none\\\)' clause" } f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } } = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 } !$acc end kernels - !$acc parallel default (none) ! { dg-message "enclosing OpenACC .parallel. construct" } + !$acc parallel default (none) ! { dg-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" } f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } !$acc end parallel + + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc kernels ! { dg-note "enclosing OpenACC 'kernels' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 } + !$acc end kernels + !$acc end data + + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + + !$acc data default (none) + !$acc parallel default (none) ! { dg-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc data + !$acc data + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + !$acc end data + !$acc end data + + !$acc data + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc data + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + !$acc end data + !$acc end data + + !$acc data + !$acc data + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + !$acc end data + !$acc end data + + !$acc data + !$acc data default (none) + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } } + = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } + ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 } + !$acc end parallel + !$acc end data + !$acc end data + !$acc end data + end subroutine f1 diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f index 30f411f70ab..4e89b6859ba 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-4.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f @@ -38,6 +38,24 @@ !$ACC END DATA END SUBROUTINE F2 + SUBROUTINE F2_ + IMPLICIT NONE + INTEGER :: F2__A = 2 + REAL, DIMENSION (2) :: F2__B + +!$ACC DATA DEFAULT (NONE) COPYIN (F2__A) COPYOUT (F2__B) +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2__a \[^\\)\]+\\) map\\(from:f2__b \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } } +!$ACC KERNELS +! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } + F2__B(1) = F2__A; +!$ACC END KERNELS +!$ACC PARALLEL +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } + F2__B(1) = F2__A; +!$ACC END PARALLEL +!$ACC END DATA + END SUBROUTINE F2_ + SUBROUTINE F3 IMPLICIT NONE INTEGER :: F3_A = 2 @@ -55,3 +73,21 @@ !$ACC END PARALLEL !$ACC END DATA END SUBROUTINE F3 + + SUBROUTINE F3_ + IMPLICIT NONE + INTEGER :: F3__A = 2 + REAL, DIMENSION (2) :: F3__B + +!$ACC DATA DEFAULT (PRESENT) COPYIN (F3__A) COPYOUT (F3__B) +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3__a \[^\\)\]+\\) map\\(from:f3__b \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } } +!$ACC KERNELS +! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } + F3__B(1) = F3__A; +!$ACC END KERNELS +!$ACC PARALLEL +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } + F3__B(1) = F3__A; +!$ACC END PARALLEL +!$ACC END DATA + END SUBROUTINE F3_ diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f index 9dc83cbe601..2cb07a8cbca 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-5.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f @@ -4,8 +4,8 @@ SUBROUTINE F1 IMPLICIT NONE - INTEGER :: F1_A = 2 - REAL, DIMENSION (2) :: F1_B + INTEGER :: F1_A = 2, F1_C = 3 + REAL, DIMENSION (2) :: F1_B, F1_D !$ACC KERNELS DEFAULT (PRESENT) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } @@ -15,4 +15,19 @@ ! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } F1_B(1) = F1_A; !$ACC END PARALLEL + +!$ACC DATA DEFAULT (PRESENT) +!$ACC KERNELS +! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } + F1_D(1) = F1_C; +!$ACC END KERNELS +!$ACC END DATA +!$ACC DATA DEFAULT (NONE) +!$ACC DATA DEFAULT (PRESENT) +!$ACC PARALLEL DEFAULT (PRESENT) +! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } } + F1_D(1) = F1_C; +!$ACC END PARALLEL +!$ACC END DATA +!$ACC END DATA END SUBROUTINE F1 -- 2.34.1