Hi Chung-Lin!
On 2023-06-06T23:11:55+0800, Chung-Lin Tang <chunglin.t...@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 <clt...@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
