[OpenACC,2.7] Implement default clause support for data constructs

  Hi Thomas,
this patch implements the OpenACC 2.7 addition of default(none|present) support
for data constructs.

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

Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect
no surprises). Is this okay for trunk?

Thanks,
Chung-Lin
gcc/c/ChangeLog:

	* c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/cp/ChangeLog:

	* parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/fortran/ChangeLog:

	* openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT.

gcc/ChangeLog:

	* gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.
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).

gcc/c/ChangeLog:

	* c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/cp/ChangeLog:

	* parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/fortran/ChangeLog:

	* openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT.

gcc/ChangeLog:

	* gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.
---
 gcc/c/c-parser.cc                             |  1 +
 gcc/cp/parser.cc                              |  1 +
 gcc/fortran/openmp.cc                         |  3 ++-
 gcc/gimplify.cc                               | 20 +++++++++++++++++++
 gcc/testsuite/c-c++-common/goacc/default-3.c  | 15 +++++++++++++-
 gcc/testsuite/c-c++-common/goacc/default-5.c  | 18 +++++++++++++++--
 gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 15 ++++++++++++++
 gcc/testsuite/gfortran.dg/goacc/default-5.f   | 17 ++++++++++++++--
 8 files changed, 84 insertions(+), 6 deletions(-)
  

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
  

Hi Thomas,

On 2023/6/23 6:47 PM, Thomas Schwinge wrote:
>> +
>>    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?

I'm not sure if this is conceptually simpler, but using 'oacc_default_kind'
is definitely faster computationally :)

However, as you mention below...

> 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.)

Noticed that we will need to track the actually lexically enclosing OpenACC construct
with the user set default-clause somewhere in 'ctx', in order to satisfy the current
diagnostics in oacc_default_clause().

(the UNKNOWN_LOCATION for the internally created default-clause probably doesn't
matter, that one is just for reminder in internal dumps, probably never plays role
in user diagnostics)

> 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.

Okay, will add more testcases.

Thanks,
Chung-Lin