From patchwork Wed Dec 15 15:54:44 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48979 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9912C3857C4B for ; Wed, 15 Dec 2021 16:18:45 +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 7FF2A3858420 for ; Wed, 15 Dec 2021 15:57:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 7FF2A3858420 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 9Eh0D7v/9JzNQug90Znhsvuoi2pK0r/jCk04MxjWoUY3bXcZ9fx2CHsk3Wq5hwXdaZDl22ZD6I w0ltOprbnT4Awk20VcSQzAnnrIRfemfDbvst5M7vQ1wW0IlETeokKUi639h16htDRP9Iu/RH1v CayaCptvJ8ayfyFQrHgPNPRkwe8ss8kACXKhrHWaGBDYlHuU2UrirfCtMlr5j+kKJ4Sj5oZy1A Ndw/wtYMVB1LGSGDgf74Z9ttturuPxhd8R8O+RxD8LcUtlZZ/xAUs2aQ+MmlODTGqLmysBvEH8 FKDpAv8UUir9Ii9GjZh8e2Sc X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69584655" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:57:25 -0800 IronPort-SDR: CE7Q0vcU39ZxLn3BrYMMdkn+eGTdSPm+7C3ncMi6Pt1t3578DZezQZko6uhaetcx/KZGQOkTlg grGUFXFGUeXI+XQCS6MDvIKsdMM811Or/CT15A1ZRQnmo1wF81hqIIUyKMr2qkcl3JzuQmNTHa ab1GJIIzkETNzRMKVdRVjW2ujCOM0jEDONCcHI/BHLUfnv7viaKIRt7ZOFNOuc+1615nBzFxFX XGL7yymYwt06SEMJMkCHpTL5EzwzQbIhAa230I3M73WEW16yW9gPl3Fh5Qyy502e5hv15+Qp4c ruE= From: Frederik Harwath To: Subject: [PATCH 37/40] Fix for is_gimple_reg vars to 'data kernels' Date: Wed, 15 Dec 2021 16:54:44 +0100 Message-ID: <20211215155447.19379-38-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: , Cc: Tobias Burnus , thomas@codesourcery.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Tobias Burnus Nearly all variable mapping is moved from 'kernels' to a surrounding 'data kernels' and then 'force_present' mapped for the 'kernels'. However, as libgomp.oacc-c-c++-common/declare-vla.c shows, moving 'int i, N' will fail as there is a special case for is_gimple_reg in mapping and that fails badly if outside a target region (e.g. offloading = false). As those are transferred by value and not as a pointer, it makes more sense to only map them at 'kernels' and ignore them for 'data kernels'. Additionally, as e.g. libgomp.oacc-c-c++-common/kernels-decompose-1.c shows, one still additionally to handle 'kernels'-declared variables which now are declared in 'kernels data' and and can be handled as is_gimple_reg. gcc/ * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region): is_gimple_reg vars are not yet mapped, fall through to map is as before the transformation. (omp_oacc_kernels_decompose_1): Don't map is_gimple_reg vars. (decompose_kernels_region_body): Use tofrom for is_gimple_reg vars. (omp_oacc_kernels_decompose_1): Handle is_gimple_reg vars as without data kernels. gcc/testsuite/ * gfortran.dg/goacc/declare-3.f95: Update scan-tree-dump-times. --- gcc/omp-oacc-kernels-decompose.cc | 9 +++++++-- gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 2 +- 2 files changed, 8 insertions(+), 3 deletions(-) -- 2.33.0 ----------------- 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 diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index c96207d96250..a6be1f1ed238 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -873,7 +873,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body, else inner_bind_vars = next; } - else + else if (!is_gimple_reg (v)) { /* Otherwise, build the map clause. */ tree new_clause = build_omp_clause (loc, OMP_CLAUSE_MAP); @@ -1222,7 +1222,9 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses) if (!DECL_ARTIFICIAL (var) && TREE_CODE (var) != CONST_DECL) { tree present_clause = build_omp_clause (loc, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (present_clause, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_SET_MAP_KIND (present_clause, + is_gimple_reg (var) + ? GOMP_MAP_TOFROM : GOMP_MAP_FORCE_PRESENT); OMP_CLAUSE_DECL (present_clause) = var; OMP_CLAUSE_SIZE (present_clause) = DECL_SIZE_UNIT (var); OMP_CLAUSE_CHAIN (present_clause) = present_clauses; @@ -1437,6 +1439,9 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt) region causes runtime errors. */ break; + if (is_gimple_reg (decl)) + break; + /* For non-artificial variables, and for non-declaration expressions like A[0:n], copy the clause to the data region. */ diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 index 9127cba6600d..2a1fe0a68465 100644 --- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 @@ -39,7 +39,7 @@ program test use mod_d use mod_e - ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_to:b\) map\(force_alloc:a\)$} original } } + ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(to:b\) map\(alloc:a\)$} original } } end program test ! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } }