From: Kwok Cheung Yeung <kcy@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>
Cc: GCC Patches <gcc-patches@gcc.gnu.org>
Subject: Re: [WIP, OpenMP] OpenMP metadirectives support
Date: Mon, 26 Jul 2021 20:28:16 +0100 [thread overview]
Message-ID: <0000a35d-75a5-3067-b59c-b4f5bde0ea58@codesourcery.com> (raw)
In-Reply-To: <20210726142902.GW2380545@tucnak>
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
next prev parent reply other threads:[~2021-07-26 19:28 UTC|newest]
Thread overview: 29+ messages / expand[flat|nested] mbox.gz Atom feed top
2021-07-09 11:16 Kwok Cheung Yeung
2021-07-26 11:38 ` Kwok Cheung Yeung
2021-07-26 14:29 ` Jakub Jelinek
2021-07-26 19:28 ` Kwok Cheung Yeung [this message]
2021-07-26 19:56 ` Jakub Jelinek
2021-07-26 21:19 ` Kwok Cheung Yeung
2021-07-26 21:23 ` Jakub Jelinek
2021-07-26 21:27 ` Kwok Cheung Yeung
2022-01-28 16:33 ` [PATCH] openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly Kwok Cheung Yeung
2021-12-10 17:27 ` [WIP, OpenMP] OpenMP metadirectives support Kwok Cheung Yeung
2021-12-10 17:29 ` [PATCH 0/7] openmp: " Kwok Cheung Yeung
2021-12-10 17:31 ` [PATCH 1/7] openmp: Add C support for parsing metadirectives Kwok Cheung Yeung
2022-02-18 21:09 ` [PATCH] openmp: Improve handling of nested OpenMP metadirectives in C and C++ (was: Re: [PATCH 1/7] openmp: Add C support for parsing metadirectives) Kwok Cheung Yeung
2022-02-18 21:26 ` [og11][committed] openmp: Improve handling of nested OpenMP metadirectives in C and C++ Kwok Cheung Yeung
2022-05-27 17:44 ` [PATCH 1/7] openmp: Add C support for parsing metadirectives Jakub Jelinek
2021-12-10 17:33 ` [PATCH 2/7] openmp: Add middle-end support for metadirectives Kwok Cheung Yeung
2022-05-30 10:54 ` Jakub Jelinek
2021-12-10 17:35 ` [PATCH 3/7] openmp: Add support for resolving metadirectives during parsing and Gimplification Kwok Cheung Yeung
2022-05-30 11:13 ` Jakub Jelinek
2021-12-10 17:36 ` [PATCH 4/7] openmp: Add support for streaming metadirectives and resolving them after LTO Kwok Cheung Yeung
2022-05-30 11:33 ` Jakub Jelinek
2021-12-10 17:37 ` [PATCH 5/7] openmp: Add C++ support for parsing metadirectives Kwok Cheung Yeung
2022-05-30 11:52 ` Jakub Jelinek
2021-12-10 17:39 ` [PATCH 6/7] openmp, fortran: Add Fortran " Kwok Cheung Yeung
2022-02-14 15:09 ` Kwok Cheung Yeung
2022-02-14 15:17 ` Kwok Cheung Yeung
2021-12-10 17:40 ` [PATCH 7/7] openmp: Add testcases for metadirectives Kwok Cheung Yeung
2022-05-27 13:42 ` Jakub Jelinek
2022-01-24 21:28 ` [PATCH] openmp: Metadirective patch fixes Kwok Cheung Yeung
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=0000a35d-75a5-3067-b59c-b4f5bde0ea58@codesourcery.com \
--to=kcy@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).