From patchwork Tue Aug 1 15:35:16 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1815552 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: legolas.ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha256 header.s=default header.b=jR1KC4UE; dkim-atps=neutral Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4RFfMX6K4yz1yZl for ; Wed, 2 Aug 2023 01:35:51 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 547FD385841E for ; Tue, 1 Aug 2023 15:35:49 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 547FD385841E DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1690904149; bh=9dNAot2cKcx1nIF+APx1WJMc6TxIbxEKHD40LzSLdH8=; h=Date:Subject:To:List-Id:List-Unsubscribe:List-Archive:List-Post: List-Help:List-Subscribe:From:Reply-To:Cc:From; b=jR1KC4UEUyPnhEhwnaIOcy5LUxHMK5X1zqGu8K6f3BHjR0RKt36365Y2sG15Ru9mk CXy++l5wWI4kG6daebieMpX4X4q7QPGXoxybVYyjpfW+Q1pGUax3H6rFngAxFC/xYA LZNxFokCNKvgz7ywBjuPGnBESsVXquw06lxeMSNc= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from APC01-TYZ-obe.outbound.protection.outlook.com (mail-tyzapc01on2072.outbound.protection.outlook.com [40.107.117.72]) by sourceware.org (Postfix) with ESMTPS id 678B43858D28 for ; Tue, 1 Aug 2023 15:35:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 678B43858D28 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=FQb5mE7sA4eTSBANVpXTWgcKUxyJ3vAIAdR6vaXmRwJEZBTQrw3vI59DlyPoVz+YY1nbWT3z2PHVB9O3dhHqWKHjauOA81hjqVXtlzbNNkdrl5nrQV+1mVXjLljzc/sgvINwYr84f0F4Y6c3zGEoCvdN/0pG540c7agOaKdI+TLnMFdEfUESPsTWqYrSufz6uwu4/4Fg9ej/p2MGfx7Ldjc9l/YUMTOvwajaJudMgW4EvpUtNd/nuGTFSp3h94k/5tFFZTR/xnbMAAVOyL/KtS0j+U/6DV44GD9uZ7mTricAPDB6HnyW9B2+fQ9etLPFoZRJdmno/k/8hBagnp2K7A== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector9901; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=9dNAot2cKcx1nIF+APx1WJMc6TxIbxEKHD40LzSLdH8=; b=NgFNd5e02ZlIkVLOqHvNuUrHyv653MVZ3pE4XejVewu+u9BlcYPZq0PvM4RokX/Av641rEChNc0pJKljQsAj2wGtkdfb04tAdV0XlZAEtQFQlI4KDzhF8aNXVx+UbhZ6lPfBw1dXxvf5cXCkx1MvQW/h5oDl+mUixnGoxeWGE9gTEPhuBLqVPL+s5aTuh5PM8Wkj+1nFPN/fL2JYRGWdUNkhxry4WJOZRDDh89ONNiopGC/X4lDIMcjJza6BRABUkXTtwOlWOqdf2R0UlFm6wz8gJ687W6Tr9MNcNdDTrpCeZfu3ejb1Ov2V8pBko03a3rMZ6Z6Gap4iE1kGsArIPw== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=siemens.com; dmarc=pass action=none header.from=siemens.com; dkim=pass header.d=siemens.com; arc=none Received: from SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) by KL1PR0601MB4260.apcprd06.prod.outlook.com (2603:1096:820:79::10) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6631.45; Tue, 1 Aug 2023 15:35:21 +0000 Received: from SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e720:fef7:1199:a273]) by SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e720:fef7:1199:a273%6]) with mapi id 15.20.6631.043; Tue, 1 Aug 2023 15:35:21 +0000 Message-ID: <3f274de5-cc52-39c7-c399-f85a1a1a4640@siemens.com> Date: Tue, 1 Aug 2023 23:35:16 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.15; rv:102.0) Gecko/20100101 Thunderbird/102.13.0 Content-Language: en-US Subject: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs To: gcc-patches , Thomas Schwinge , Catherine Moore X-ClientProxiedBy: TYCP286CA0317.JPNP286.PROD.OUTLOOK.COM (2603:1096:400:3b7::7) To SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: SG2PR06MB5430:EE_|KL1PR0601MB4260:EE_ X-MS-Office365-Filtering-Correlation-Id: baec7127-a41e-4cbc-f630-08db92a4ea4d X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: nDFDEOZ3YLHbEE6gQyb5P/OPkMtIr5WnybQmMOpq59q728T0akCdXrWldRaOy+aiGrjDIHJG7mH+EwCcJgDylCU6FO/J6M6yk6LVQxYS2UD2SC918o6U3GM6b/uDzf5sJI3U+M05IcGT7/U+YOL8N4BrC9kkniOWT1/H6vqsf0te7HAUHsJgxBHkxcmCwerlKKLzrEikCuE9NBQ9NmiI9ydoMNJZ3D6+laraGsQSwj/hOqSJntoiX5BtEW1iTTg4GZ4NkRx+2h3PKL7rU2hb9rqdUbwkuVrJaSIqT4gb40oAgq5uqEZTQI+TKT8qL4WqR///hItl+WTalu1q5BYt5P9fvGpY9umOawX8a3H8K4DFDkRJbx3ht+V2wDKChP8GNcYns0dRbaZepPNHHRKjBn6nnvSXteZSBnPo2TbMNQ36pK1Gjz2BsrRkaSY/4JTLOCPQhQvQRy5ZHKqt1B0Oqn+dwGEaz6lPT5m4DB4AG2ucSfKPC18uMMgnS+5pKyB4mssnCJWQ+x0FJ27JYh7m5GI+730dtz9npAx9okvLecIOYsOD8EwxYIZ5QxfBAd7GVRmVyPCdvmN9vmO/aDPlQBn+M7xlD3dRcOGZVizI3l8Q9/36hWYAwi/Hb8SV9+7pG46rbPwP9dADTY5yP1lxZg== X-Forefront-Antispam-Report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:SG2PR06MB5430.apcprd06.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230028)(4636009)(366004)(396003)(39860400002)(136003)(346002)(376002)(451199021)(66946007)(66476007)(66556008)(83380400001)(2616005)(36756003)(82960400001)(31696002)(86362001)(38100700002)(110136005)(478600001)(6486002)(6512007)(6666004)(33964004)(6506007)(186003)(26005)(2906002)(31686004)(5660300002)(41300700001)(235185007)(316002)(84970400001)(8936002)(8676002)(43740500002)(45980500001); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?q?BAsdUCwdBdZMbpQPyaN62fOxRoKq?= =?utf-8?q?UVB3NaEaRP+wzvfFCr0/M+VEVoxD1zung/C0dJdhoppMHV79rvA0JHjVphSAWLIlV?= =?utf-8?q?3tQ0cj8smKJwTNtVZs38WeZcwEm9v2fdan0FgYs00/6OG/VQhpELW6L/gLdeLto0H?= =?utf-8?q?KE4f4uN3hDHRsB6wqAfJ4tceL4a9m09YvyDZbmPP6bSKxy3aGBR0GyCoq2/EOgaU5?= =?utf-8?q?z/vXTLNTcHt0/QTsAfRZzkYI+LIT82UKQaQkBEbcuORRUrNvmG4nsAA7z+eG2ISCW?= =?utf-8?q?kbzWotidudkY3F7sJAANSny0jV7HquQQ63/kQGRcd77lLnVqc7O2fNCKyMvqJREkG?= =?utf-8?q?fhca/fsgnrfnyRSjloa9XLZNUBLxlE1NXAwlneAYPBISgvHTMSbSDu3o9weSOcD75?= =?utf-8?q?6jUquzebMKA7KgVlRcAelOJpZN0UJghMfK9CHLqY5c9NDGqj4/qH9TAprc6FfOgOt?= =?utf-8?q?gxoqh9siE5obPJkIIytDQ6sHtTiw0C8sgGOvIqsVKVa2RbAxxhTO3Zh+rGD6js9Gq?= =?utf-8?q?IlkHgslg6doSrXvFocFejl5/RGecJaT77qeDZX6MUoBQyrbAzjAHwcvLd/p5Jr7NJ?= =?utf-8?q?7cuVu4AsxRDKoO7WsOgtPauLlqq9U0ICJm6fpnPzpEYZ1oxXwSdEenYLwNlMfangJ?= =?utf-8?q?mMsVWZNgOMCvQb2TMt727Pe+3zCLeJaCiTpTai78CPePmU7xTMktXJPuYZjYuFqFk?= =?utf-8?q?ue7Mn3Xwjq6WUFNJ6pLrAIBlHBSUpHVl7t7Ddo1/+/vCa99jANo7cQ+1qlQBlkhoE?= =?utf-8?q?yZJmFowNBGBrpnBwpV39pyPf4lxTQgr6ISal0/iSzBa4h2TJVioNZoT+Ai2foasJ+?= =?utf-8?q?f6mIPWgHexOdtbFbQ9gi6zAGfDlLPLtX9qLnY7td6pARi8591RqW/yznbzL4WKnBj?= =?utf-8?q?WwG7VCYFI1OECvQu+4qo3MBRZt4e3NhaZvSGtJD3LtVRJA2biCGvDu6CK1TrZRSEx?= =?utf-8?q?igHh0elGhNkHy+M/ZZwdYiwSY25ixgkqT8vDwd7DiI4M6T7g1+mL2fcAAxcnewFEp?= =?utf-8?q?Y+wfJQ4ntAHJvOImuoJ+QILaobGk+rlEPTN+EYSwy9Q14iGwbjaYlOJpVweTQ0CR+?= =?utf-8?q?zYe87JXgF7zJZqS3OaiWgjSUy9TayHlMIcBeuKT6016aZJenGbKqpsw15OZ9TVDbS?= =?utf-8?q?szZeoCkxU2X0BLD5WSKqqH/DHINffhcdPW08oISw3yozxhdZ6JurvpOXhQbEbIoSQ?= =?utf-8?q?CckBK9i9a2HonojCbcrwFXC6lT5QVVNrac8ThINDF7g6xUoG+yhmFDa3EdDPqMqNC?= =?utf-8?q?SADU4VlKPXVx0NVJex3P/O0fRro1xLQCoMc0FP1tMcbpNsn8n+ZcIf/Lgr6FpYhuH?= =?utf-8?q?7wGmryqKz/40V5b7M2jICmWdmfHPcjQGYlF4UyBR4jZ0QfW1exudbTZ6zZ7FZ+auY?= =?utf-8?q?u7XZj/cAlHmnFso1B5XWKom/QKqXuNK7HCa0x5EK2t96J/lsFFFscvNf4Ly2yJfPx?= =?utf-8?q?VMoaoZA0IyXcDtStCco/m4RIRXY/GE29wvIKYtYB0tMBtt5+nkPLZPPJ88zNSjhxD?= =?utf-8?q?SE1beERhitUDIcNqZMgwIvvzV3Ex+HCx2A=3D=3D?= X-OriginatorOrg: siemens.com X-MS-Exchange-CrossTenant-Network-Message-Id: baec7127-a41e-4cbc-f630-08db92a4ea4d X-MS-Exchange-CrossTenant-AuthSource: SG2PR06MB5430.apcprd06.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 01 Aug 2023 15:35:21.1139 (UTC) X-MS-Exchange-CrossTenant-FromEntityHeader: Hosted X-MS-Exchange-CrossTenant-Id: 38ae3bcd-9579-4fd4-adda-b42e1495d55a X-MS-Exchange-CrossTenant-MailboxType: HOSTED X-MS-Exchange-CrossTenant-UserPrincipalName: sgeywVGbav3+CvFdU/sXZVOXcZDJhIgi1BKcR2urf3rbNvxFaT7PZn6jHz1JbpTny5COkWXqMMOcuq6/7+WlhO7CvOVozr+61MePDpZhKGk= X-MS-Exchange-Transport-CrossTenantHeadersStamped: KL1PR0601MB4260 X-Spam-Status: No, score=-7.0 required=5.0 tests=BAYES_00, DKIMWL_WL_MED, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FORGED_SPF_HELO, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H2, SPAM_BODY, SPF_HELO_PASS, SPF_NONE, TXREP, T_SCC_BODY_TEXT_LINE 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: , X-Patchwork-Original-From: Chung-Lin Tang via Gcc-patches From: Chung-Lin Tang Reply-To: cltang@codesourcery.com Cc: Chung-Lin Tang Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi Thomas, this is v2 of the patch for implementing the OpenACC 2.7 addition of default(none|present) support for data constructs. 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. This supports displaying the location/type of OpenACC construct where the default-clause is in the error messages. 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. Note, I got rid of the dummy OMP_CLAUSE_DEFAULT creation in this version, since it seemed not really needed. Re-tested on master on powerpc64le-linux/nvptx. Okay to commit? Thanks, Chung-Lin 2023-08-01 Chung-Lin Tang 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 (struct gimplify_omp_ctx): Add oacc_default_clause_ctx field. (new_omp_context): Initialize oacc_default_clause_ctx field. (oacc_region_type_name): New function. (oacc_default_clause): Lookup current default_kind value from ctx->oacc_default_clause_ctx, adjust default(none) error and inform message dumping. (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-5.c: Adjust testcase. * gfortran.dg/goacc/default-3.f95: Adjust testcase. * gfortran.dg/goacc/default-5.f: Adjust testcase. diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 24a6eb6e459..974f0132787 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18196,6 +18196,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 d7ef5b34d42..bc59fbeac20 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45860,6 +45860,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 2952cd300ac..c37f843ec3b 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..ec0ccc67da8 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -225,6 +225,7 @@ struct gimplify_omp_ctx vec loop_iter_var; location_t location; enum omp_clause_default_kind default_kind; + struct gimplify_omp_ctx *oacc_default_clause_ctx; enum omp_region_type region_type; enum tree_code code; bool combined_loop; @@ -459,6 +460,10 @@ new_omp_context (enum omp_region_type region_type) c->default_kind = OMP_CLAUSE_DEFAULT_SHARED; else c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; + if (gimplify_omp_ctxp) + c->oacc_default_clause_ctx = gimplify_omp_ctxp->oacc_default_clause_ctx; + else + c->oacc_default_clause_ctx = c; c->defaultmap[GDMK_SCALAR] = GOVD_MAP; c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP; c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP; @@ -7699,6 +7704,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 +7730,6 @@ 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; bool on_device = false; bool is_private = false; bool declared = is_oacc_declared (decl); @@ -7735,17 +7758,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) flags |= GOVD_MAP_TO_ONLY; } + /* Use the enclosing construct with a default clause to set the current + default kind. */ + enum omp_clause_default_kind default_kind + = ctx->oacc_default_clause_ctx->default_kind; + 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 (default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7758,8 +7784,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 +7791,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 (default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7785,16 +7809,20 @@ 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 (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)); + } + else if (default_kind == OMP_CLAUSE_DEFAULT_PRESENT) ; /* Handled above. */ else - gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED); + gcc_checking_assert (default_kind == OMP_CLAUSE_DEFAULT_SHARED); return flags; } @@ -12124,6 +12152,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_DEFAULT: ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); + if (flag_openacc) + /* For OpenACC, set current context to the 'active' one for + default-clause lookup. */ + ctx->oacc_default_clause_ctx = ctx; break; case OMP_CLAUSE_INCLUSIVE: diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c index ac169a903e9..e68f33606fd 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-3.c +++ b/gcc/testsuite/c-c++-common/goacc/default-3.c @@ -4,7 +4,7 @@ void f1 () { int f1_a = 2; float f1_b[2]; - + #pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */ { f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */ @@ -15,4 +15,49 @@ void f1 () 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-message "enclosing OpenACC .data. construct" } */ +#pragma acc kernels + { + 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 data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */ +#pragma acc parallel + { + 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-message "enclosing OpenACC .parallel. construct" } */ + { + 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-message "enclosing OpenACC .data. construct" } */ +#pragma acc data +#pragma acc data +#pragma acc parallel + { + 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-message "enclosing OpenACC .data. construct" } */ +#pragma acc data +#pragma acc parallel + { + 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-message "enclosing OpenACC .data. construct" } */ +#pragma acc parallel + { + 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-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c index 37e3c3555cd..76526396124 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,18 @@ 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 (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..09b8c1bbd87 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 @@ -15,4 +15,65 @@ subroutine f1 = 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-message "enclosing OpenACC .data. construct" } + !$acc kernels + 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-message "enclosing OpenACC .data. construct" } + !$acc parallel + 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-message "enclosing OpenACC .parallel. construct" } + 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-message "enclosing OpenACC .data. construct" } + !$acc data + !$acc data + !$acc parallel + 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-message "enclosing OpenACC .data. construct" } + !$acc data + !$acc parallel + 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-message "enclosing OpenACC .data. construct" } + !$acc parallel + 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-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f index 9dc83cbe601..9260db07a1a 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,17 @@ ! { 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 (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 END SUBROUTINE F1