From patchwork Thu Jun 4 13:40:53 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1303552 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 49d6PY4G8Jz9sSy for ; Thu, 4 Jun 2020 23:41:21 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 712F13890409; Thu, 4 Jun 2020 13:41:17 +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 E4035388F077 for ; Thu, 4 Jun 2020 13:41:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org E4035388F077 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: xBXsn8Su09l7YfEzc4zrEW+2buI8Vbhs8SlAQGUCf4QMcZNutWuyN/K8Mu/OyYDwvI/4prWznN yFwM3hctkmwyyiM9GCb5eDKeY+/GPdEH9N1LZzPj+qwb/r8QYLHvYdayyR7SSubBbyVP/qo8iV oaYRzGOaf4LIBKSLWzKIjfMG7EbyidnEYpvP+jJj45FzRl5l2bc4/MSUBa4UpaWEacGiNznPb8 PhD6lgIqBK/FVtWRaCf4Lc+4zwYM8nE29JbU9zqgb9VqTB8BPl028+P9O3D7P8pbyRk1fDLjkt AWU= X-IronPort-AV: E=Sophos;i="5.73,472,1583222400"; d="scan'208";a="49505642" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 04 Jun 2020 05:41:13 -0800 IronPort-SDR: 8BNDatRKd856KUBolS1cR9F+Os/66GujekGzUwPuxGVpDmOqRNjwEqA8sv/fCZTn9HA2HSvUyH D1K3laDhsXREiUpETOPzXzFeaWhPq+cW0Q2qSz19iIHOdDlHeKOgGlXATaVJ1QDxQ50rS/nYoP xeA2GKwtnvfejAKq5DrFiCyfY9NuzDxtys5t6NB+qu3Ote3NKNGsH+Mq/psK7o3SZT69vV0uz+ HgNvVKfnyskOIzzj8aAFcyyjJ6eRIXMghti43lWNlpWG+v1iSUzhQTI0noo/AA3fRRQKL37umQ L2w= From: Julian Brown To: Subject: [PATCH 1/3] OpenACC "exit data" copyout for struct members Date: Thu, 4 Jun 2020 06:40:53 -0700 Message-ID: <6538c388d22e016bd01737e8c2f80437f432c70f.1591276990.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: <87ftbw9kqh.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) 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_NONE, 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: Catherine_Moore@mentor.com, Tobias Burnus , Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch removes unnecessary special-case code in oacc-mem.c which prevented copy-back of structure members on OpenACC "exit data" directives. I've added a couple of new testcases to verify the behaviour (which fail without the patch). (On editing this mail, I notice I omitted to add a comment about GOMP_MAP_STRUCT being a no-op, and/or no longer emitted for OpenACC exit data with the patch later in this series. I'll add such a comment before commit.) OK? Thanks, Julian libgomp/ * oacc-mem.c (goacc_exit_data_internal): Remove special-case (no-copyback) handling for GOMP_MAP_STRUCT. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. --- libgomp/oacc-mem.c | 32 -------------- .../struct-copyout-1.c | 38 ++++++++++++++++ .../struct-copyout-2.c | 44 +++++++++++++++++++ 3 files changed, 82 insertions(+), 32 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2d4bba78efd..232683a85f0 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, break; case GOMP_MAP_STRUCT: - { - int elems = sizes[i]; - for (int j = 1; j <= elems; j++) - { - struct splay_tree_key_s k; - k.host_start = (uintptr_t) hostaddrs[i + j]; - k.host_end = k.host_start + sizes[i + j]; - splay_tree_key str; - str = splay_tree_lookup (&acc_dev->mem_map, &k); - if (str) - { - if (finalize) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount -= str->virtual_refcount; - str->virtual_refcount = 0; - } - if (str->virtual_refcount > 0) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount--; - str->virtual_refcount--; - } - else if (str->refcount > 0 - && str->refcount != REFCOUNT_INFINITY) - str->refcount--; - if (str->refcount == 0) - gomp_remove_var_async (acc_dev, str, aq); - } - } - i += elems; - } break; default: diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c new file mode 100644 index 00000000000..b86f1c921a9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c @@ -0,0 +1,38 @@ +#include + +struct str1 { + int a; + int b; +}; + +struct str2 { + int c; + int d; + struct str1 s; +}; + +int +main (int argc, char *argv[]) +{ + struct str2 t; + + t.c = 1; + t.d = 2; + t.s.a = 3; + t.s.b = 4; + + #pragma acc enter data copyin(t.s) + + #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + t.s.a = 5; + t.s.b = 6; + } + + #pragma acc exit data copyout(t.s) + + assert (t.s.a == 5); + assert (t.s.b == 6); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c new file mode 100644 index 00000000000..4dd8a3a7e17 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c @@ -0,0 +1,44 @@ +#include +#include + +struct str1 { + int a; + int b; + int *c; +}; + +#define N 1024 + +int +main (int argc, char *argv[]) +{ + struct str1 s; + + s.a = 1; + s.b = 2; + s.c = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + s.c[i] = i + 10; + + #pragma acc enter data copyin(s.a, s.b, s.c[0:N]) + + #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + s.a = 3; + s.b = 4; + for (int i = 0; i < N; i++) + s.c[i] = i + 20; + } + + #pragma acc exit data copyout(s.a, s.b, s.c[0:N]) + + assert (s.a == 3); + assert (s.b == 4); + for (int i = 0; i < N; i++) + assert (s.c[i] == i + 20); + + free (s.c); + + return 0; +}