From patchwork Fri May 22 22:16:04 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1296481 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 49TLTN6Bggz9sRf for ; Sat, 23 May 2020 08:17:44 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 6A50F3983064; Fri, 22 May 2020 22:17:40 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id F3964386F83F for ; Fri, 22 May 2020 22:17:37 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org F3964386F83F 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: vO3OT/6lx4d7EhyNA9MAeZ0N9KeuFQ5w7vsRgmAg11FNd99XIHEzNVXB39gY9N3L/EaAiHHzrZ Cyj2mBe43fYl2QV2b84o8cdffg2GMWRPZnewgKGA28BxPogzjAs24Gv0WpcA3lfQBIF93xf9xz YDc8efhUIhqeuOb9Fs4D5H48yrH5+qUK9+P6V/1klwo8R0mOV4Gy+QllRmMFMoG7VyHSiI2Jju CgNosHOJL5Y6uyCapxR1X2ijvEsP1eW/LxNi8p3INvv+pSXGAs5+nDVDIa3SM97TdOfyHUGRVm DxA= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="51165107" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 22 May 2020 14:17:21 -0800 IronPort-SDR: +kKfPwFQWAWXdoItH2WbqyNT0LhBITKELbq4Ud00xfbGL7qpqXdfRTHEoMJu0dB83+8/JVUSQS BvZueZYRwzIcPzLuTzrkveWAlhV7IhRbEArcM9LdMLlNXZyfwrQsERKLZVR/i8J2nS2KoyyaA9 /dPnI/zqvx9LzPj1WI/hXL3GvFneTD3LwOIuiP98Ki4PhCy/Cp+OvMzAcIZeIsp8Qo0BXPbtZj zxYG2VgAkKkOmabgZmho5OxFIpCBQ/mnxjowu5StnpmRL9hsGH31DefI8RZ4TO4ZW0sK+lQQTe X2U= From: Julian Brown To: Subject: [PATCH 1/7] [OpenACC] Missing unlocking on error paths in attach/detach code Date: Fri, 22 May 2020 15:16:04 -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-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.0 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch adds some missing unlocking from error paths in the OpenACC attach/detach code, noticed during development of other patches in this series. OK? Julian ChangeLog libgomp/ * oacc-mem.c (acc_attach_async): Add missing gomp_mutex_unlock on error path. (goacc_detach_internal): Likewise. --- libgomp/oacc-mem.c | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2d4bba78efd..c06b7341cbb 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -887,7 +887,10 @@ acc_attach_async (void **hostaddr, int async) n = splay_tree_lookup (&acc_dev->mem_map, &cur_node); if (n == NULL) - gomp_fatal ("struct not mapped for acc_attach"); + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("struct not mapped for acc_attach"); + } gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, 0, NULL); @@ -920,7 +923,10 @@ goacc_detach_internal (void **hostaddr, int async, bool finalize) n = splay_tree_lookup (&acc_dev->mem_map, &cur_node); if (n == NULL) - gomp_fatal ("struct not mapped for acc_detach"); + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("struct not mapped for acc_detach"); + } gomp_detach_pointer (acc_dev, aq, n, (uintptr_t) hostaddr, finalize, NULL); From patchwork Fri May 22 22:16:05 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1296484 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 49TLZb2xpQz9sRf for ; Sat, 23 May 2020 08:22:15 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 198003851C0C; Fri, 22 May 2020 22:21:33 +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 5819C3851C0C for ; Fri, 22 May 2020 22:21:29 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 5819C3851C0C 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: 0YvvyxeUCpYfSgE0qFoF6x0MvnFPXJ1SFfqggFB61EOMYFVCJ96DPfV1EZB6t8cNoZshinMopN s7i+rsdLY7fUNMCI30E6Gh75WvQkfkAcpNePxM88eyliHtviEUTaK4fLp+J3XMlIkaIiqbJ2bG zLn8GXi9DYe5MFK+uAji9WIrrJZ1ESzkSoMWzoqbwiyqwlPHiD3Iu4hHKh2BFrPau1tZSVU2j5 m0kdLR19Y3nVF+QGyd5iY2aA3Jo8ITccsSx839ciNSkqNbYt0Q4l1fpnX+gB5ZQV2lmAYf7NhS pnE= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127036" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:21:28 -0800 IronPort-SDR: Fm1fUrk8vBH6oUOfA2CqVe86eGkY0uEj1O5d99ENTTBd2ERzjmqcprddwtrmsUHnP3yP89U5T4 RuPb4aqgkHEA4YHiecu9qod6OSALd+5iK+jXNlSVAm97jaoc8S2cSZxCYZDZgYaSdt9tyntcdD GFb2T9qj8EW/D734ADjl2N11eQY62NrkidDmFXAuPbs5EmKCCb/VE3cFzP6NX41jpZstgIBC8B toa1ZWzSqtJTBPe3XHdLZI2/uT0wbY6901MhHOrWoa8fj8O8+hYQ7KliGWjGmMUTEO4IkMdZCo JnA= From: Julian Brown To: Subject: [PATCH 2/7] [OpenACC] Adjust dynamic reference count semantics Date: Fri, 22 May 2020 15:16:05 -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-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.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch adjusts the semantics of dynamic reference counts, as described in the parent email. There are also two new test cases derived from Thomas's test in the email: https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546166.html that work now. 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. (acc_unmap_data): Replace open-coded refcount handling with call to gomp_remove_var. (goacc_enter_datum): Adjust for dynamic_refcount semantics. Use tgt returned from gomp_map_vars_async. Update assertions. (goacc_exit_datum): Re-add some error checking. Adjust for dynamic_refcount semantics. Fix is_tgt_unmapped test for struct mappings. (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. Re-introduce error checking for tgt unmapping when appropriate. * 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. --- libgomp/libgomp.h | 8 +- libgomp/oacc-mem.c | 241 ++++++++++++------ libgomp/target.c | 38 +-- .../libgomp.oacc-c-c++-common/refcounting-1.c | 31 +++ .../libgomp.oacc-c-c++-common/refcounting-2.c | 31 +++ 5 files changed, 243 insertions(+), 106 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 c06b7341cbb..fff0d573f59 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -407,7 +407,7 @@ acc_map_data (void *h, void *d, size_t s) assert (tgt); splay_tree_key n = tgt->list[0].key; assert (n->refcount == 1); - assert (n->virtual_refcount == 0); + assert (n->dynamic_refcount == 0); /* Special reference counting behavior. */ n->refcount = REFCOUNT_INFINITY; @@ -454,7 +454,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.)? */ @@ -475,14 +475,19 @@ acc_unmap_data (void *h) gomp_mutex_unlock (&acc_dev->lock); gomp_fatal ("cannot unmap target block"); } - else if (tgt->refcount > 1) - tgt->refcount--; - else + + if (tgt->refcount == 1) { - free (tgt->array); - free (tgt); + /* This is the last reference. Nullifying these fields prevents + 'gomp_unmap_tgt' via 'gomp_remove_var' from freeing the target + memory. */ + tgt->tgt_end = 0; + tgt->to_free = NULL; } + bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); + assert (is_tgt_unmapped); + gomp_mutex_unlock (&acc_dev->lock); if (profiling_p) @@ -540,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); } @@ -555,16 +558,18 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) goacc_aq aq = get_goacc_asyncqueue (async); - gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, - true, GOMP_MAP_VARS_OPENACC_ENTER_DATA); + struct target_mem_desc *tgt + = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, + 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->dynamic_refcount == 0); + n->dynamic_refcount++; - gomp_mutex_lock (&acc_dev->lock); - n = lookup_host (acc_dev, hostaddrs[0], sizes[0]); - assert (n != NULL); - assert (n->tgt_offset == 0); - assert ((uintptr_t) hostaddrs[0] == n->host_start); - d = (void *) n->tgt->tgt_start; - gomp_mutex_unlock (&acc_dev->lock); + d = (void *) tgt->tgt_start; } if (profiling_p) @@ -683,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) { @@ -722,8 +732,16 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async) gomp_remove_var_async (acc_dev, n, aq); else { + int num_mappings = 0; + /* If the target_mem_desc represents a single data mapping, we can + check that it is freed when this splay tree key's refcount + reaches zero. Otherwise (e.g. for a struct mapping with multiple + members), fall back to skipping the test. */ + for (int i = 0; i < n->tgt->list_count; i++) + if (n->tgt->list[i].key) + num_mappings++; bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); - assert (is_tgt_unmapped); + assert (num_mappings > 1 || is_tgt_unmapped); } } @@ -1018,13 +1036,102 @@ 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; - 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); + 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_mutex_unlock (&acc_dev->lock); + } + else if (n && groupnum > 1) + { + assert (n->refcount != REFCOUNT_INFINITY + && n->refcount != REFCOUNT_LINK); + + 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; } @@ -1115,18 +1222,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)) @@ -1137,45 +1241,40 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, cur_node.host_end - cur_node.host_start); if (n->refcount == 0) - gomp_remove_var_async (acc_dev, n, aq); - } - 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 (aq) { - if (finalize) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount -= str->virtual_refcount; - str->virtual_refcount = 0; - } - if (str->virtual_refcount > 0) + /* TODO The way the following code is currently + implemented, we need the 'is_tgt_unmapped' return + value from 'gomp_remove_var', so can't use + 'gomp_remove_var_async' here -- see the + 'gomp_unref_tgt' comment in + ; + PR92881 -- so have to synchronize here. */ + if (!acc_dev->openacc.async.synchronize_func (aq)) { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount--; - str->virtual_refcount--; + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("synchronize failed"); } - else if (str->refcount > 0 - && str->refcount != REFCOUNT_INFINITY) - str->refcount--; - if (str->refcount == 0) - gomp_remove_var_async (acc_dev, str, aq); } + int num_mappings = 0; + /* If the target_mem_desc represents a single data mapping, we + can check that it is freed when this splay tree key's + refcount reaches zero. Otherwise (e.g. for a struct + mapping with multiple members), fall back to skipping the + test. */ + for (int j = 0; j < n->tgt->list_count; j++) + if (n->tgt->list[j].key) + num_mappings++; + bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); + assert (num_mappings > 1 || is_tgt_unmapped); } - i += elems; } break; + case GOMP_MAP_STRUCT: + continue; + default: gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x", kind); 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; +} + From patchwork Fri May 22 22:16:06 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1296482 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 49TLZV0bdqz9sPK for ; Sat, 23 May 2020 08:22:09 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A018E388B02E; Fri, 22 May 2020 22:21:32 +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 42465386F824 for ; Fri, 22 May 2020 22:21:30 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 42465386F824 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: bUFty+3hRDWJPh4pd/EXT3roi1YW83LzeSochZ0qPitkJvPn6+OCm7Dhfx+tqQWTCRzwg3lRaq mrEbA5bHpu9/asg9iHm8Kh4Hqm1ViUDkPFLpzFgG35TDzS/c39bEZpfkJoctGhwhRDiKRQP3ix bWFtXYhEu8ETq6cLRQ4MZApH5xxbDe0z27pAaqxwvahAFtQ8GiSz5hAsIqB/w6qJKW/22KkIDg ddjOKa7qjleX9m7ClqspWtKfM//xSXYiaKuWJRc2BwctyENN6SyqxmzdUDk9fvx/dyUwV+JdQ/ Nic= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127038" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:21:28 -0800 IronPort-SDR: as6Ski29x4X0DOcxzAO8VrlYxOwPhECeucFm1N5z7aS7npBNyyp/+ZrHow93eETzaLKEf9VqMP CuAv8vF5pFsLTdIGPWUJ8M2o56UwplLWhid5kG4hSg36pj8SnqI/D82SHprihdHUreRkD6/wUv BlAhvEThhr6med3+4ZQ40hkjM7CRApxL0Whk+PA7ILcj6iIJ1YlOAHQQTQSrqX6ir7JZ8BsOD0 g/mJol/PRWK0V5lHFw2331JUqxcovztN6kUWJU2Q0o2JTsw063I037sJUQq9eLVm/hx2kKEwJy GJw= From: Julian Brown To: Subject: [PATCH 3/7] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum Date: Fri, 22 May 2020 15:16:06 -0700 Message-ID: <93a8c26ad510454f6326705ecb20f99fd8582ca5.1590182783.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-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.4 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Since goacc_enter_datum only maps a single data item now, there is no need to pass "kinds" as an array. Passing as a scalar allows for some simplification in the function's callers. 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 fff0d573f59..20d241382a8 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 Fri May 22 22:16:07 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1296483 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 49TLZX5zvyz9sPK for ; Sat, 23 May 2020 08:22:12 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 0DCB03984010; Fri, 22 May 2020 22:21:36 +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 D1D12386F824 for ; Fri, 22 May 2020 22:21:31 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org D1D12386F824 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: q2ujycTYFGc3JraPaMf888UgC17i7L0cFhhAFxqBNje5ahGQ6QV0TfH0dzjWJl/H3w4CSt8oFl TEkssNFspOIdN2pjfa98Mv6ss8flAUP6uoHauIUmhV/M+tBJSMN5jVkFR+EUfExVhovcMo2jnA MKCIXgBP5dBPVgWSIKtNmjjMPNw5QyUxuxw+073sOOwF4Q3/GSYpTVOFpWnBluCXnU23Zs3Olm OrMtfW1qoWe+qrup1r3qBNRdim+UMWjlY1f1EshDVL+s5rfpE1WxP4mGCYY856U5jlxHeM4wxl /D0= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127040" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:21:29 -0800 IronPort-SDR: M0EWEvVK3pyAlsZ+arXnF6rTfBatHuNtQD5DjsE49Lsu3NUQ6Gsr2hPDV5PJO5FmcAUn0mjlpP 0cx0ZBjVrH8zasPcoz7ElgCPCII1ZKQb9gR23XMRBKAFBLf8E37mhydQwrf1acIYeoMSAWXvqP d3fwYt4VclZjh6COVzXoHhzmEwBc8Rhlz5MZPQLA1xqOEz1/yBBc4Q4NCIStqciQLTImP7b+Qq hhPSWUiur6WjRNXFG2TsDVKhDQXoplb41sHR02e429jqAb4Qd92O6/Wl48oz/oFe/Oh1HA0fDN ABk= From: Julian Brown To: Subject: [PATCH 4/7] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843) Date: Fri, 22 May 2020 15:16:07 -0700 Message-ID: <7b2f54faa0a25a7a445ad86bf4726202a1190a4f.1590182783.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-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.9 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch provides a bug fix (on top of previous patches in this series) that allows the PR92843 test case to pass. Data mapped in with "acc_map_data" is not copied out by an "exit data" directive. OK? Julian 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. --- 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 20d241382a8..c2b4a131a5f 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1230,6 +1230,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 Fri May 22 22:21: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: 1296485 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 49TLds1BJMz9sPK for ; Sat, 23 May 2020 08:25:05 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7D50A38708F9; Fri, 22 May 2020 22:25:02 +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 0E4D13851C0C for ; Fri, 22 May 2020 22:24:58 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 0E4D13851C0C 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: Z/xo+iBK+PcvFfMZ1TjlICDDUEXTpt+QJFmCswFoV++W5yT/DEtN0zNnouaIb3rfpaCek/wIW5 3EeQ8kUsibloQ/9SxvuNeY/nyfdue0irbxMIlNpr7ZhPKKyJfiRck90lNzMRadRW0uAh7VU5vO od2lpRvYyYZ188Wpl6Rbw5G5ytIVfq+985F3An9L3GLt7OeyC5YvsAgYlRdj4ShHeLQ0dhEFSx p1nVifmBlCd498drV1xT/nWnWikdIgdyNz3W1wnr/kncUgNJ4JBqxvU/tn0tlJUJ8Rph6klGuS sDg= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127126" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:24:58 -0800 IronPort-SDR: 2blM/CambvTDmT1HVhOEml10dfeA9WQ+DbJTxZzgOOCqJXdK4W8r/WywFgmNtjyCSsumnL95Zv oCLtGVMWw0qT5wmdyWYtDk3CiDU2wcUVk5VMKHid5CvcRPS9OE2CaGup4Vw8VYAGradITcN55l z5+WsYsoMlbpX5BTNPjgVi0WJtkxPqOZGo64w3BNfhhZG86YSlZkZPw6iyQYeWnb/lEDT5xYwC sXEMnIKtyycAy9U5UFOiocvOz4GUepPPa5F+KNjsN+yxt0dYte8nh9kiIw1e6ahIxfwQzUkBrR eWg= From: Julian Brown To: Subject: [PATCH 5/7] [OpenACC] Distinguish structural/dynamic mappings in libgomp Date: Fri, 22 May 2020 15:21:43 -0700 Message-ID: <0e04de91b83d3590e2dd35cf37e190cc771af823.1590182783.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-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-13.1 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch provides support for distinguishing target_mem_descs introduced via structured data lifetimes from those arising from dynamic data lifetimes. This is a prerequisite for the following reference-count self-checking patch. This patch (and those following it) are not vital for this patch series, but are "nice-to-have" additions. OK? Julian libgomp/ * libgomp.h (struct target_mem_desc): Update comment on prev field. * oacc-int.h (goacc_mark_dynamic, goacc_marked_dynamic_p): Add prototypes. * oacc-mem.c (dyn_tgt_sentinel): New. (goacc_mark_dynamic, goacc_marked_dynamic_p): New functions. (goacc_enter_datum): Call goacc_mark_dynamic. (goacc_enter_data_internal): Likewise. * target.c (gomp_unmap_vars_internal): Convert a target_mem_desc from a structural mapping to dynamic when appropriate. --- libgomp/libgomp.h | 3 ++- libgomp/oacc-int.h | 3 +++ libgomp/oacc-mem.c | 28 ++++++++++++++++++++++++++++ libgomp/target.c | 8 +++++++- 4 files changed, 40 insertions(+), 2 deletions(-) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 7b52ce7d5c2..0d1978ffb13 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -971,7 +971,8 @@ struct target_mem_desc { uintptr_t tgt_end; /* Handle to free. */ void *to_free; - /* Previous target_mem_desc. */ + /* Previous target_mem_desc. Also used in OpenACC to indicate that this + target_mem_desc is used only for an "enter data" mapping. */ struct target_mem_desc *prev; /* Number of items in following list. */ size_t list_count; diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h index 3c2c9b84b2f..2d8d3eb5a4b 100644 --- a/libgomp/oacc-int.h +++ b/libgomp/oacc-int.h @@ -165,6 +165,9 @@ bool _goacc_profiling_setup_p (struct goacc_thread *, void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *, acc_api_info *); +extern void goacc_mark_dynamic (struct target_mem_desc *); +extern bool goacc_marked_dynamic_p (struct target_mem_desc *tgt); + #ifdef HAVE_ATTRIBUTE_VISIBILITY # pragma GCC visibility pop #endif diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index c2b4a131a5f..038ab68e8a2 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -497,6 +497,30 @@ acc_unmap_data (void *h) } } +/* Indicate (via storing its address in the "prev" field) a target_mem_desc + that is used for an "enter data" mapping. */ +const static struct target_mem_desc dyn_tgt_sentinel; + +attribute_hidden void +goacc_mark_dynamic (struct target_mem_desc *tgt) +{ + tgt->prev = (struct target_mem_desc *) &dyn_tgt_sentinel; +} + +attribute_hidden bool +goacc_marked_dynamic_p (struct target_mem_desc *tgt) +{ + return tgt->prev == (struct target_mem_desc *) &dyn_tgt_sentinel; +} /* Enter dynamic mapping for a single datum. Return the device pointer. */ @@ -563,6 +587,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, unsigned short kind, = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, &kind, true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); + goacc_mark_dynamic (tgt); + assert (tgt->list_count == 1); n = tgt->list[0].key; assert (n); @@ -1122,6 +1148,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, &sizes[i], &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); + goacc_mark_dynamic (tgt); + for (size_t j = 0; j < tgt->list_count; j++) { n = tgt->list[j].key; diff --git a/libgomp/target.c b/libgomp/target.c index 3f2becdae0e..1d60d0cb573 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1447,7 +1447,13 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, bool do_unmap = false; if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) - k->refcount--; + { + k->refcount--; + /* If we only have dynamic references left, mark the tgt_mem_desc + appropriately. */ + if (k->refcount == k->dynamic_refcount) + goacc_mark_dynamic (k->tgt); + } else if (k->refcount == 1) { k->refcount--; From patchwork Fri May 22 22:21: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: 1296486 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 49TLdx1whmz9sPK for ; Sat, 23 May 2020 08:25:09 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7EBC2398403F; Fri, 22 May 2020 22:25:05 +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 591903851C0C for ; Fri, 22 May 2020 22:25:01 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 591903851C0C 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: jrKUMmXawHTiQ7HYq/EFpgscrmOpi+jmR6VnBCuCb0hiwnVWQEY/cnoV48qwNysq35vIyKHPlE ok/F0OpCV3ikbECAXB9Cxnf5ju/GBTic//GfFmb2MB3ImWCHLGDkuiEeZxx/p73U3INtm3/tq9 dPzz7qE/GCRnjMYPStAUpBJS4AXqs3B6u4BZJWD/629pplJHJUL+eZLZd7RJeDJCK3SugnSImb zU2mxq66GkmdXrwpDr+4W1AeqXndt5rX2uG7S1NxS7VvF0rwb+knL9DUdYuo7vFhmJZfC5AZYY JYc= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127128" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:25:00 -0800 IronPort-SDR: n5TZRCN1E6EQ7J7ac7Ywymp4cyUndCY+joQatpudOnNWKWPXh1OhMOf+hNAbRMV5DAPtgv29zm o2v9JSxvXK8a19+njrIy6oK7j+9sEWKgdL0rEQTKhymCcTYDSNEij8OJNHI67yygB37auwKaFd tONOIiLBXvmJxqme3EvDGZgpf4gl0boP/l7cUQioUVAQYYA04WuGJxtPLX0H+dREd+ETtzjd7p d/xdYlp1LieJGaze5KYiWmRe3PFeqfVHsnCqgcdAwBrtfjBEutOKLxjJ3bNOCpLbLSV/HNUmxK d4Q= From: Julian Brown To: Subject: [PATCH 6/7] [OpenACC] Reference count self-checking (dynamic_refcount version) Date: Fri, 22 May 2020 15:21:44 -0700 Message-ID: <1cc9ea4c5807c5da2df9f17a0a11935e78b0c721.1590182783.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-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-13.2 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This is a new version of the reference count self-checking code, adjusted to work with the new (old) dynamic_refcount counting scheme. The key observation is that a target_mem_desc that was created from a dynamic data lifetime should not contribute to the structured refcount for splay tree keys in its variable list. We can figure out which target_mem_descs that applies to using the information recorded in the previous patch. In a sense, this takes the "awkward corner cases" from the virtual_refcount ("overhaul") patch, and moves them to the optional self-test code, where they can potentially do less harm. With this, we still have a formal-ish model of what refcounts mean and some confidence that they remain consistent (at least throughout execution of a test run), which I think is a good thing. OK? (We probably want a way of configuring-in this testing automatically, as mentioned previously.) Julian ChangeLog libgomp/ * libgomp.h (RC_CHECKING): New macro, disabled by default, guarding all hunks in this patch. (target_mem_desc): Add refcount_chk, mark fields. (splay_tree_key_s): Add refcount_chk field. (dump_tgt, gomp_rc_check): Add prototypes. * oacc-mem.c (GOACC_enter_exit_data): Add refcount self-check code. * oacc-parallel.c (GOACC_parallel_keyed_internal): Add refcount self-check code. (GOACC_data_start, GOACC_data_end, GOACC_enter_exit_data): Likewise. * target.c (stdio.h): Include. (dump_tgt, rc_check_clear, rc_check_count, rc_check_verify, gomp_rc_check): New functions to consistency-check reference counts. --- libgomp/libgomp.h | 18 ++++ libgomp/oacc-mem.c | 6 ++ libgomp/oacc-parallel.c | 27 ++++++ libgomp/target.c | 185 ++++++++++++++++++++++++++++++++++++++++ 4 files changed, 236 insertions(+) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 0d1978ffb13..eaa7c6ebb4c 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -960,9 +960,17 @@ struct target_var_desc { uintptr_t length; }; +/* Uncomment to enable reference-count consistency checking (for development + use only). */ +//#define RC_CHECKING 1 + struct target_mem_desc { /* Reference count. */ uintptr_t refcount; +#ifdef RC_CHECKING + uintptr_t refcount_chk; + bool mark; +#endif /* All the splay nodes allocated together. */ splay_tree_node array; /* Start of the target region. */ @@ -1019,6 +1027,10 @@ struct splay_tree_key_s { uintptr_t refcount; /* Dynamic reference count. */ uintptr_t dynamic_refcount; +#ifdef RC_CHECKING + /* The recalculated reference count, for verification. */ + uintptr_t refcount_chk; +#endif struct splay_tree_aux *aux; }; @@ -1174,6 +1186,12 @@ extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); +#ifdef RC_CHECKING +extern void dump_tgt (const char *, struct target_mem_desc *); +extern void gomp_rc_check (struct gomp_device_descr *, + struct target_mem_desc *); +#endif + extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *, size_t, void **, void **, size_t *, void *, bool, diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 038ab68e8a2..c8ec3c9a7dd 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1450,4 +1450,10 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr->prof_info = NULL; thr->api_info = NULL; } + +#ifdef RC_CHECKING + gomp_mutex_lock (&acc_dev->lock); + gomp_rc_check (acc_dev, thr->mapped_data); + gomp_mutex_unlock (&acc_dev->lock); +#endif } diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index c7e46e35bd6..0774cdc7e4f 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -301,6 +301,15 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), &api_info); } +#ifdef RC_CHECKING + gomp_mutex_lock (&acc_dev->lock); + assert (tgt); + dump_tgt (__FUNCTION__, tgt); + tgt->prev = thr->mapped_data; + gomp_rc_check (acc_dev, tgt); + gomp_mutex_unlock (&acc_dev->lock); +#endif + devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, i); @@ -347,6 +356,12 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), thr->prof_info = NULL; thr->api_info = NULL; } + +#ifdef RC_CHECKING + gomp_mutex_lock (&acc_dev->lock); + gomp_rc_check (acc_dev, thr->mapped_data); + gomp_mutex_unlock (&acc_dev->lock); +#endif } /* Legacy entry point (GCC 5). Only provide host fallback execution. */ @@ -481,6 +496,12 @@ GOACC_data_start (int flags_m, size_t mapnum, thr->prof_info = NULL; thr->api_info = NULL; } + +#ifdef RC_CHECKING + gomp_mutex_lock (&acc_dev->lock); + gomp_rc_check (acc_dev, thr->mapped_data); + gomp_mutex_unlock (&acc_dev->lock); +#endif } void @@ -554,6 +575,12 @@ GOACC_data_end (void) thr->prof_info = NULL; thr->api_info = NULL; } + +#ifdef RC_CHECKING + gomp_mutex_lock (&thr->dev->lock); + gomp_rc_check (thr->dev, thr->mapped_data); + gomp_mutex_unlock (&thr->dev->lock); +#endif } void diff --git a/libgomp/target.c b/libgomp/target.c index 1d60d0cb573..9a51e1c70f6 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -38,6 +38,9 @@ #include #include #include +#ifdef RC_CHECKING +#include +#endif #ifdef PLUGIN_SUPPORT #include @@ -347,6 +350,188 @@ gomp_free_device_memory (struct gomp_device_descr *devicep, void *devptr) } } +#ifdef RC_CHECKING +void +dump_tgt (const char *where, struct target_mem_desc *tgt) +{ + if (!getenv ("GOMP_DEBUG_TGT")) + return; + + fprintf (stderr, "%s: %s: tgt=%p\n", __FUNCTION__, where, (void*) tgt); + fprintf (stderr, "refcount=%d\n", (int) tgt->refcount); + fprintf (stderr, "tgt_start=%p\n", (void*) tgt->tgt_start); + fprintf (stderr, "tgt_end=%p\n", (void*) tgt->tgt_end); + fprintf (stderr, "to_free=%p\n", tgt->to_free); + fprintf (stderr, "list_count=%d\n", (int) tgt->list_count); + for (int i = 0; i < tgt->list_count; i++) + { + fprintf (stderr, "list item %d:\n", i); + fprintf (stderr, " key: %p\n", (void*) tgt->list[i].key); + if (tgt->list[i].key) + { + fprintf (stderr, " key.host_start=%p\n", + (void*) tgt->list[i].key->host_start); + fprintf (stderr, " key.host_end=%p\n", + (void*) tgt->list[i].key->host_end); + fprintf (stderr, " key.tgt=%p\n", (void*) tgt->list[i].key->tgt); + fprintf (stderr, " key.offset=%d\n", + (int) tgt->list[i].key->tgt_offset); + fprintf (stderr, " key.refcount=%d\n", + (int) tgt->list[i].key->refcount); + fprintf (stderr, " key.dynamic_refcount=%d\n", + (int) tgt->list[i].key->dynamic_refcount); + if (tgt->list[i].key->aux) + { + fprintf (stderr, " key.aux.link_key=%p\n", + (void*) tgt->list[i].key->aux->link_key); + fprintf (stderr, " key.aux.attach_count=%p\n", + (void*) tgt->list[i].key->aux->attach_count); + } + } + } + fprintf (stderr, "\n"); +} + +static void +rc_check_clear (splay_tree_node node) +{ + splay_tree_key k = &node->key; + + k->refcount_chk = 0; + k->tgt->refcount_chk = 0; + k->tgt->mark = false; + + if (node->left) + rc_check_clear (node->left); + if (node->right) + rc_check_clear (node->right); +} + +static void +rc_check_count (splay_tree_node node) +{ + splay_tree_key k = &node->key; + struct target_mem_desc *t; + + /* Add dynamic reference counts ("acc enter data", etc.) for this key. */ + k->refcount_chk += k->dynamic_refcount; + + t = k->tgt; + t->refcount_chk++; + + /* Do not count references from tgt_mem_descs that arise from dynamic data + lifetimes: those are counted already by their keys' dynamic_refcount. */ + if (!t->mark && goacc_marked_dynamic_p (t)) + t->mark = true; + else if (!t->mark) + { + for (int i = 0; i < t->list_count; i++) + if (t->list[i].key) + t->list[i].key->refcount_chk++; + + t->mark = true; + } + + if (node->left) + rc_check_count (node->left); + if (node->right) + rc_check_count (node->right); +} + +static bool +rc_check_verify (splay_tree_node node, bool noisy, bool errors) +{ + splay_tree_key k = &node->key; + struct target_mem_desc *t; + + if (k->refcount != REFCOUNT_INFINITY) + { + if (noisy) + fprintf (stderr, "key %p (%p..+%d): rc=%d/%d, dyn_rc=%d\n", k, + (void *) k->host_start, (int) (k->host_end - k->host_start), + (int) k->refcount, (int) k->refcount_chk, + (int) k->dynamic_refcount); + + if (k->refcount != k->refcount_chk) + { + if (noisy) + fprintf (stderr, " -- key refcount mismatch!\n"); + errors = true; + } + + t = k->tgt; + + if (noisy) + fprintf (stderr, "tgt %p: rc=%d/%d\n", t, (int) t->refcount, + (int) t->refcount_chk); + + if (t->refcount != t->refcount_chk) + { + if (noisy) + fprintf (stderr, + " -- target memory descriptor refcount mismatch!\n"); + errors = true; + } + } + + if (node->left) + errors |= rc_check_verify (node->left, noisy, errors); + if (node->right) + errors |= rc_check_verify (node->right, noisy, errors); + + return errors; +} + +/* Call with device locked. */ + +attribute_hidden void +gomp_rc_check (struct gomp_device_descr *devicep, struct target_mem_desc *tgt) +{ + splay_tree sp = &devicep->mem_map; + + bool noisy = getenv ("GOMP_DEBUG_TGT") != 0; + + if (noisy) + fprintf (stderr, "\n*** GOMP_RC_CHECK ***\n\n"); + + if (sp->root) + { + rc_check_clear (sp->root); + + for (struct target_mem_desc *t = tgt; t; t = t->prev) + { + t->refcount_chk = 0; + t->mark = false; + } + + /* Add references for interconnected splay-tree keys. */ + rc_check_count (sp->root); + + /* Add references for the tgt for a currently-executing kernel and/or + any enclosing data directives. */ + for (struct target_mem_desc *t = tgt; t; t = t->prev) + { + t->refcount_chk++; + + if (!t->mark) + { + for (int i = 0; i < t->list_count; i++) + if (t->list[i].key) + t->list[i].key->refcount_chk++; + + t->mark = true; + } + } + + if (rc_check_verify (sp->root, noisy, false)) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("refcount checking failure"); + } + } +} +#endif + /* Handle the case where gomp_map_lookup, splay_tree_lookup or gomp_map_0len_lookup found oldn for newn. Helper function of gomp_map_vars. */ From patchwork Fri May 22 22:21: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: 1296487 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 49TLf13Hmwz9sPK for ; Sat, 23 May 2020 08:25:13 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E36DB3984059; Fri, 22 May 2020 22:25:05 +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 B025A3851C0C for ; Fri, 22 May 2020 22:25:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org B025A3851C0C 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: JGZ0eVXgBYMkHA+TlMYc/0yMm72yjvD+W8J3q0oA7q+vAeGmrVRydiIMD65AF3S5pFNcQrr2PW aOaSsxHaue9wKI5Zi8U5nXX4rNiThxASN3jIkurWXOnGc6isUZbx3gHthnAEyZ436qxj6VfGdZ uWklSoFcwXBMO6JMHb94l4bVdgI4yFRr09+ih5Srtji0iAlfg1uz2mX68xK3ycPt8SUoGrQVfq zZnzrpzpMVBWF1bvyRpZB15R3UVg/1DEidv3Y6SfoAl96/dc+nr6szPfzpCaCVf0DiNxp32QOX bo0= X-IronPort-AV: E=Sophos;i="5.73,423,1583222400"; d="scan'208";a="49127132" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 22 May 2020 14:25:03 -0800 IronPort-SDR: akaqjmJx37S/Qk6IDPxjuvuAjWbgDgMZUkO20RrP/qRg4UZviOIdceiEYgci3ER7xYsfGrk0kF 72LOqWewGZ7HzncakEdhNrZYLCvpBpoKPNAZXo+L2vaAvAmUeow6xVSqCArs7zW/ulS8kawJSf qt8QpAcZXTCoyKIvQLmI4Wlrtk6JmwDWQAk3XLGG6IAiSc/EjhHFUdJ53QWwvdGeQjkNvz0uYz xDZGDqbSZTlbg7CTOlI+1ghoboEeImQhyGE5ndoUNzzDSobMAOj+CqViAEjevR/SADV/AqIQUY Bos= From: Julian Brown To: Subject: [PATCH 7/7] [OpenACC] Stricter dynamic data unmapping testing (WIP) Date: Fri, 22 May 2020 15:21:45 -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-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-13.3 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, jakub@redhat.com, thomas@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Using the ability to distinguish structural from dynamic mappings' target_mem_descs, we can adjust how the assertions in goacc_exit_datum and goacc_exit_data_internal work. This is possibly a slightly stronger test than the one introduced earlier in this patch series -- though actually I haven't quite convinced myself of that. Anyway, this passes a regression run, with the refcount self-checking code enabled also. OK, or any comments? Julian ChangeLog libgomp/ * oacc-mem.c (goacc_exit_datum): Adjust self-test code. (goacc_exit_data_internal): Likewise. * target.c (gomp_unmap_vars_internal): Clear target_mem_desc variable list keys on unmapping. --- libgomp/oacc-mem.c | 43 ++++++++++++++++++++++++------------------- libgomp/target.c | 8 +++++++- 2 files changed, 31 insertions(+), 20 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index c8ec3c9a7dd..d7a1d87c9ef 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -755,16 +755,19 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async) gomp_remove_var_async (acc_dev, n, aq); else { - int num_mappings = 0; - /* If the target_mem_desc represents a single data mapping, we can - check that it is freed when this splay tree key's refcount - reaches zero. Otherwise (e.g. for a struct mapping with multiple - members), fall back to skipping the test. */ - for (int i = 0; i < n->tgt->list_count; i++) - if (n->tgt->list[i].key) - num_mappings++; + int remaining_mappings = 0; + bool dynamic = goacc_marked_dynamic_p (n->tgt); + if (dynamic) + { + /* For dynamic mappings, we may have more than one live splay + tree in the target_mem_desc's variable list. That's not an + error. */ + for (int i = 0; i < n->tgt->list_count; i++) + if (n->tgt->list[i].key) + remaining_mappings++; + } bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); - assert (num_mappings > 1 || is_tgt_unmapped); + assert ((dynamic && remaining_mappings > 0) || is_tgt_unmapped); } } @@ -1283,17 +1286,19 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_fatal ("synchronize failed"); } } - int num_mappings = 0; - /* If the target_mem_desc represents a single data mapping, we - can check that it is freed when this splay tree key's - refcount reaches zero. Otherwise (e.g. for a struct - mapping with multiple members), fall back to skipping the - test. */ - for (int j = 0; j < n->tgt->list_count; j++) - if (n->tgt->list[j].key) - num_mappings++; + int remaining_mappings = 0; + bool dynamic = goacc_marked_dynamic_p (n->tgt); + if (dynamic) + { + /* For dynamic mappings, we may have more than one live + splay tree in the target_mem_desc's variable list. + That's not an error. */ + for (int j = 0; j < n->tgt->list_count; j++) + if (n->tgt->list[j].key) + remaining_mappings++; + } bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); - assert (num_mappings > 1 || is_tgt_unmapped); + assert ((dynamic && remaining_mappings > 0) || is_tgt_unmapped); } } break; diff --git a/libgomp/target.c b/libgomp/target.c index 9a51e1c70f6..f072e050cc1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1630,6 +1630,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, if (k == NULL) continue; + bool clear_mapping = true; bool do_unmap = false; if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) { @@ -1637,7 +1638,10 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, /* If we only have dynamic references left, mark the tgt_mem_desc appropriately. */ if (k->refcount == k->dynamic_refcount) - goacc_mark_dynamic (k->tgt); + { + goacc_mark_dynamic (k->tgt); + clear_mapping = false; + } } else if (k->refcount == 1) { @@ -1662,6 +1666,8 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, assert (!is_tgt_unmapped || k_tgt != tgt); } + if (clear_mapping) + tgt->list[i].key = NULL; } if (aq)