From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 6E0C5385383E for ; Mon, 26 Jul 2021 19:28:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 6E0C5385383E 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: m/eLQXKvebmP15dEQvO2hFBhDppOwIwVRXJq7VV5HeAgsxsnT4hDftOEGZ716p3UnI8KtpUURU s6oxnAJTxDfnBdMNJwLSk04o+IShX9pCCxDXDnPb9ye5XPCq5VQuY1SL7HpYEDcPb66aa3RliB 7f4VbObdp7pozuWPh1OXhN+uTpUHgxC7qnQV+pluEYMAf001/NzAgh+4Ce9rfN/u1RQxZkNbAB 5anyYCgXREwKXxyYGNiAtOm1tcQuIf/7ZCTYo2XDSumsIBh6/JYWDVUrY4sEdZo6DcezmBjdVn Uwii+8LF2VVVQ9jWHDunEwgZ X-IronPort-AV: E=Sophos;i="5.84,270,1620720000"; d="scan'208";a="64116408" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 26 Jul 2021 11:28:24 -0800 IronPort-SDR: kTnUlqoV5PM2axoVtSrDp1TFqzQm211ycFVawtkFtToft/fJu6CXUZ7+4/SHSCRg4DxW6MMiKS OIyazDqM01RR2KdzvJM1rzG0BIdIs0SUIcFPzSHgydv+B/w3TQHWsybniufUjRLuDkp2tJhJ9D ysHw3moePAzCbx+jgKq0jE+9E7LDB5GpOVnfdttOm2ssE04XeRAIWmQhXhA7va0HfKyxn3H8iB eIG/YqvkbQfXA0rhp35eFq24mvGiMnKH5l0MWQ7G8UjKoZ1auqFFsWAR7U1kch4WsaiR0idwVG oq8= Subject: Re: [WIP, OpenMP] OpenMP metadirectives support To: Jakub Jelinek CC: GCC Patches References: <8d413974-0068-3a31-6ae5-d36c1be06d06@codesourcery.com> <20210726142902.GW2380545@tucnak> From: Kwok Cheung Yeung Message-ID: <0000a35d-75a5-3067-b59c-b4f5bde0ea58@codesourcery.com> Date: Mon, 26 Jul 2021 20:28:16 +0100 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:78.0) Gecko/20100101 Thunderbird/78.12.0 MIME-Version: 1.0 In-Reply-To: <20210726142902.GW2380545@tucnak> Content-Type: text/plain; charset="utf-8"; format=flowed Content-Language: en-GB Content-Transfer-Encoding: 7bit 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=-6.4 required=5.0 tests=BAYES_00, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: Mon, 26 Jul 2021 19:28:27 -0000 Hello Thanks for your reply. On 26/07/2021 3:29 pm, Jakub Jelinek wrote: > On Fri, Jul 09, 2021 at 12:16:15PM +0100, Kwok Cheung Yeung wrote: >> 3) In the OpenMP examples (version 5.0.1), section 9.7, the example >> metadirective.3.c does not work as expected. >> >> #pragma omp declare target >> void exp_pi_diff(double *d, double my_pi){ >> #pragma omp metadirective \ >> when( construct={target}: distribute parallel for ) \ >> default( parallel for simd) >> ... >> int main() >> { >> ... >> #pragma omp target teams map(tofrom: d[0:N]) >> exp_pi_diff(d,my_pi); >> ... >> exp_pi_diff(d,my_pi); > > The spec says in this case that the target construct is added to the > construct set because of the function appearing in between omp declare target > and omp end declare target, so the above is something that resolves > statically to distribute parallel for. > It is true that in OpenMP 5.1 the earlier > For functions within a declare target block, the target trait is added to the beginning of the > set as c 1 for any versions of the function that are generated for target regions so the total size > of the set is increased by 1. > has been mistakenly replaced with: > For device routines, the target trait is added to the beginning of the set as c 1 for any versions of > the procedure that are generated for target regions so the total size of the set is increased by 1. > by that has been corrected in 5.2: > C/C++: > For functions that are declared in a code region that is delimited by a declare target directive and > its paired end directive, the target trait is added to the beginning of the set as c 1 for any target > variants that result from the directive so the total size of the set is increased by one. > Fortran: > If a declare target directive appears in the specification part of a procedure or in the > specification part of a procedure interface body, the target trait is added to the beginning of the > set as c 1 for any target variants that result from the directive so the total size of the set is > increased by one. > > So, it is really a static decision that can be decided already during > parsing. In Section 1.2.2 of the OpenMP TR10 spec, 'target variant' is defined as: A version of a device routine that can only be executed as part of a target region. So isn't this really saying the same thing as the previous versions of the spec? The target trait is added to the beginning of the construct set _for any target variants_ that result from the directive (implying that it shouldn't be added for non-target variants). In this example, the same function exp_pi_diff is being used in both a target and non-target context, so shouldn't the metadirective resolve differently in the two contexts, independently of the function being declared in a 'declare target' block? If not, there does not seem to be much point in that example (in section 9.7 of the OpenMP Examples v5.0.1). From reading the spec, I infer that they expect the device and non-device versions of a function with 'declare target' to be separate, but that is not currently the case for GCC - on the host compiler, the same version of the function gets called in both target and non-target regions (though in the target region case, it gets called indirectly via a compiler-generated function with a name like main._omp_fn.0). The offload compiler gets its own streamed version, so there is no conflict there - by definition, its version must be in a target context. Thanks, Kwok