From patchwork Tue Jun 16 22:38:31 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310759 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjmf5ffKz9sSc for ; Wed, 17 Jun 2020 08:39:14 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E8DB4388E800; Tue, 16 Jun 2020 22:38:59 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 56AF83858D34; Tue, 16 Jun 2020 22:38:56 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 56AF83858D34 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: 3eH+Sh9T+OKuomPFhFVvvPAcEVyW6bvaXLxIaJFN4l7niZjvitFZSbprevtLHYBmUVxj7Iu2Yj kzL4zL27tNFZt0pWyCmUq7t1++pYdrbArTq3RyVOrVLYE2cv79rtLekdNi8BeJxsFqp2u2x8Pn 5J/+lN2xf29/FZ7gF1f/HO2s5dVD0o1ojWoSfJUuvqhgDMLHlZpfVn7+vYqMV6igmW/CYccyaE sIJxH5cFBvjJMxP8QkUXMQL0qM3l4bE1fiGFGWqtP4dtZ+P1xVVgXcT5T0hHvpt73CbHI5q5LD weo= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="50003695" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:38:56 -0800 IronPort-SDR: d48PIm3j4Asi2/yCj437B6YQcaeGRcWN2GXja+aXpQx17knb+XHQx7oErXxDB+V0dz8bB3XUfi jHDBOb1TqA7H+yHmE9oSAd2qNZhh3vIhPBPQczMXtvnDzbp9hehE28ipXDqGy8AaK/fT5x8Orj SRRBeWOWbkcIK5D2uP2J6URoWx+GfP0CdzhDIc7vsjQ8zyAhgbOZbn9otKqAkhlqon4Cy6ndVz k5v1Lgxw+T6jsybEOLKb/KOADjkA9ziFzs4myqFphfQ0tY8nC2OsuduHQZYWyOB81oxglcNjHr gpo= From: Julian Brown To: Subject: [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix Date: Tue, 16 Jun 2020 15:38:31 -0700 Message-ID: <54822b42a9e8e1d04d447ba5276ed6c7510e51d7.1592343756.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a slightly-updated version of the patch sent here, with some of Thomas's suggestions incorporated: https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547407.html I'm still assuming this is approved, but including for completeness. Julian ChangeLog gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Use 'inner' not 'decl' for derived type members which themselves have derived types. gcc/testsuite/ * gfortran.dg/goacc/mapping-tests-3.f90: New test. * gfortran.dg/goacc/mapping-tests-4.f90: New test. --- gcc/fortran/trans-openmp.c | 4 ++-- .../gfortran.dg/goacc/mapping-tests-3.f90 | 15 +++++++++++++++ .../gfortran.dg/goacc/mapping-tests-4.f90 | 17 +++++++++++++++++ 3 files changed, 34 insertions(+), 2 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 7e2f6256c43..02c40fdc660 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2774,9 +2774,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } else { - OMP_CLAUSE_DECL (node) = decl; + OMP_CLAUSE_DECL (node) = inner; OMP_CLAUSE_SIZE (node) - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); } } else if (lastcomp->next diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode 100644 index 00000000000..890ca781967 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 @@ -0,0 +1,15 @@ +! { dg-additional-options "-fdump-tree-gimple" } + +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + + !$acc enter data copyin(x%A) +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } +end diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode 100644 index 00000000000..17cc4841d4e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 @@ -0,0 +1,17 @@ +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + +! This is accepted at present, although it represents a probably-unintentional +! overlapping subcopy. + !$acc enter data copyin(x%A, x%A%i) +! But this raises an error. + !$acc enter data copyin(x%A, x%A%i, x%A%i) +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target *-*-* } .-1 } +end From patchwork Tue Jun 16 22:38:32 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310760 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjmk4BkQz9sSc for ; Wed, 17 Jun 2020 08:39:18 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id DCB20388E839; Tue, 16 Jun 2020 22:39:01 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 9756F388A836; Tue, 16 Jun 2020 22:38:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 9756F388A836 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: T7QLzcGY/na/lyBUke0p8JKWadXQmRXqTJm1RgtCYYxaDb/ExU9cZLzUlh6Byo18vrkW9PVhS3 aIAvVjmMHNUD38b/wF8neKQctAhWvaVt1/4v2qYEgfaDlHX9u4873owudiiyKPjYxAJSYEOuMa GpRZwmejDI6IAAhvFPY2DWjndoQOpPN59vGDHJtioMjX5Ye6aQ4epvDWKi3CK6YkgBRe6MQC7I UBXkRQ1jn1s593U733CS21/HOsZyaMxd/H0gJ9K5aESUaKJTWk7ORllpBExQcU0f0gf0d0MXqx pKc= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="50003698" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:38:59 -0800 IronPort-SDR: 4Gsy7IetxftBkkkkrVmP9kZ9jGw3lCxFJwps5PsZz52D5/q1pPQwzrYjNHT7462BjDSkd/IvHf fUY749nV0uteOB+191Q+wH3t4WdsIwfovBbVNH4VdX4v4eaWNv17HBCAOwqcHKWx5zc83c0Ax5 OC/24BJdEvSF7Nn2wttdIOu5Sm1NLS+VdP3CND1MCytYRUwdVosODcv07Dalak3SECiWLWNIKm 2Si2FrPo2N2Ubv757Gm2HIMY4xhLfyXZyvr6efEtWo9GcIbUWL1tWyYgROxcmZ/uxIFazHFSbB cPE= From: Julian Brown To: Subject: [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last Date: Tue, 16 Jun 2020 15:38:32 -0700 Message-ID: <037fa9b35dfbe459776be8ef7b01eca95e3dc7a3.1592343756.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.4 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.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Later patches in this series assume that GOMP_MAP_ATTACH will be grouped together with a preceding GOMP_MAP_TO_PSET or other "to" data movement clause, except in cases where an explicit "attach" clause is used. This patch arranges for that to be so. OK? Julian ChangeLog libgomp/ * oacc-mem.c (find_group_last): Group data-movement clauses (GOMP_MAP_TO_PSET, GOMP_MAP_TO, etc.) together with a subsequent GOMP_MAP_ATTACH. Allow standalone GOMP_MAP_ATTACH also. --- libgomp/oacc-mem.c | 22 +++++++++++++++++++--- 1 file changed, 19 insertions(+), 3 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 936ae649dd9..be7f8d600eb 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -985,9 +985,15 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) switch (kind0) { case GOMP_MAP_TO_PSET: - while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER) + if (pos + 1 < mapnum + && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH) + return pos + 1; + + while (pos + 1 < mapnum + && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER) pos++; - /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */ + /* We expect at least one GOMP_MAP_POINTER (if not a single + GOMP_MAP_ATTACH) after a GOMP_MAP_TO_PSET. */ assert (pos > first_pos); break; @@ -1002,6 +1008,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) gomp_fatal ("unexpected mapping"); break; + case GOMP_MAP_ATTACH: + return pos; + default: /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other mapping. */ @@ -1012,9 +1021,16 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) return pos + 1; } + /* We can have a single GOMP_MAP_ATTACH mapping after a to/from + mapping. */ + if (pos + 1 < mapnum + && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH) + return pos + 1; + /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from (etc.) mapping. */ - while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER) + while (pos + 1 < mapnum + && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER) pos++; } From patchwork Tue Jun 16 22:38:33 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310761 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjms3NL1z9sSc for ; Wed, 17 Jun 2020 08:39:25 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 8DF4C388F056; Tue, 16 Jun 2020 22:39:06 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 3AD873851C0D; Tue, 16 Jun 2020 22:39:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 3AD873851C0D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: xNoa/wJQcOXUl5vVHgt8XlEZ0ZsxJy8Xy7+fG+/erfoG6McFtB2duceyj72wG5bYW25gXo3UCB MuR8UZXAOsSuaVLNaN63yHfi2oC5grlgkwPTNHSlAhqjsHZl8VpmB/4XVzsqa4g9MKi5M4Matt fqfO2lLd8epYuOzRyTjqTu6vYZakQC6Wr6FH6ZvzyurqGXLV0vt8JkrsZ024A/vp+I688KIuI7 tmgFiDDjjdR4Fw3O8hQijPVcDJB/h5w3gSXa8c9Gqm+wKJ68R2WzXUlEFOf1y8N+1zUpvbel5r mnM= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="50003701" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:39:02 -0800 IronPort-SDR: 0bhnJI/H1JGu4ToK47EeAwNzExiIR7paNc2FGtGe94EiVZFq9RXdeqS1iJ27eHR323/nFFFq8q rn2omEm7TwCcDkuYxELMl8nbDVSBiu4hPanreEv0Dn1Plhu29d6wVppIkYsyeSiBrnML3hzi/r whsrsl6HzMepbqBrbXV/1TlRcdtNpRTh0G62rgmbbgqFGDx4Xk0bTRTtCfwUyHJiW5Pjogjtsn wGRUXy6V4KnX8d1FH+dCaAGnwZeU7Fg7gG7sjOOmCFPsh7J4BArB8IRmzyq5rEVPL2fLOR9d2y UY0= From: Julian Brown To: Subject: [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics Date: Tue, 16 Jun 2020 15:38:33 -0700 Message-ID: <5e9472b80dc475214a4a082ef54ee919d7f9dcff.1592343756.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a new version of the patch last sent here: https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546332.html Minus the bits that Thomas has committed already (thanks!), and with adjustments to allow for GOMP_MAP_ATTACH being grouped together with a preceding clause. OK? Julian ChangeLog libgomp/ * libgomp.h (struct splay_tree_key_s): Change virtual_refcount to dynamic_refcount. (struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA. * oacc-mem.c (acc_map_data): Substitute virtual_refcount for dynamic_refcount. (goacc_enter_datum): Adjust for dynamic_refcount semantics. (goacc_exit_datum): Re-add some error checking. Adjust for dynamic_refcount semantics. (goacc_enter_data_internal): Implement "present" case of dynamic memory-map handling here. Update "non-present" case for dynamic_refcount semantics. (goacc_exit_data_internal): Update for dynamic_refcount semantics. * target.c (gomp_map_vars_internal): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA handling. Update for dynamic_refcount handling. (gomp_unmap_vars_internal): Remove virtual_refcount handling. (gomp_load_image_to_device): Substitute dynamic_refcount for virtual_refcount. libgomp/ * testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs. --- libgomp/libgomp.h | 8 +- libgomp/oacc-mem.c | 155 ++++++++++++++---- libgomp/target.c | 38 +---- .../libgomp.oacc-c-c++-common/refcounting-1.c | 31 ++++ .../libgomp.oacc-c-c++-common/refcounting-2.c | 31 ++++ .../libgomp.oacc-fortran/deep-copy-6.f90 | 6 +- 6 files changed, 201 insertions(+), 68 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index ca42e0de640..7b52ce7d5c2 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1016,11 +1016,8 @@ struct splay_tree_key_s { uintptr_t tgt_offset; /* Reference count. */ uintptr_t refcount; - /* Reference counts beyond those that represent genuine references in the - linked splay tree key/target memory structures, e.g. for multiple OpenACC - "present increment" operations (via "acc enter data") referring to the same - host-memory block. */ - uintptr_t virtual_refcount; + /* Dynamic reference count. */ + uintptr_t dynamic_refcount; struct splay_tree_aux *aux; }; @@ -1153,7 +1150,6 @@ struct gomp_device_descr enum gomp_map_vars_kind { GOMP_MAP_VARS_OPENACC, - GOMP_MAP_VARS_OPENACC_ENTER_DATA, GOMP_MAP_VARS_TARGET, GOMP_MAP_VARS_DATA, GOMP_MAP_VARS_ENTER_DATA diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index be7f8d600eb..bc64bebe6c1 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -409,7 +409,7 @@ acc_map_data (void *h, void *d, size_t s) splay_tree_key n = tgt->list[0].key; assert (n); assert (n->refcount == 1); - assert (n->virtual_refcount == 0); + assert (n->dynamic_refcount == 0); /* Special reference counting behavior. */ n->refcount = REFCOUNT_INFINITY; @@ -456,7 +456,7 @@ acc_unmap_data (void *h) (void *) n->host_start, (int) host_size, (void *) h); } /* TODO This currently doesn't catch 'REFCOUNT_INFINITY' usage different from - 'acc_map_data'. Maybe 'virtual_refcount' can be used for disambiguating + 'acc_map_data'. Maybe 'dynamic_refcount' can be used for disambiguating the different 'REFCOUNT_INFINITY' cases, or simply separate 'REFCOUNT_INFINITY' values per different usage ('REFCOUNT_ACC_MAP_DATA' etc.)? */ @@ -545,10 +545,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) assert (n->refcount != REFCOUNT_LINK); if (n->refcount != REFCOUNT_INFINITY) - { - n->refcount++; - n->virtual_refcount++; - } + n->refcount++; + n->dynamic_refcount++; gomp_mutex_unlock (&acc_dev->lock); } @@ -562,13 +560,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) struct target_mem_desc *tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, - kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA); + kinds, true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); assert (tgt->list_count == 1); n = tgt->list[0].key; assert (n); assert (n->refcount == 1); - assert (n->virtual_refcount == 0); + assert (n->dynamic_refcount == 0); + n->dynamic_refcount++; d = (void *) tgt->tgt_start; } @@ -689,23 +688,28 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async) (void *) h, (int) s, (void *) n->host_start, (int) host_size); } + assert (n->refcount != REFCOUNT_LINK); + if (n->refcount != REFCOUNT_INFINITY + && n->refcount < n->dynamic_refcount) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("Dynamic reference counting assert fail\n"); + } + bool finalize = (kind == GOMP_MAP_DELETE || kind == GOMP_MAP_FORCE_FROM); if (finalize) { if (n->refcount != REFCOUNT_INFINITY) - n->refcount -= n->virtual_refcount; - n->virtual_refcount = 0; + n->refcount -= n->dynamic_refcount; + n->dynamic_refcount = 0; } - - if (n->virtual_refcount > 0) + else if (n->dynamic_refcount) { if (n->refcount != REFCOUNT_INFINITY) n->refcount--; - n->virtual_refcount--; + n->dynamic_refcount--; } - else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY) - n->refcount--; if (n->refcount == 0) { @@ -1048,13 +1052,111 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, { for (size_t i = 0; i < mapnum; i++) { - int group_last = find_group_last (i, mapnum, sizes, kinds); + splay_tree_key n; + size_t group_last = find_group_last (i, mapnum, sizes, kinds); + bool struct_p = false; + size_t size, groupnum = (group_last - i) + 1; + + switch (kinds[i] & 0xff) + { + case GOMP_MAP_STRUCT: + { + int last = i + sizes[i]; + size = (uintptr_t) hostaddrs[last] + sizes[last] + - (uintptr_t) hostaddrs[i]; + struct_p = true; + } + break; + + case GOMP_MAP_ATTACH: + size = sizeof (void *); + break; + + default: + size = sizes[i]; + } + + n = lookup_host (acc_dev, hostaddrs[i], size); + + if (n && struct_p) + { + if (n->refcount != REFCOUNT_INFINITY) + n->refcount += groupnum - 1; + n->dynamic_refcount += groupnum - 1; + gomp_mutex_unlock (&acc_dev->lock); + } + else if (n && groupnum == 1) + { + void *h = hostaddrs[i]; + size_t s = sizes[i]; + + /* A standalone attach clause. */ + if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, + (uintptr_t) h, s, NULL); + else if (h + s > (void *) n->host_end) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s); + } + + assert (n->refcount != REFCOUNT_LINK); + if (n->refcount != REFCOUNT_INFINITY) + n->refcount++; + n->dynamic_refcount++; - gomp_map_vars_async (acc_dev, aq, - (group_last - i) + 1, - &hostaddrs[i], NULL, - &sizes[i], &kinds[i], true, - GOMP_MAP_VARS_OPENACC_ENTER_DATA); + gomp_mutex_unlock (&acc_dev->lock); + } + else if (n && groupnum > 1) + { + assert (n->refcount != REFCOUNT_INFINITY + && n->refcount != REFCOUNT_LINK); + + for (size_t j = i + 1; j <= group_last; j++) + if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH) + { + splay_tree_key m + = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, + (uintptr_t) hostaddrs[j], sizes[j], NULL); + } + + bool processed = false; + + struct target_mem_desc *tgt = n->tgt; + for (size_t j = 0; j < tgt->list_count; j++) + if (tgt->list[j].key == n) + { + for (size_t k = 0; k < groupnum; k++) + if (j + k < tgt->list_count && tgt->list[j + k].key) + { + tgt->list[j + k].key->refcount++; + tgt->list[j + k].key->dynamic_refcount++; + } + processed = true; + } + + gomp_mutex_unlock (&acc_dev->lock); + if (!processed) + gomp_fatal ("dynamic refcount incrementing failed for " + "pointer/pset"); + } + else if (hostaddrs[i]) + { + gomp_mutex_unlock (&acc_dev->lock); + + struct target_mem_desc *tgt + = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, + &sizes[i], &kinds[i], true, + GOMP_MAP_VARS_ENTER_DATA); + assert (tgt); + for (size_t j = 0; j < tgt->list_count; j++) + { + n = tgt->list[j].key; + if (n) + n->dynamic_refcount++; + } + } i = group_last; } @@ -1148,18 +1250,15 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if (finalize) { if (n->refcount != REFCOUNT_INFINITY) - n->refcount -= n->virtual_refcount; - n->virtual_refcount = 0; + n->refcount -= n->dynamic_refcount; + n->dynamic_refcount = 0; } - - if (n->virtual_refcount > 0) + else if (n->dynamic_refcount) { if (n->refcount != REFCOUNT_INFINITY) n->refcount--; - n->virtual_refcount--; + n->dynamic_refcount--; } - else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY) - n->refcount--; if (copyfrom && (kind != GOMP_MAP_FROM || n->refcount == 0)) diff --git a/libgomp/target.c b/libgomp/target.c index 36425477dcb..3f2becdae0e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -666,8 +666,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; - tgt->refcount = (pragma_kind == GOMP_MAP_VARS_ENTER_DATA - || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1; + tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1; tgt->device_descr = devicep; tgt->prev = NULL; struct gomp_coalesce_buf cbuf, *cbufp = NULL; @@ -1094,7 +1093,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[i].copy_from = false; tgt->list[i].always_copy_from = false; tgt->list[i].do_detach - = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA); + = (pragma_kind != GOMP_MAP_VARS_ENTER_DATA); n->refcount++; } else @@ -1155,7 +1154,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[i].offset = 0; tgt->list[i].length = k->host_end - k->host_start; k->refcount = 1; - k->virtual_refcount = 0; + k->dynamic_refcount = 0; tgt->refcount++; array->left = NULL; array->right = NULL; @@ -1294,20 +1293,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, /* If the variable from "omp target enter data" map-list was already mapped, tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or gomp_exit_data. */ - if ((pragma_kind == GOMP_MAP_VARS_ENTER_DATA - || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) - && tgt->refcount == 0) - { - /* If we're about to discard a target_mem_desc with no "structural" - references (tgt->refcount == 0), any splay keys linked in the tgt's - list must have their virtual refcount incremented to represent that - "lost" reference in order to implement the semantics of the OpenACC - "present increment" operation properly. */ - if (pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) - for (i = 0; i < tgt->list_count; i++) - if (tgt->list[i].key) - tgt->list[i].key->virtual_refcount++; - + if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0) + { free (tgt); tgt = NULL; } @@ -1459,14 +1446,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, continue; bool do_unmap = false; - if (k->tgt == tgt - && k->virtual_refcount > 0 - && k->refcount != REFCOUNT_INFINITY) - { - k->virtual_refcount--; - k->refcount--; - } - else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) + if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) k->refcount--; else if (k->refcount == 1) { @@ -1631,7 +1611,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt = tgt; k->tgt_offset = target_table[i].start; k->refcount = REFCOUNT_INFINITY; - k->virtual_refcount = 0; + k->dynamic_refcount = 0; k->aux = NULL; array->left = NULL; array->right = NULL; @@ -1665,7 +1645,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt = tgt; k->tgt_offset = target_var->start; k->refcount = is_link_var ? REFCOUNT_LINK : REFCOUNT_INFINITY; - k->virtual_refcount = 0; + k->dynamic_refcount = 0; k->aux = NULL; array->left = NULL; array->right = NULL; @@ -2935,7 +2915,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr, k->tgt = tgt; k->tgt_offset = (uintptr_t) device_ptr + device_offset; k->refcount = REFCOUNT_INFINITY; - k->virtual_refcount = 0; + k->dynamic_refcount = 0; k->aux = NULL; array->left = NULL; array->right = NULL; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c new file mode 100644 index 00000000000..4e6d06d48d5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c @@ -0,0 +1,31 @@ +/* Test dynamic unmapping of separate structure members. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include + +struct s +{ + char a; + char b; +}; + +int main () +{ + struct s s; + +#pragma acc enter data create(s.a, s.b) + + assert (acc_is_present (&s.a, sizeof s.a)); + assert (acc_is_present (&s.b, sizeof s.b)); + +#pragma acc exit data delete(s.a) +#pragma acc exit data delete(s.b) + + assert (!acc_is_present (&s.a, sizeof s.a)); + assert (!acc_is_present (&s.b, sizeof s.b)); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c new file mode 100644 index 00000000000..5539fd8d57f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c @@ -0,0 +1,31 @@ +/* Test dynamic unmapping of separate structure members. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include + +struct s +{ + char a; + char b; +}; + +int main () +{ + struct s s; + +#pragma acc enter data create(s.a, s.b) + + assert (acc_is_present (&s.a, sizeof s.a)); + assert (acc_is_present (&s.b, sizeof s.b)); + + acc_delete (&s.a, sizeof s.a); + acc_delete (&s.b, sizeof s.b); + + assert (!acc_is_present (&s.a, sizeof s.a)); + assert (!acc_is_present (&s.b, sizeof s.b)); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 index 5837a403910..eb7d3ca160e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 @@ -43,12 +43,8 @@ program dtype print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_get_device_type() .ne. acc_device_host) then if (acc_is_present(var%a(5:n - 5))) stop 21 if (acc_is_present(var%b(5:n - 5))) stop 22 From patchwork Tue Jun 16 22:38:34 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310762 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjn13dMvz9sSc for ; Wed, 17 Jun 2020 08:39:33 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7E3E6388F070; Tue, 16 Jun 2020 22:39:07 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 7E50D388B011; Tue, 16 Jun 2020 22:39:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 7E50D388B011 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: OxafrI9wppGmaEawav/AARkE5S/SvS9lp6a3LTirt/HMfGa3z1cBmthC5CqT6rZUJegmPvHQ5U yDwLMNwaqHr0DsNSSk4cTdMmZWAcylfYBwQ4LKbYR3kWBfHacawcl06z6kkseiLiIzRVSB2la2 sMERR6ESLL7ClK8HCQ5e5Fhs9D6y83nfG/6R0mBcyQFhMpVUR4VmX2Aaqyf/TAfKFYp/aLUTAv J0usoDLcCsPNh7geI8fBJuF0OWK40X9MdtgFDYGh299xbK4wo99RsW9W7VEDslYg34Rl/8h+9q XAw= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="50003703" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:39:03 -0800 IronPort-SDR: 4IsSV9IqOwvwoGa08qEWugJ19BGniKZC1svlQ0tmwDv1kcCprtj5BvD5Jgi/xa5UVsucWKjhMI Le4135qID/X4YWTUbsODPI4v74sSJ5P9jZ6ngW/FG1KJT8zPcCIhPjRK/oSe0EDBL3Z0w6r6kT il2ZBNjkgmgKX5L3ak56yS6ZZeNBHDJx8iOhG00Jjk59JXCc4263YGTFhCyhCE3dm9J2KwTW8U Ty8TrSfwjHtmJTTaWHEhas3bJ0GEOv3YUBAsIE1S5rxc0dsSi3STECyLuVRc9ADXZbuDF0aXCp 3Y0= From: Julian Brown To: Subject: [PATCH 4/9] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum Date: Tue, 16 Jun 2020 15:38:34 -0700 Message-ID: <57f10eb6f6449cb2b1676ef839c9b667e028f8de.1592343756.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a minor cleanup for goacc_enter_datum. Unchanged from previous posting, but including for completeness: https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546331.html OK? Julian ChangeLog libgomp/ * oacc-mem.c (goacc_enter_datum): Use scalar kind argument instead of kinds array. (acc_create, acc_create_async, acc_copyin, acc_copyin_async): Update calls to goacc_enter_datum. --- libgomp/oacc-mem.c | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index bc64bebe6c1..05998ebc6de 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -501,7 +501,8 @@ acc_unmap_data (void *h) /* Enter dynamic mapping for a single datum. Return the device pointer. */ static void * -goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) +goacc_enter_datum (void **hostaddrs, size_t *sizes, unsigned short kind, + int async) { void *d; splay_tree_key n; @@ -560,7 +561,7 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) struct target_mem_desc *tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, - kinds, true, GOMP_MAP_VARS_ENTER_DATA); + &kind, true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); assert (tgt->list_count == 1); n = tgt->list[0].key; @@ -584,15 +585,13 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) void * acc_create (void *h, size_t s) { - unsigned short kinds[1] = { GOMP_MAP_ALLOC }; - return goacc_enter_datum (&h, &s, &kinds, acc_async_sync); + return goacc_enter_datum (&h, &s, GOMP_MAP_ALLOC, acc_async_sync); } void acc_create_async (void *h, size_t s, int async) { - unsigned short kinds[1] = { GOMP_MAP_ALLOC }; - goacc_enter_datum (&h, &s, &kinds, async); + goacc_enter_datum (&h, &s, GOMP_MAP_ALLOC, async); } /* acc_present_or_create used to be what acc_create is now. */ @@ -617,15 +616,13 @@ acc_pcreate (void *h, size_t s) void * acc_copyin (void *h, size_t s) { - unsigned short kinds[1] = { GOMP_MAP_TO }; - return goacc_enter_datum (&h, &s, &kinds, acc_async_sync); + return goacc_enter_datum (&h, &s, GOMP_MAP_TO, acc_async_sync); } void acc_copyin_async (void *h, size_t s, int async) { - unsigned short kinds[1] = { GOMP_MAP_TO }; - goacc_enter_datum (&h, &s, &kinds, async); + goacc_enter_datum (&h, &s, GOMP_MAP_TO, async); } /* acc_present_or_copyin used to be what acc_copyin is now. */ From patchwork Tue Jun 16 22:39:41 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310763 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjnZ2Pz9z9sSc for ; Wed, 17 Jun 2020 08:40:02 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E4E973885C03; Tue, 16 Jun 2020 22:39:59 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id CADEB386F447; Tue, 16 Jun 2020 22:39:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org CADEB386F447 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: PriPuIoh68cS4J2Lk3q93m69QsIFJAeH1bRCvrTjwP43M+bxQ0sOs4HZrW0V00CDahVcHjNsBF YFZOkvLX3zvjiWW/868CETvsRYE+B2yBrWm+BTu9ZvJABeJpIXU/hU8gPSm/Kq0HDhwl+DxhJ/ 6momurw7XGir7HUz/hrlZDaDIu+NLOa5JdqPZzCXw29BSoehJ+O5orGx9jP0TMyxETv5xr+OYN Z/j5Pym0S775So05+xM6PtexC+jBXp9Xxy7NdQGTz2dcSpEHaYVSHGenY2swyME41opwi5gbpB oB0= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="49874498" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:39:56 -0800 IronPort-SDR: Ra4Qs9doJR09E3nEUqeUJn0D/scO9QnUlZZu1LuZgrw/Io6iEPgsaxoRtaSWdI0s9fpgsZ3XVe hY6v9y/MM7/lLzXELK6IajSXbuJkHSq0aFQ0hDYZAd5YLJ9GhgoE9xVgCAuWkbSu9yYddzs8Wd SbEO91gBlAPxrFPDx/nFJc5H8koFjLa69FSZOihFTJ2e5s+IcA08Ijw2f5xKdvADt6OfD+u7Kh c9ZwYKPJxBsooIjM4dOVNcCubFfNPoQDWlq/7ji/Y9urCfbe/IfDwNa9VqkAHICjHltNtqKoKw EkA= From: Julian Brown To: Subject: [PATCH 5/9] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843) Date: Tue, 16 Jun 2020 15:39:41 -0700 Message-ID: <0585b4fda1ba5c76dce2ac5053a55e2ceef06041.1592343756.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a repost of the following: https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546333.html I made a minor edit to the ChangeLog, but the patch is otherwise unchanged. Including for completeness. OK? ChangeLog PR libgomp/92843 libgomp/ * oacc-mem.c (goacc_exit_data_internal): Don't copyout data mapped with acc_map_data in exit data directive. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL. --- libgomp/oacc-mem.c | 1 + libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 1 - 2 files changed, 1 insertion(+), 1 deletion(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 05998ebc6de..745cb132621 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1258,6 +1258,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, } if (copyfrom + && n->refcount != REFCOUNT_INFINITY && (kind != GOMP_MAP_FROM || n->refcount == 0)) gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start, (void *) (n->tgt->tgt_start + n->tgt_offset diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c index f16c46a37bf..db5b35b08d9 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c @@ -1,7 +1,6 @@ /* Verify that 'acc_copyout' etc. is a no-op if there's still a structured reference count. */ -/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ #include From patchwork Tue Jun 16 22:39:42 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310764 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjnh0PGsz9sSc for ; Wed, 17 Jun 2020 08:40:08 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 951ED388E825; Tue, 16 Jun 2020 22:40:02 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id AFFAA388B03F; Tue, 16 Jun 2020 22:40:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AFFAA388B03F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: 20toNzHAubzrh12r7KT5pMaMfqd4o3c5D6CwiWJZIyaxDb1Zwm6yYHfGPQ2rlr0VfQKX9k/wQA 1LmrhWrBNPvhLQcJHe9zcNg1kfpkseViMlcMwp79RgH0oJeT7AGASw6AVuDzhkq1OxN/k2Lo8E Yofj39Zoh05bmj0Io8Z0JvfecpMvj9B3esdHUU3SqV4BOqwTL97zvQx8wGzgTeS0iP/JQdMF8l 58jFKTWTOtdcddt+B10XBbJ07vngfdChMpMx+UrkEQfOVyFvBHl5qpHANaRFXVY2x2h1csp+A5 Wh4= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="49874500" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:40:00 -0800 IronPort-SDR: opJVQDqrswTw5iKixg16yBhRxhrNNgYvCk3dJvuseY6E1PcFS2GrN3LW9mNVV0hn0V71ZdACXw bbToZrFIIZdRvs8uAa71rcx+5/q8kLAk5YS1wzG86AEXAn4ybOGBZREtmLbCn3SfN0wMaEjQiE t+VpScnzP28GxeECME6dqU/gF722/JjJdnLYP4KZBNSNNYCMngspAt3gcpe4vN+yYn6OzVgtp0 6mO0Hig+GQIMi5fxQselpcLB/3Abc6l90uFbSWBZx1FYxL1n2QuCUFfGJ++S0k3dwqeEbDO0rK xCc= From: Julian Brown To: Subject: [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++ Date: Tue, 16 Jun 2020 15:39:42 -0700 Message-ID: X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-2.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, UNWANTED_LANGUAGE_BODY autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a fix for the pointer (or array) size inadvertently being used for the bias of attach and detach clauses (PR95270), for C and C++. OK? Julian ChangeLog PR middle-end/95270 gcc/c/ * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero for standalone attach/detach clauses. gcc/cp/ * semantics.c (finish_omp_clauses): Likewise. gcc/testsuite/ * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero bias. --- gcc/c/c-typeck.c | 8 ++++++++ gcc/cp/semantics.c | 8 ++++++++ gcc/testsuite/c-c++-common/goacc/mdc-1.c | 14 +++++++------- 3 files changed, 23 insertions(+), 7 deletions(-) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 385bf3a1c7b..134f1520239 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -14533,6 +14533,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (c_oacc_check_attachments (c)) remove = true; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + OMP_CLAUSE_SIZE (c) = size_zero_node; break; } if (t == error_mark_node) @@ -14546,6 +14550,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + OMP_CLAUSE_SIZE (c) = size_zero_node; if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 64587c791c6..77e6ff7fb0d 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7334,6 +7334,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (cp_oacc_check_attachments (c)) remove = true; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + OMP_CLAUSE_SIZE (c) = size_zero_node; break; } if (t == error_mark_node) @@ -7347,6 +7351,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + OMP_CLAUSE_SIZE (c) = size_zero_node; if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index fb5841a709d..337c1f7cc77 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -45,12 +45,12 @@ t1 () /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */ From patchwork Tue Jun 16 22:39:43 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310765 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjnr6JR1z9sSc for ; Wed, 17 Jun 2020 08:40:16 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9D8E8388F064; Tue, 16 Jun 2020 22:40:07 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 08438388B03F; Tue, 16 Jun 2020 22:40:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 08438388B03F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: exLq0CjM2nuNMt9R0n9fuxveKWHNBlFu+rYkF49XM8JZbKEHKC3JjuKoGYPa+5jpYX9WcovzoG IVd4640s6immmI/Xu6+HCVUuJ0LuWDEwKkuMaH0b0qlrOoyYx2MI7ybP+0PJx2+JCdQjRHLxPD PoS+Gu6lus1Z6Qsy3xGXO3hIpRlljrrSg2F9vkpuU7ZjPdjAGeV0Rh0yQLYa7TDU4P0iRfzyeq MmmkuFMR4RpRQUan0WZtl5Vay+oTP85WS0zE4ICwXEc3fEIspWuAOIOAd4kqeHbHxnS+dUKX0F wZY= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="49874502" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:40:03 -0800 IronPort-SDR: YMwmH6AarcBO7gXtVPjnr75G0AZFQ1NMkx/IMY9H7oXVILhhxmT3JjhHZDSe5elIgnQGeZA7N5 z+HJ3FHtzw4Q9Nn9YEhxvMEzUoXzi8B1v3Uq6NS/MiVgvVxCaNX/jf/sLv5oHy/UioNPpkYyzk Tw1EqbRU0hnOU8sJqK2mOrxxl8nHy5xmauiP3KRp/21nl4x/oThLp9YHtFRGa6Ej65GFx9euUm a+TMRVTPqTo+ce8JCW5HBlKvRUXHf9V09DIxCkKcA1XoAu2MW2M/ss3aWCFe89JNH1SDno6DBf 5GA= From: Julian Brown To: Subject: [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives Date: Tue, 16 Jun 2020 15:39:43 -0700 Message-ID: X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-2.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, UNWANTED_LANGUAGE_BODY autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" When attaching pointers in Fortran, OpenACC 2.6 specifies that a descriptor must be copied to the target at the same time (see next patch). That means that stripping GOMP_MAP_TO_PSET (and lesserly, GOMP_MAP_POINTER), which was behaviour introduced by the manual deep-copy middle-end support patch, was probably wrong. That arguably answers some of the questions at the end of: https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547424.html It appears that the user can (but certainly should not!) map a synthesized array descriptor using an "enter data" operation that can go out of scope before that data is unmapped. It would be nice to give a warning for an attempt to do such a thing, though I have no idea if that's possible in practice. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data directives. gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update expected dump output. --- gcc/gimplify.c | 11 ++--------- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++-- 2 files changed, 4 insertions(+), 11 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 9851edfc4db..aa6853f0dcc 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8767,6 +8767,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -8775,15 +8777,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, mapped, but not the pointer to it. */ remove = true; break; - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - remove = true; - break; default: break; } diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 1e2e3e94b8a..ca642156e9f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -21,7 +21,7 @@ !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5)) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } @@ -33,5 +33,5 @@ !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } END SUBROUTINE f From patchwork Tue Jun 16 22:39:44 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310766 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjny6V0nz9sSc for ; Wed, 17 Jun 2020 08:40:22 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 0B5E5388F077; Tue, 16 Jun 2020 22:40:10 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 36221383E82C; Tue, 16 Jun 2020 22:40:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 36221383E82C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: 0K1auYXV4EseN8Q6PJ4iKHFSwJ7WLAr/4mXLp2xxng1K0CFE6z3eIn+b0aZri3+5n5T22DAX00 B3nhwkHCd1F5IjB0RpVqpvHnR/3mwyZKAGkADAqxIqAuBpOqL36sXKHQ2AuE86qE7R5ExKa+28 ZXQqP7JiiQMPI5nLD++zPQ07siVROC4P1vgcPB0m2kZKGjJYT35qMRd70klDREqOXbjFbQNVdp nneYKM1FUtKytRKDBzn0DTCImgEOYyKFpcqT8OXBaUlPAZ0UjFF0jdCAD84iXu8yCOrtZKH8Nx XZM= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="49874503" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:40:06 -0800 IronPort-SDR: AZo0T8b9HJ2aUHXnp4fprZZfMttmvkXwh0rVmcxXiorw9IBZ+Xot6Bky05/Rpqn+fGLg5Cp269 Th1tHql828rP4Cuyz9nTfToXBoovnk5gFwuVfUMJcoraMCWWGlIKFV9rUqX1tkjOcmUSCRlfUe Q8RgV8oAClrK4Nyt0EMrZm4Q+6ZBpbvMTiyVY1/APxbjAtQq1A54boLrEUEKXIINpUkAl+WQFS pKhyiatmh9RMgDMHD+xbdlvCbismFxyYQrkHn2YASYwefyI0StHtHyaisd298BrAiJycHlwzp8 V7M= From: Julian Brown To: Subject: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Date: Tue, 16 Jun 2020 15:39:44 -0700 Message-ID: <0193ff08d4a4a2b6ca86c7a891b8ff35203fa440.1592343757.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.6 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.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" As mentioned in the blurb for the previous patch, an "attach" operation for a Fortran pointer with an array descriptor must copy that array descriptor to the target. This patch arranges for that to be so. OK? Julian ChangeLog gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Copy array descriptor to target for attach clauses when appropriate. libgomp/ * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: New test. --- gcc/fortran/trans-openmp.c | 40 ++++++++++++++- .../attach-descriptor-1.f90 | 51 +++++++++++++++++++ 2 files changed, 89 insertions(+), 2 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 02c40fdc660..909a86795e0 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } } if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)) - && n->u.map_op != OMP_MAP_ATTACH - && n->u.map_op != OMP_MAP_DETACH) + && (n->u.map_op == OMP_MAP_ATTACH + || n->u.map_op == OMP_MAP_DETACH)) + { + tree type = TREE_TYPE (decl); + tree data = gfc_conv_descriptor_data_get (decl); + if (present) + data = gfc_build_cond_assign_expr (block, present, + data, + null_pointer_node); + tree ptr + = fold_convert (build_pointer_type (char_type_node), + data); + ptr = build_fold_indirect_ref (ptr); + /* Standalone attach clauses used with arrays with + descriptors must copy the descriptor to the target, + else they won't have anything to perform the + attachment onto (see OpenACC 2.6, "2.6.3. Data + Structures with Pointers"). */ + OMP_CLAUSE_DECL (node) = ptr; + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (node2) = decl; + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); + node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); + if (n->u.map_op == OMP_MAP_ATTACH) + { + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); + n->u.map_op = OMP_MAP_ALLOC; + } + else /* OMP_MAP_DETACH. */ + { + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); + n->u.map_op = OMP_MAP_RELEASE; + } + OMP_CLAUSE_DECL (node3) = data; + OMP_CLAUSE_SIZE (node3) = size_int (0); + } + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) { tree type = TREE_TYPE (decl); tree ptr = gfc_conv_descriptor_data_get (decl); diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 new file mode 100644 index 00000000000..2dd1a6fa5b6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -0,0 +1,51 @@ +program att + use openacc + implicit none + type t + integer :: arr1(10) + integer, allocatable :: arr2(:) + end type t + integer :: i + type(t) :: myvar + integer, target :: tarr(10) + integer, pointer :: myptr(:) + + allocate(myvar%arr2(10)) + + do i=1,10 + myvar%arr1(i) = 0 + myvar%arr2(i) = 0 + tarr(i) = 0 + end do + + call acc_copyin(myvar) + call acc_copyin(myvar%arr2) + call acc_copyin(tarr) + + myptr => tarr + + !$acc enter data attach(myvar%arr2, myptr) + + ! FIXME: This warning is emitted on the wrong line number. + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 } + !$acc serial present(myvar%arr2) + do i=1,10 + myvar%arr1(i) = i + myvar%arr2(i) = i + end do + myptr(3) = 99 + !$acc end serial + + !$acc exit data detach(myvar%arr2, myptr) + + call acc_copyout(myvar%arr2) + call acc_copyout(myvar) + call acc_copyout(tarr) + + do i=1,10 + if (myvar%arr1(i) .ne. i) stop 1 + if (myvar%arr2(i) .ne. i) stop 2 + end do + if (tarr(3) .ne. 99) stop 3 + +end program att From patchwork Tue Jun 16 22:39:45 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1310767 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49mjp66JxDz9sSd for ; Wed, 17 Jun 2020 08:40:30 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7D2913890405; Tue, 16 Jun 2020 22:40:14 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id BB5773890401; Tue, 16 Jun 2020 22:40:10 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org BB5773890401 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: PfWgj2M+WGbAKJmehD/5SS2rhqWDnDdLcZnR9/8a/otSY2t4p7GbSpLE2Ykblw2sZaySqDQxYl INK3aU/s1xjMe1uEA5TMnQsvbp7AwucmmeFR+YNaLIQgf4k+C3vvSTDkV87ClWS5QJd/xVz2Pv XlCa0sHy5hDxvdBAZ0S0ng+VeAzlOjoDojEo3LMhWMAh51TTs4jtzwsU6JlrLkVAttjBA8AIMg 5AxwPB/5kSMbmSrGUHtCbaKi7WuYluDH3kh+gZlOTYHY3xput94o88l/yYB9oGl8BcI0pNvZl4 7r4= X-IronPort-AV: E=Sophos;i="5.73,520,1583222400"; d="scan'208";a="49874506" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 16 Jun 2020 14:40:10 -0800 IronPort-SDR: i9AjmT7xqF8zXP0rtNvEqurPMUMA+plm9Uav64c4iUbQEYYOxR/7t6t/dpB4oHv9i+lU+obA5k 8TKtLtU4noDzA7vabph4w1UDFnC9WUXAGF9sJaDX+FcZQcpwKXP40W/Kpt9YNr9qkPATFPm/+f 3enbxV8JhxUOG3X4SoxnlPMvXkWtJNS3Te3qQQPjOBNq4r9nQPlLZcxAYoqbcPkQmwR6iC2jbu +on+zm1ffAvhRzTMPpIElPH48nwayYPxAQT1SOHyPMXBHScYIeZdoEaJsqoXszv+oz4zWNyQaL uiA= From: Julian Brown To: Subject: [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount Date: Tue, 16 Jun 2020 15:39:45 -0700 Message-ID: <5309b57f3707783cf65c2b6c5b9b784d8e61b760.1592343757.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.7 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.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Jakub Jelinek , Tobias Burnus , "Moore, Catherine" , Thomas Schwinge , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch fixes a set of XFAILs in some recently-added patches by skipping a detach operation on "no-op" exit data operations for blocks with zero dynamic refcount. This takes advantage of the ordering of detach clauses with respect to associated data-movement clauses: i.e., they are grouped together adjacently. OK? Julian ChangeLog libgomp/ * oacc-mem.c (find_group_last): Handle detach operations. (goacc_exit_data_internal): Detect detachments that are part of copyout operations, and suppress them if dynamic refcount is zero. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Fix typo. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Remove XFAILs. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90: Likewise. --- libgomp/oacc-mem.c | 54 ++++++++++++++++--- .../mdc-refcount-1-1-1.f90 | 6 +-- .../mdc-refcount-1-1-2.F90 | 2 +- .../mdc-refcount-1-2-1.f90 | 6 +-- .../mdc-refcount-1-2-2.f90 | 6 +-- .../mdc-refcount-1-3-1.f90 | 6 +-- .../mdc-refcount-1-3-2.f90 | 5 +- .../mdc-refcount-1-4-1.f90 | 6 +-- .../mdc-refcount-1-4-2.f90 | 5 +- 9 files changed, 55 insertions(+), 41 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 745cb132621..f852652c048 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -987,7 +987,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) { case GOMP_MAP_TO_PSET: if (pos + 1 < mapnum - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH) + && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH + || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH + || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH)) return pos + 1; while (pos + 1 < mapnum @@ -1010,6 +1012,8 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) break; case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: return pos; default: @@ -1025,7 +1029,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) /* We can have a single GOMP_MAP_ATTACH mapping after a to/from mapping. */ if (pos + 1 < mapnum - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH) + && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH + || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH + || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH)) return pos + 1; /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from @@ -1168,15 +1174,43 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, { gomp_mutex_lock (&acc_dev->lock); - /* Handle "detach" before copyback/deletion of mapped data. */ - for (size_t i = 0; i < mapnum; ++i) + /* Handle "detach" before copyback/deletion of mapped data. If this isn't a + standalone "detach" clause, take care to skip the "detach" operation if + the dynamic refcount of the data to be detached is zero. */ + for (size_t grp = 0; grp < mapnum; grp++) { - unsigned char kind = kinds[i] & 0xff; + size_t i = grp, group_last = find_group_last (grp, mapnum, sizes, kinds); + unsigned char kind = kinds[grp] & 0xff; bool finalize = false; + switch (kind) { + case GOMP_MAP_TO_PSET: + case GOMP_MAP_TOFROM: + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + { + if (i + 1 >= mapnum) + break; + kind = kinds[i + 1] & 0xff; + if (kind != GOMP_MAP_FORCE_DETACH && kind != GOMP_MAP_DETACH) + break; + splay_tree_key n = lookup_host (acc_dev, hostaddrs[i], sizes[i]); + if (n == NULL) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("target data not mapped for detach operation"); + } + i++; + if (n->dynamic_refcount == 0) + break; + } + /* Fallthrough. */ + case GOMP_MAP_FORCE_DETACH: - finalize = true; + finalize = (kind == GOMP_MAP_FORCE_DETACH); /* Fallthrough. */ case GOMP_MAP_DETACH: @@ -1197,9 +1231,15 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_detach_pointer (acc_dev, aq, n, hostaddr, finalize, NULL); } break; + case GOMP_MAP_STRUCT: + case GOMP_MAP_POINTER: + /* Ignore. */ + break; default: - ; + gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x", + kind); } + grp = group_last; } for (size_t i = 0; i < mapnum; ++i) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 index 445cbabb8ca..7171affb9f0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 @@ -24,12 +24,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) finalize - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 index 7b206ac2042..2aa46189e9a 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 @@ -6,4 +6,4 @@ #include "mdc-refcount-1-1-1.f90" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } -! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" } +! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 index 8554534b2f2..9a10aa5a781 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 @@ -26,12 +26,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) finalize - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 index 8e696cc70e8..f506adf8e91 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 @@ -26,12 +26,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 index 070a6f8e149..450d95d3686 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 @@ -27,12 +27,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) finalize - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90 index 3c4bbda7f66..35efad4138a 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90 @@ -27,11 +27,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) - !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 index b22e411567f..816562fc055 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 @@ -26,12 +26,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data detach(var%a) finalize - !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed. - !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } !$acc exit data delete(var%a) if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90 index 476cd5c1bee..b98bfd74924 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90 @@ -27,11 +27,8 @@ program main print *, "CheCKpOInT1" ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } !$acc exit data delete(var%a) - !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing). - !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. - !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. print *, "CheCKpOInT2" - ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } if (acc_is_present(var%a)) stop 3 if (.not. acc_is_present(var)) stop 4