From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 6A29D3952007 for ; Wed, 20 May 2020 09:37:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 6A29D3952007 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: 0eJrohbHRGE9/kH1xsy20212L2uSoiB29gXBA+5HBAg9sYxyiZNseazw7AkKXPuV8DAF4AB2Qd 50MFuA1u6tQU7U3YBGmnBlegk2DNt1OpbukoQ37THt4WyXhA5plBLhF8aKGpx7dxEP1BvnRwTo cXWQEHW9ayN75AWe6PzCiRaoMubooJQ825yUUJU18SzdQYXavZ8O9Ysv88Ryy4YUg7xTDQpTUk fq2yx0gYVS/bGyKEAGJe8vVu5IPTrLx3XVgXPWBAEsd+cbBz4XxnwD1TOiK2B7KpzgD8SplzGn tJ8= X-IronPort-AV: E=Sophos;i="5.73,413,1583222400"; d="scan'208";a="48960777" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 20 May 2020 01:37:44 -0800 IronPort-SDR: qmjVzOJ9owpupJZJ9pZfTgfjIGXaZO/ChYWh+LaLv7BYaHyefaQkAL5k24heBIgW5MORSKvNWR R4/3ib/eKD5omYHMBKIJpsln0ZoRRD/+v7Kvfq7fEpkoYDgGj6DUwg5o3rYzLrspC9mZzKAJsM HRgMC1L+kulNKNBOXUGODppqBUznRiOdsUQiPB8SXHgfySB3zXp8daQfbYa2CqbgM9FzE1e2Xr CjRWU+lIIfr6MXpjllSzdRNEBAZ23yP6U9aVUHQ1zFe9FSHsozPHfm9fJBPuive6oEXmoAU20m 6RI= From: Thomas Schwinge To: Julian Brown CC: , Jakub Jelinek Subject: Re: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts In-Reply-To: <65540b92dff74db1f15af930f87f7096d03e7efe.1576648001.git.julian@codesourcery.com> References: <65540b92dff74db1f15af930f87f7096d03e7efe.1576648001.git.julian@codesourcery.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 20 May 2020 11:37:35 +0200 Message-ID: <87r1vf7xr4.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-Spam-Status: No, score=-7.3 required=5.0 tests=BAYES_00, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=no 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: , X-List-Received-Date: Wed, 20 May 2020 09:37:48 -0000 Hi Julian! 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 =3D sizes[i]; > + for (int j =3D 1; j <=3D elems; j++) > + { > + struct splay_tree_key_s k; > + k.host_start =3D (uintptr_t) hostaddrs[i + j]; > + k.host_end =3D k.host_start + sizes[i + j]; > + splay_tree_key str; > + str =3D splay_tree_lookup (&acc_dev->mem_map, &k); > + if (str) > + { > + if (finalize) > + { > + str->refcount -=3D str->virtual_refcount; > + str->virtual_refcount =3D 0; > + } > + if (str->virtual_refcount > 0) > + { > + str->refcount--; > + str->virtual_refcount--; > + } > + else if (str->refcount > 0) > + str->refcount--; > + if (str->refcount =3D=3D 0) > + gomp_remove_var_async (acc_dev, str, aq); > + } > + } > + i +=3D 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, so that we don't run into 'default:' case 'goacc_exit_data_internal UNHANDLED kind'" (my words/interpretation). 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: --- libgomp/oacc-mem.c +++ libgomp/oacc-mem.c @@ -1163,6 +1165,7 @@ goacc_exit_data_internal break; case GOMP_MAP_STRUCT: + assert (!"GOMP_MAP_STRUCT"); break; default: gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x"= , kind); ... regresses: - 'libgomp.oacc-c-c++-common/deep-copy-7.c' - 'libgomp.oacc-c-c++-common/deep-copy-8.c' - 'libgomp.oacc-fortran/classtypes-2.f95' - 'libgomp.oacc-fortran/deep-copy-4.f90' - 'libgomp.oacc-fortran/deep-copy-5.f90' - 'libgomp.oacc-fortran/deep-copy-6.f90' - 'libgomp.oacc-fortran/derivedtype-2.f95' | Otherwise, | please add a comment why "do nothing" is appropriate for | 'GOMP_MAP_STRUCT'. I suppose we still need to unmap the "'GOMP_MAP_STRUCT' components", but can do that individually, outside of the 'GOMP_MAP_STRUCT' context. That'd then also explain... | In particular, for both scenarios, why we don't need | to skip the following 'sizes[i]' mappings? ... this question. 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, 'libgomp.oacc-c-c++-common/deep-copy-7.c': #pragma acc exit data copyout(v.b[:n]) finalize #pragma acc exit data delete(v.a) 'deep-copy-7.c.004t.original': #pragma acc exit data finalize map(from:*v.b [len: (sizetype) n * 4]) m= ap(attach_detach:v.b [bias: 0]); #pragma acc exit data map(release:v.a); 'deep-copy-7.c.005t.gimple': #pragma omp target oacc_enter_exit_data finalize map(struct:v [len: 1])= map(delete:v.b [len: 8]) map(force_from:*_15 [len: _14]) map(force_detach:= v.b [bias: 0]) #pragma omp target oacc_enter_exit_data map(struct:v [len: 1]) map(rele= ase:v.a [len: 4]) 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': --- gcc/gimplify.c +++ gcc/gimplify.c @@ -10406,7 +10406,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,= gimple_seq body, tree *list_p, } } else if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT - && code =3D=3D OMP_TARGET_EXIT_DATA) + && (code =3D=3D OMP_TARGET_EXIT_DATA + || code =3D=3D OACC_EXIT_DATA)) remove =3D true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) !=3D INTEGER_CST ..., 'deep-copy-7.c.005t.gimple' gets simplified as expected: - #pragma omp target oacc_enter_exit_data finalize map(struct:v = [len: 1]) map(delete:v.b [len: 8]) map(force_from:*_15 [len: _14]) map(forc= e_detach:v.b [bias: 0]) + #pragma omp target oacc_enter_exit_data finalize map(delete:v.= b [len: 8]) map(force_from:*_15 [len: _14]) map(force_detach:v.b [bias: 0]) - #pragma omp target oacc_enter_exit_data map(struct:v [len: 1])= map(release:v.a [len: 4]) + #pragma omp target oacc_enter_exit_data map(release:v.a [len: = 4]) ..., and all the "'assert' regressions" mentioned above again disappear, 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... Gr=C3=BC=C3=9Fe Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter