From patchwork Wed Jul 29 14:14:38 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Andrew Stubbs X-Patchwork-Id: 1338347 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 4BGwY90VS2z9sSd for ; Thu, 30 Jul 2020 00:15:08 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 02F93388CC2E; Wed, 29 Jul 2020 14:15:05 +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 2BE7938708C5 for ; Wed, 29 Jul 2020 14:15:01 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 2BE7938708C5 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Andrew_Stubbs@mentor.com IronPort-SDR: rvdVzsp3HSX5PcIolEDPYjzrxc+a09jVTjEWiFlPo2x+VB43MCHK9VdjZQuwRFgMbbGRVTXvda w2w7bdZxfVzh1OVNUVNF8FKcHKkAkho5IGoSZJN4upjDsRdDufb/H276Hv6yNhdEw4Aoou9GCj a4jXrTRI79ci2zcWhb4i1DNhezOWffcn8Ka26D6oVqvmA2OB07oFJNYoT69lclz2QNwdlgA1JN P3sHL8LlEC9uDlLektQIgUoKSEj2+u3yZ1g+z41xfjOJo/9ZRxU8/7qL8SMIkiMySXFy00FVmh /ok= X-IronPort-AV: E=Sophos;i="5.75,410,1589270400"; d="scan'208";a="51534182" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 29 Jul 2020 06:14:43 -0800 IronPort-SDR: 8wUw/ORCpPK+ZUZujPUMODWKzkaubyQkBpyicfmKo0beq/Pogi1Ot/cLLjKjTJdZ8PcrnC5hvm KwD+/wubVxZEIAtW//pqD9PyvzohjO3EH0zZYdTtP6k+TyJwJdvncY1tQ/2gEhoH5U639UtpAm ybM00ZvA7Lrk8mt+HNYRj1LmP0gHYkBQXuAWE0k+AEDN5NRr6o+1CkXMFaH+OLSxNJkV/I0kQ1 bud6BAbAmRw7YYC5w/m632/TMpbusgPuYNz6xsKbazVRgh++kdNhDW76n9GGn8EWnqfJXD2xUR AZs= From: Andrew Stubbs Subject: [PATCH] OpenACC: Support GOMP_MAP_ZERO_LEN_ARRAY_SECTION To: "gcc-patches@gcc.gnu.org" Message-ID: Date: Wed, 29 Jul 2020 15:14:38 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.10.0 MIME-Version: 1.0 Content-Language: en-GB X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.3 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: , Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch adds support for zero-length arrays in OpenACC data transfers. Previously, trying to use an array section with zero length would cause a fatal error at runtime. This patch requires that my other patch "OpenACC: Separate enter/exit data APIs" is already applied. Unfortunately, because the reference counting is handled by the code shared with OpenMP, and because the semantics there appear to be a little bit different (or broken?), I've been unable to get acc_is_present to return true for zero-length arrays created by pragmas (those created via acc_create are fine). That issue will require a another patch, probably with more invasive changes. The test case should cover all the main uses of zero-length arrays, and I've added an xfail message to highlight the known deficiency. OK for mainline (and backport to OG10)? Andrew OpenACC: Support GOMP_MAP_ZERO_LEN_ARRAY_SECTION The shared code with OpenMP use special map kinds for zero-length arrays (detected at runtime), but the OpenACC specific code doesn't know what to do with them. This patch implements support for GOMP_MAP_ZERO_LEN_ARRAY_SECTION and GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION throughout. The last remaining problem case -- acc_is_present not reporting the array present -- is highlighted in the testcase so it doesn't get forgotten, but will need to be solved another time. libgomp/ChangeLog: * libgomp.h (splay_compare): Ensure that distinct zero-length mappings aren't confused. * oacc-mem.c (acc_is_present): Don't reject zero-sized queries. (goacc_enter_datum): Likewise. (update_dev_host): Don't actual copy zero-length arrays. (goacc_enter_data_internal): Allow tgt to be null. (goacc_exit_data_internal): Handle GOMP_MAP_ZERO_LEN_ARRAY_SECTION and GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION. * oacc-parallel.c (GOACC_update): Handle GOMP_MAP_ZERO_LEN_ARRAY_SECTION. * testsuite/libgomp.oacc-c/zerolengtharray.c: New test. diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index f9080e9f70f..e0426acdbfe 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1026,7 +1026,8 @@ struct splay_tree_key_s { static inline int splay_compare (splay_tree_key x, splay_tree_key y) { - if (x->host_start == x->host_end + if (x->host_start == y->host_start + && x->host_start == x->host_end && y->host_start == y->host_end) return 0; if (x->host_end <= y->host_start) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 45162d24786..965c81ddbd7 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -322,7 +322,7 @@ acc_is_present (void *h, size_t s) { splay_tree_key n; - if (!s || !h) + if (!h) return 0; goacc_lazy_initialize (); @@ -534,7 +534,7 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) void *d; splay_tree_key n; - if (!hostaddrs[0] || !sizes[0]) + if (!hostaddrs[0]) gomp_fatal ("[%p,+%d] is a bad range", hostaddrs[0], (int) sizes[0]); goacc_lazy_initialize (); @@ -849,6 +849,10 @@ update_dev_host (int is_dev, void *h, size_t s, int async) if (h == NULL) return; + /* Zero length arrays registered via gomp_map_vars don't show as mapped. */ + if (s == 0) + return; + acc_prof_info prof_info; acc_api_info api_info; bool profiling_p = GOACC_PROFILING_SETUP_P (thr, &prof_info, &api_info); @@ -1203,16 +1207,17 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, &sizes[i], &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); - assert (tgt); gomp_mutex_lock (&acc_dev->lock); - for (size_t j = 0; j < tgt->list_count; j++) - { - n = tgt->list[j].key; - if (n && !tgt->list[j].is_attach) - n->dynamic_refcount++; - } + /* TGT can be null for zero-length arrays. */ + if (tgt) + for (size_t j = 0; j < tgt->list_count; j++) + { + n = tgt->list[j].key; + if (n && !tgt->list[j].is_attach) + n->dynamic_refcount++; + } } i = group_last; @@ -1276,6 +1281,8 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, case GOMP_MAP_POINTER: case GOMP_MAP_DELETE: case GOMP_MAP_RELEASE: + case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: + case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION: { struct splay_tree_key_s cur_node; size_t size; diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index bca31b51427..d3277e60404 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -647,6 +647,7 @@ GOACC_update (int flags_m, size_t mapnum, { case GOMP_MAP_POINTER: case GOMP_MAP_TO_PSET: + case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: break; case GOMP_MAP_ALWAYS_POINTER: diff --git a/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c b/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c new file mode 100644 index 00000000000..cae102cb580 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c @@ -0,0 +1,78 @@ +/* Ensure that GOMP_MAP_ZERO_LEN_ARRAY_SECTION is supported. */ + +/* { dg-do run } */ + +#include +#include + +#ifndef ACC_MEM_SHARED +#define ACC_MEM_SHARED 0 +#endif + +int a[100]; +int b[100]; +int c[100]; +int d[100]; +int e[100]; + +int s = 0; // runtime size is zero + +int main () +{ + /* Check it's not present too early. + Host fallback and shared-memory systems will show the data present. */ + if (acc_is_present (a, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (b, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (d, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (e, s) == !ACC_MEM_SHARED) + abort (); + + /* Test GOMP_MAP_ZERO_LEN_ARRAY_SECTION inputs. */ +#pragma acc enter data create(a[0:s]) +#pragma acc enter data copyin(b[0:s]) + acc_create (d, s); + acc_copyin (e, s); +#pragma acc update device(a[0:s]) + acc_update_device (b, s); + + // FIXME: update these when they work correctly + // TODO { dg-output "acc_is_present does not currently work for zero-length arrays created via pragmas" { xfail *-*-* } } + if (/*!*/acc_is_present (a, s) == !ACC_MEM_SHARED) + abort (); + if (/*!*/acc_is_present (b, s) == !ACC_MEM_SHARED) + abort (); + if (!acc_is_present (d, s)) + abort (); + if (!acc_is_present (e, s)) + abort (); + +#pragma acc parallel copy(c[0:s]) + { + ; + } + + /* Test GOMP_MAP_ZERO_LEN_ARRAY_SECTION outputs. */ + acc_update_self (a, s); +#pragma acc update self(b[0:s]) + acc_copyout (e, s); +#pragma acc exit data copyout(b[0:s]) + + /* Test GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION. */ + acc_delete (d, s); +#pragma acc exit data delete(a[0:s]) finalize + + /* Check it's no longer present. */ + if (acc_is_present (a, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (b, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (d, s) == !ACC_MEM_SHARED) + abort (); + if (acc_is_present (e, s) == !ACC_MEM_SHARED) + abort (); + + return 0; +}