From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 2A2FB38A9430; Tue, 11 Jan 2022 13:28:10 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 2A2FB38A9430 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: LjeRmKeXijV9K7T9NpN4gUeD7tVGSGM5+xQHhNOhhQqQkBvAyMtRLQf34fQEEoZ7uHpotTxSM1 9b3/rByMlY1Go22aM6wwtsavm/pFL7kuyPUm6JBeX+0GjdkaKIEHXLnn2LlBQS0AZkJvCv1Sb3 KDjn45xxdjtC0KEBVZegAfZ9lxWarfS49droksCatO42qN3NMvTo4xt4vM0JUS9sWNj4DkWAfB NgiYSmSgVnTGgqiZ6cMU8ZSXKIv79K+os6TI3KdgUs0NSLpdAmtILnvxGtqmUbKo/ryXJV0E4z 6HIfVtlHzk89nP7S4yN0/csy X-IronPort-AV: E=Sophos;i="5.88,279,1635235200"; d="scan'208";a="70488969" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 11 Jan 2022 05:28:09 -0800 IronPort-SDR: YQRY154fHa/p6WQFYsVtRcK8xGIU4QQmDjAYGiIgRBCcuykDjDw5/GsZ4dNvG3dWtjgE4DXx+y ///xhhSxfbDeXbbdItKgr/GYYIR83pyhYENV+nQhzw/0yztQDk3cyizC6TRZ/GxMVa/K09EGj5 UcNZBUmTs4fOvfgS4oExzZe49iY/w5hDRbsTE10eu9Vilrkv68DDROBJYstbI912E0b6Mh5cND nYB7dZkQnnDnqbl8rHXIjWiQbL5GIGe9DHNTYtwKOhQHspcemZiFqc3TiNMMztWKmz0X7kTB4d er4= Message-ID: Date: Tue, 11 Jan 2022 14:27:57 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.4.1 Subject: Re: [PATCH] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct Content-Language: en-US To: Jakub Jelinek , Marcel Vollweiler CC: , References: <20211020123844.GU304296@tucnak> <25633d5a-d94d-0231-b626-97d28159237f@codesourcery.com> <20220111115324.GV2646553@tucnak> From: Tobias Burnus In-Reply-To: <20220111115324.GV2646553@tucnak> Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-5.6 required=5.0 tests=BAYES_00, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=no autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: fortran@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Fortran mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 11 Jan 2022 13:28:11 -0000 Hi Jakub, hi all, let me quickly comment on 'has_device_addr' with Fortran arrays and with an array section (i.e. regarding your comment quoted at the very bottom of this email). Unfortunately, the OpenMP specification is rather unclear what has_device_addr means for C/C++ array sections and in general for Fortran, especially when arrays, allocatables/pointers, and type parameters like nonconst string lengths are involved. Thus, I opened a spec issue =E2=80=93 after some discussions (lang-spec meeting, C++/affinity (=E2=86=92 Fortran) meeting), it starts to converge: https://github.com/OpenMP/spec/issues/3180 If I understood it correctly, for C/C++, using has_device_addr with an array section implies firstprivate, while it does not without array section. For Fortran, the discussion converged to * array descriptor must be the same at beginning/end of target region, i.e: - if allocated before target region, allocation may not be changed inside the target region. If unallocated, it can be allocated but must be deallocated before the target region is left. - if a pointer, pointer association must be the same at the end of the target region as when it was entered (unmodified or reset) * meta data (array bounds, alloc status, len type parameters like char length) are accessible on the host - and are then made accessible to the device (firstprivate, direct access on shared memory or whatever means. Due to the conditions above, it is indistinguishable). However, the discussion has not yet fully settled and spec updates are still needed. Side remark: I note that use_device_addr permits array sections, but GCC does not support them yet. (Useful when doing a partial map of an array + 'omp data use_device_addr()' on the partially mapped array.) I have not checked what Marcel's implementation does and whether that's compatible with the incomplete OpenMP 5.{1,2} spec and the on-going discussion of/in OpenMP lang-spec Issue 3180. On 11.01.22 12:53, Jakub Jelinek via Fortran wrote: >> +++ b/gcc/fortran/trans-openmp.c >> ... >> + /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize t= he >> + descriptor such that the bounds are available; its data comp= onent >> + is unmodified; it is handled as device address inside target= . */ ... > Not sure about the above, > >> --- a/gcc/gimplify.c >> +++ b/gcc/gimplify.c >> ... > but this looks weird. > If decl after stripping the ARRAY_REFs is a var with pointer type, sure, > firstprivatizing it is the way to go. > But it can be also a variable with ARRAY_TYPE, can't it? Something like: > int a[64]; > #pragma omp target data map(a) use_device_addr(a) > { > #pragma omp target has_device_addr(a[3:16]) > a[3] =3D 1; > } > and in this case firstprivatization of a looks wrong. use_device_addr > should replace (but only at omp-low.c time I think) a used in the block > with the remapped a (i.e. *device_address_of_a). > Or perhaps it could be a non-static data member with array type > inside of a C++ method. Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955