public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Kwok Cheung Yeung <kcy@codesourcery.com>
To: GCC Patches <gcc-patches@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>
Subject: Re: [WIP, OpenMP] OpenMP metadirectives support
Date: Mon, 26 Jul 2021 12:38:23 +0100	[thread overview]
Message-ID: <eb43dd4e-d6c0-ecdd-8df5-3a00e2058459@codesourcery.com> (raw)
In-Reply-To: <8d413974-0068-3a31-6ae5-d36c1be06d06@codesourcery.com>

Ping? Does anyone have any opinions on how this issue should be resolved?

On 09/07/2021 12:16 pm, 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);
> 
> In the first call to exp_pi_diff in an '#pragma omp target' construct, the 
> metadirective is expected to expand to 'distribute parallel for', but in the 
> second (without the '#pragma omp target'), it should expand to 'parallel for simd'.
> 
> During OMP expansion of the 'omp target', it creates a child function that calls 
> exp_pi_diff:
> 
> __attribute__((omp target entrypoint))
> void main._omp_fn.0 (const struct .omp_data_t.12 & restrict .omp_data_i)
> {
>    ...
>    <bb 4> :
>    __builtin_GOMP_teams (0, 0);
>    exp_pi_diff (d.13, my_pi);
> 
> This is not a problem on the offload compiler (since by definition its copy of 
> exp_pi_diff must be in a 'target'), but if the host device is used, the same 
> version of exp_pi_diff is called in both target and non-target contexts.
> 
> What would be the best way to solve this? Offhand, I can think of two solutions:
> 
> (a) Recursively go through all functions that can be reached via a target region 
> and create clones for each, redirecting all function calls in the clones to the 
> new cloned versions. Resolve the metadirectives in the clones and originals 
> separately.
> 

Maybe this could be done at the same time as when marking functions implicitly 
'declare target'? It seems a lot of work for one special case though...

> (b) Make the construct selector a dynamic selector when OpenMP 5.1 metadirective 
> support is implemented. Keep track of the current construct list every time an 
> OpenMP construct is entered or exited, and make the decision at runtime.
> 

I think this would be easier to implement at runtime (assuming that the 
infrastructure for OpenMP 5.1 was already in place) since this a host-side 
issue, but it probably goes against the intent of the specification, given that 
the 'construct' selector set appeared in the 5.0 specification before dynamic 
replacements became available.

Thanks

Kwok

  reply	other threads:[~2021-07-26 11:38 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 [this message]
2021-07-26 14:29 ` Jakub Jelinek
2021-07-26 19:28   ` Kwok Cheung Yeung
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=eb43dd4e-d6c0-ecdd-8df5-3a00e2058459@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).