From patchwork Fri Jan 17 21:18:20 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1224976 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=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-517619-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha1 header.s=default header.b=uPmgVdvG; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 47zv805HdMz9sR1 for ; Sat, 18 Jan 2020 08:19:12 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-transfer-encoding:content-type; q=dns; s= default; b=QefERj62EYpO2oH7WBzjdK61nc/H5P8jee8M+vxC0W2/uvPTnMw1D SbyeKC6GyrOzG2Wlfnry0q/KoyQgveZxsFQB22EWYO5MzeeVRX0w/htPF+Ev6kL6 mu3mRBB6NbwvKEhBsq5Q6xRu3qMNVppyhOUSX3RgIMAq3jGT/EAx3A= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-transfer-encoding:content-type; s=default; bh=rBpRqqGPmWZFEX+t1xAjYuN4A9I=; b=uPmgVdvGsdEvt5rrdxd6kau4U1Lm sXW/BYkjk0Q7zj21Y+Kly8namOlRvGcFrliGKMTSTWjhLclNtwNeYymZt/5+2FWe rOC5sq1F9y9ZSI5UOYAUB34JrNVGxqBOF0Z8Juh0qdC1v1KZcAi7HX53QS/m7HC4 z4g+52W4FvF1asM= Received: (qmail 128402 invoked by alias); 17 Jan 2020 21:18:53 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 128343 invoked by uid 89); 17 Jan 2020 21:18:48 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.3 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3 autolearn=ham version=3.3.1 spammy= X-HELO: esa4.mentor.iphmx.com Received: from esa4.mentor.iphmx.com (HELO esa4.mentor.iphmx.com) (68.232.137.252) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 17 Jan 2020 21:18:46 +0000 IronPort-SDR: IPiH4cQ0gm3JkdOjf4vq1qxlY/MZvc4r8EfCxnAi+z0RkUYK7E3LE2ga3uVW2StfD+e8cUbOZv JGo+UyOTJICemvE6GRFO9JXq0lrNeGxD0qNb37P2j/oxh4uuFZ0R+afc3ZfdWnwMJmKJQpmbID yJzlHFpgFCD4rLKxRdb0AyC2TVOjeP7bdjVqtgqv3TVsbNwRan2d4bOnaBIbncdINChZEQJQCH OoKiRDYUtLI2fDfCNybGHE1/6rh6zxiYzki/UpO09pq1nlrLgpBdn130dndKsqaTunA/pq8uaz JVU= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 17 Jan 2020 13:18:45 -0800 IronPort-SDR: swdSgnJLouro1gZhixSoV9wFZVgF/jcZ4CiCN1rXHNu/1dL1OoQ2dN0aLewou9+HE8sEdKlXGO v3U/QJ+swZNA== From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek Subject: [PATCH 2/3] Don't copy back vars mapped with acc_map_data Date: Fri, 17 Jan 2020 13:18:20 -0800 Message-ID: <4bbb4970cd48424873573a50e627786cc9cf3378.1579292772.git.julian@codesourcery.com> In-Reply-To: References: MIME-Version: 1.0 X-IsSubscribed: yes This patch prevents "exit data" directives from copying back data that was mapped with an acc_map_data API call. This matches the behaviour expected by the pr92843-1.c test, and together with the previous patch in this series, allows that test to pass (with no other regressions). Tested alongside other patches in this series with offloading to NVPTX (with and without the third & final patch). OK? Thanks, Julian ChangeLog PR libgomp/92843 libgomp/ * oacc-mem.c (goacc_exit_data_internal): Don't copy-back data mapped with acc_map_data on an "exit data" directive. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL. Add explanatory comment. --- libgomp/oacc-mem.c | 1 + libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 4 +++- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 45ab2b169d7..783e7f363fb 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, n->refcount--; 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..786a12a8504 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 @@ -96,6 +95,9 @@ test_acc_map_data () verify_array (h, N, c1); assign_array (h, N, c1); + /* Note that we're not expecting this (nor the copyouts below) to perform + an actual "finalize" or copyout since the data was mapped with + acc_map_data. */ #pragma acc exit data copyout (h[0:N]) finalize assert (acc_is_present (h, N)); verify_array (h, N, c1);