public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Thomas Schwinge <thomas@codesourcery.com>
To: Chung-Lin Tang <cltang@codesourcery.com>
Cc: <gcc-patches@gcc.gnu.org>, Catherine Moore <clm@codesourcery.com>
Subject: Re: [PATCH, OpenACC 2.7] Implement default clause support for data constructs
Date: Fri, 23 Jun 2023 12:47:24 +0200	[thread overview]
Message-ID: <87y1kacv5v.fsf@euler.schwinge.homeip.net> (raw)
In-Reply-To: <f2b7a229-0ec8-5404-fea2-4612ffb73bf2@siemens.com>

Hi Chung-Lin!

On 2023-06-06T23:11:55+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote:
> this patch implements the OpenACC 2.7 addition of default(none|present) support
> for data constructs.

Thanks!

It wasn't clear to me what is supposed to happen, for example, for:

    #pragma acc data default(none)
    {
      #pragma acc data // no 'default' clause
      {
        #pragma acc parallel

Specifically, does the "no 'default' clause" inner 'data' construct
invalidate the 'default(none)' clause of the outer 'data' construct?

In later revisions of the OpenACC specification, wording for 'default'
clause etc. generally has been changed; for example, OpenACC 3.3,
2.6.2 "Variables with Implicitly Determined Data Attributes" defines:

    *Visible 'default' clause*: The nearest 'default' clause appearing on the compute construct or a lexically containing 'data' construct.

Therefore, in the example above, the 'default(none)' still holds.

> Apart from adjusting the front-ends for allowed clauses masks (for acc data),
> mostly implemented in gimplify.

ACK ('s%mostly%%') -- but a little bit differently, please:

> From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 Mon Sep 17 00:00:00 2001
> From: Chung-Lin Tang <cltang@codesourcery.com>
> Date: Tue, 6 Jun 2023 03:46:29 -0700
> Subject: [PATCH 2/2] OpenACC 2.7: default clause support for data constructs
>
> This patch implements the OpenACC 2.7 addition of default(none|present) support
> for data constructs.
>
> Now, specifying "default(none|present)" on a data construct turns on same
> default clause behavior for all enclosed compute constructs (which don't
> already themselves have a default clause).

Please say "lexically enclosed" -- it's that only, not any dynamic
extent.

> --- a/gcc/gimplify.cc
> +++ b/gcc/gimplify.cc
> @@ -225,6 +225,7 @@ struct gimplify_omp_ctx
>    vec<tree> loop_iter_var;
>    location_t location;
>    enum omp_clause_default_kind default_kind;
> +  enum omp_clause_default_kind oacc_data_default_kind;
>    enum omp_region_type region_type;
>    enum tree_code code;
>    bool combined_loop;
> @@ -459,6 +460,8 @@ new_omp_context (enum omp_region_type region_type)
>      c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
>    else
>      c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
> +  if (gimplify_omp_ctxp)
> +    c->oacc_data_default_kind = gimplify_omp_ctxp->oacc_data_default_kind;
>    c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
>    c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
>    c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;

Instead of adding a new 'oacc_data_default_kind' to 'gimplify_omp_ctx',
let's please do this the other way round:

> @@ -12050,6 +12053,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>
>       case OMP_CLAUSE_DEFAULT:
>         ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);

Here, we already preserve 'default' for whichever OMP construct.

> +       if (code == OACC_DATA)
> +         ctx->oacc_data_default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
>         break;
>
>       case OMP_CLAUSE_INCLUSIVE:
> @@ -12098,6 +12103,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>       list_p = &OMP_CLAUSE_CHAIN (c);
>      }
>
> +  if ((code == OACC_PARALLEL
> +       || code == OACC_KERNELS
> +       || code == OACC_SERIAL)
> +      && ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED
> +      && ctx->oacc_data_default_kind != OMP_CLAUSE_DEFAULT_UNSPECIFIED)
> +    {
> +      ctx->default_kind = ctx->oacc_data_default_kind;
> +
> +      /* Append actual default clause on compute construct. Not really needed
> +      for omp_notice_variable to work properly, but for debug dump files.  */
> +      c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEFAULT);
> +      OMP_CLAUSE_DEFAULT_KIND (c) = ctx->oacc_data_default_kind;
> +      *list_p = c;
> +    }
> +
>    ctx->clauses = *orig_list_p;
>    gimplify_omp_ctxp = ctx;
>  }

Instead of this, in 'gimplify_omp_workshare', before the
'gimplify_scan_omp_clauses' call, do something like:

    if ((ort & ORT_ACC)
        && !omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_DEFAULT))
      {
        /* Determine effective 'default' clause for OpenACC compute construct.  */
        for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
          {
            if (ctx->region_type == ORT_ACC_DATA
                && ctx->default_kind != OMP_CLAUSE_DEFAULT_SHARED)
              {
                [Append actual default clause on compute construct.]
                break;
              }
          }
      }

That seems conceptually simpler to me?

For the 'build_omp_clause', does using 'ctx->location' instead of
'UNKNOWN_LOCATION' help diagnostics in any way?  Like if we add in
'gcc/gimplify.cc:oacc_default_clause',
'if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)' another 'inform' to
point to the 'data' construct's 'default' clause?  (But not sure if
that's easily done; otherwise don't.)

Similar to the ones you've already got, please also add a few test cases
for nested 'default' clauses, like:

    #pragma acc data // no vs. 'default(none)' vs. 'default(present)'
    {
      #pragma acc data // no vs. same vs. different 'default' clause
      {
        #pragma acc data // no vs. same vs. different 'default' clause
        {
          #pragma acc parallel

Similarly, test cases where 'default' on the compute construct overrides
'default' of an outer 'data' construct.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

  reply	other threads:[~2023-06-23 10:47 UTC|newest]

Thread overview: 5+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2023-06-06 15:11 Chung-Lin Tang
2023-06-23 10:47 ` Thomas Schwinge [this message]
2023-07-14 10:34   ` Chung-Lin Tang
2023-08-15 14:35     ` [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs) Thomas Schwinge
2023-08-01 15:35 [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs Chung-Lin Tang

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=87y1kacv5v.fsf@euler.schwinge.homeip.net \
    --to=thomas@codesourcery.com \
    --cc=clm@codesourcery.com \
    --cc=cltang@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    /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).