From patchwork Fri Jun 5 16:36:54 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1304242 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 49dpG95nGWz9sT6 for ; Sat, 6 Jun 2020 02:37:20 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 3B47B3890417; Fri, 5 Jun 2020 16:37:17 +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 AC72B3851C28 for ; Fri, 5 Jun 2020 16:37:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AC72B3851C28 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: gqyMGp00uSb6YFQ2SqWGAOKlUndv+ldBB/sNMrcy9n1rtfeqRhh+nmcAAwYT7fe/zdJy9sKF3b uzLV814qYqpdGunQtM6SUIhTp1Dvu3YWU52X76Wvz5iSe+j+8QlzIUDcqqzQVHSM6TbdHlYVEt 1CzXJE7JbtMeQV3JMf889J1vNILCWFeWO2RH2pIjs3hrKXgsKJ/Y1aoFNbBM2f16GtCe1mF8Ju bq2udWCqb8U7QNgOTOUL6fciDOHmyL8A6bLY6ussqug//+meIlLTTDkgsvQiAxof7BvhLQQVrJ PZA= X-IronPort-AV: E=Sophos;i="5.73,477,1583222400"; d="scan'208,223";a="51621337" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 05 Jun 2020 08:37:04 -0800 IronPort-SDR: mcD7qKDcx25HU3T/uwW6k1MONWrYfiDv8eJ0n2SN9r6BtQFWEpcWYRo9JnvQIv5OaX7x99yndW 9xNdh4VLNiEgnMJ7ns7dy0WXCnlO4BasgwS7BoMa/KQwWsPHRHDKNyT27CYMpEwrvyGtu6GoBI 0fmazgba/J9qGVwWgpNSPB4KSNha9Avsa8wVKVvlxSTfxjU03V9rGr+nDLAPgtNwbiAq903+nl UnAL7DHjbNWvL3CXYWPLzs4ACbXiprdOJgbE6s5z20edGk3Yh1cDnPEHCrcxtrUgBV9nffS6mx BEg= From: Thomas Schwinge To: Julian Brown , Subject: [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) In-Reply-To: <87r1vf7xr4.fsf@euler.schwinge.homeip.net> References: <65540b92dff74db1f15af930f87f7096d03e7efe.1576648001.git.julian@codesourcery.com> <87r1vf7xr4.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Fri, 5 Jun 2020 18:36:54 +0200 Message-ID: <87v9k5qxjt.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Spam-Status: No, score=-11.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, 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 Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2020-05-20T11:37:35+0200, I wrote: > Moving this over, from the "Fix component mappings with derived types for > OpenACC" thread, > , where > you propose to change this 'GOMP_MAP_STRUCT' handling code: > > On 2019-12-17T22:03:47-0800, Julian Brown wrote: >> --- a/libgomp/oacc-mem.c >> +++ b/libgomp/oacc-mem.c > >> @@ -1075,6 +1119,39 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, >> 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 (finalize) >> + { >> + str->refcount -= str->virtual_refcount; >> + str->virtual_refcount = 0; >> + } >> + if (str->virtual_refcount > 0) >> + { >> + str->refcount--; >> + str->virtual_refcount--; >> + } >> + else if (str->refcount > 0) >> + str->refcount--; >> + if (str->refcount == 0) >> + gomp_remove_var_async (acc_dev, str, aq); >> + } >> + } >> + i += elems; >> + } >> + break; >> + >> default: >> gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x", >> kind); > > ... into an "empty 'case GOMP_MAP_STRUCT:' as a no-op [...] > Further citing myself, > : > > | Is my understanding correct that "fixed" GCC won't generate such > | 'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty > | 'case GOMP_MAP_STRUCT:' only remains in here for backwards compatibility? > | In this case, please add a comment to the code, stating this. > > My guess was wrong: running a quick test, I do see that we still generate > 'GOMP_MAP_STRUCT' for OpenACC unmap: [...] > But one step back. Why generate 'GOMP_MAP_STRUCT' for unmap, if we then > just skip it in libgomp handling? Cross checking, OpenMP > 'libgomp/target.c:gomp_exit_data' also doesn't expect to see any > 'GOMP_MAP_STRUCT'. > > For example, [...] > > I haven't studied 'GOMP_MAP_STRUCT' in detail (so it may be easy to prove > me wrong), but for OpenACC 'exit data' etc., these 'map(struct:[...])' > seem "pointless" to me, given that in libgomp we (intend to) just skip > them? > > Well, and quickly I find that's exactly what OpenMP 'target exit data' is > doing, and doing the same for OpenACC 'exit data': [...] > ..., 'deep-copy-7.c.005t.gimple' gets simplified as expected: [...] > [...] and so 'GOMP_MAP_STRUCT' handling could be removed from > 'libgomp/oacc-mem.c:goacc_exit_data_internal' (and 'default:' > 'gomp_fatal' would then catch any such cases). > > But of course, given that GCC 10.1 now does generate these > 'GOMP_MAP_STRUCT's, we do have to support them in one way or another... You've picked that up in your patch submission. Thanks for the test case you've included, I changed that one just a bit: '-fdump-tree-gimple' scanning (as that's where the processing happens), removed the 'stdlib.h' dependency, made the 'struct' (and thus the 'GOMP_MAP_STRUCT' that's generated) a bit more interesting. I've pushed "[OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings" to master branch in commit 1afc4672561a41dfbf4e3f2c1f35f7a5b7a20339, and releases/gcc-10 branch in commit 27e985479e931771472663cad34f8b99c6f62627, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter From 27e985479e931771472663cad34f8b99c6f62627 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 20 May 2020 10:56:55 +0200 Subject: [PATCH] [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings These are not itself necessary for OpenACC 'exit data' directives, and are skipped over (now) in libgomp. We might as well not emit them to start with, in line with the equivalent OpenMP directive. We keep the no-op handling in libgomp for the reason of backward compatibility. gcc/ * gimplify.c (gimplify_adjust_omp_clauses): Remove 'GOMP_MAP_STRUCT' mapping from OpenACC 'exit data' directives. gcc/testsuite/ * c-c++-common/goacc/struct-enter-exit-data-1.c: New file. libgomp/ * oacc-mem.c (goacc_exit_data_internal) : Explain special handling. Co-Authored-By: Julian Brown (cherry picked from commit 1afc4672561a41dfbf4e3f2c1f35f7a5b7a20339) --- gcc/gimplify.c | 3 ++- .../goacc/struct-enter-exit-data-1.c | 27 +++++++++++++++++++ libgomp/oacc-mem.c | 5 ++-- 3 files changed, 32 insertions(+), 3 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 0bac99002102..4b3306cfc258 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10408,7 +10408,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && code == OMP_TARGET_EXIT_DATA) + && (code == OMP_TARGET_EXIT_DATA + || code == OACC_EXIT_DATA)) remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c new file mode 100644 index 000000000000..df405e448b28 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -0,0 +1,27 @@ +/* Check 'GOMP_MAP_STRUCT' mapping, and in particular that it gets removed from + OpenACC 'exit data' directives. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +struct str { + int a; + int *b; + int *c; + int d; + int *e; + int f; +}; + +#define N 1024 + +void +test (int *b, int *c, int *e) +{ + struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 }; + +#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */ + +#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */ +} diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 1e3685a073da..936ae649dd93 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1181,8 +1181,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, case GOMP_MAP_STRUCT: /* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular processing - for all its entries. TODO: don't generate these no-op - 'GOMP_MAP_STRUCT's. */ + for all its entries. This special handling exists for GCC 10.1 + compatibility; afterwards, we're not generating these no-op + 'GOMP_MAP_STRUCT's anymore. */ break; default: -- 2.26.2