[PR90742] OpenACC/OpenMP target offloading: Fortran 'allocatable' scalars in 'firstprivate' clauses
Thomas Schwinge
thomas@codesourcery.com
Mon Oct 7 09:29:00 GMT 2019
Hi!
Jakub, ping -- and/or: Kwok, Tobias, as you recently worked through that
code for related issues (Fortran optional arguments), do you happen to
have any comments?
On 2019-06-07T16:01:29+0200, I wrote:
> As I had mentioned in the PR...
>
> On Tue, 7 Aug 2018 14:55:07 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
>> This patch
>
> ... would be one component for fixing <https://gcc.gnu.org/PR90742>
> "OpenACC/OpenMP target offloading: Fortran 'allocatable' scalars in
> 'firstprivate' clauses".
>
> (Also, as mentioned there, such changes have been submitted already, a
> few times, muddled into other changes. So, thanks, that this also got
> submitted separately, to address just this one issue.)
>
>> updates the way that lower_omp_target uses firstprivate
>> pointers in OpenACC offloaded regions. On host side, when preparing
>> firstprivate data mapping for pointer type objects, not to be confused
>> with GOMP_MAP_FIRSTPRIVATE_POINTER, the compiler passes passes the
>> address of the value being pointed to and not the address of the pointer
>> itself to the runtime. Correspondingly, on the device side, the compiler
>> generates to code to dereference the remapped pointer once to copy the
>> data to a local buffer.
>>
>> While this behavior looks like it would break things, it will not affect
>> C or C++ data mappings, because those languages transfer pointers via
>> GOMP_MAP_FIRSTPRIVATE_POINTER.
>
> Not with current GCC sources, as I should eventually find out, which are
> still missing another patch or two, or three, or more.
>
>> In addition, this will not cause
>> problems with array types, because the default remapping rules for
>> OpenACC is to transfer them in via copy. Besides it really doesn't
>> make sense to allow arrays to be transferred in via firstprivate
>> because that would use up a lot of memory on the accelerator.
>
> (Huh, but the latter ought to be supported nevertheless, as far as I
> understand? Anyway, that'll be for later.)
>
>> Is this OK for trunk? I bootstrapped and regtested it for x86_64 with
>> nvptx offloading.
>
> The patch, as proposed, does introduce regressions.
>
>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -7643,15 +7643,21 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
>> if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
>> {
>> gcc_assert (is_gimple_omp_oacc (ctx->stmt));
>> - if (omp_is_reference (new_var)
>> - && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE)
>> + if (omp_is_reference (new_var))
>> {
>> /* Create a local object to hold the instance
>> value. */
>> - tree type = TREE_TYPE (TREE_TYPE (new_var));
>> + tree type = TREE_TYPE (new_var);
>> + /* Pointer types are mapped onto the device via a
>> + single level of indirection. */
>> + if (TREE_CODE (type) != POINTER_TYPE)
>> + type = TREE_TYPE (type);
>> const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var));
>> tree inst = create_tmp_var (type, id);
>> - gimplify_assign (inst, fold_indirect_ref (x), &fplist);
>> + if (TREE_CODE (TREE_TYPE (new_var)) == POINTER_TYPE)
>> + gimplify_assign (inst, fold_indirect_ref (x), &fplist);
>> + else
>> + gimplify_assign (inst, fold_indirect_ref (x), &fplist);
>> x = build_fold_addr_expr (inst);
>> }
>> gimplify_assign (new_var, x, &fplist);
>
> (It seems strange to have the same code in both branches of the 'if'
> statement?)
>
>> @@ -7879,7 +7885,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
>> else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
>> {
>> gcc_assert (is_gimple_omp_oacc (ctx->stmt));
>> - if (!omp_is_reference (var))
>> + /* Handle Fortran allocatable scalars. */
>> + if (!omp_is_reference (var)
>> + && TREE_CODE (TREE_TYPE (var)) != POINTER_TYPE)
>> {
>> if (is_gimple_reg (var)
>> && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c))
> | TREE_NO_WARNING (var) = 1;
> | var = build_fold_addr_expr (var);
> | }
> | else
> | talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
> | gimplify_assign (x, var, &ilist);
> | }
>
> That's what's causing regressions, for example for 'firstprivate' clauses
> even in non-offloading situation ('if(0)' clause, for example):
>
> Program received signal SIGSEGV, Segmentation fault.
> 0x0000000000402f8a in main._omp_fn.1 () at source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c:59
> 59 b[ii] = a[ii] + 1;
> (gdb) list 10, 10
> 10 float *a, *b, *d_a, *d_b, exp, exp2;
> (gdb) list 16, 17
> 16 a = (float *) malloc (N * sizeof (float));
> 17 b = (float *) malloc (N * sizeof (float));
> (gdb) list 52, 63
> 52 #pragma acc parallel if(0)
> 53 {
> 54 int ii;
> 55
> 56 for (ii = 0; ii < N; ii++)
> 57 {
> 58 if (acc_on_device (acc_device_host))
> 59 b[ii] = a[ii] + 1;
> 60 else
> 61 b[ii] = a[ii];
> 62 }
> 63 }
>
> So we got here implicit 'firstprivate(a, b)' (which in this scenario
> means no-op, given that the host pointer values are just passed through).
> (On x86_64) these used to have eight bytes alignment, now they have four
> bytes. But worse, the code on the "sending" side is changed as follows
> ('omplower' dump):
>
> - b.57 = b;
> - .omp_data_arr.54.b = &b.57;
> + .omp_data_arr.54.b = b;
> - a.58 = a;
> - .omp_data_arr.54.a = &a.58;
> + .omp_data_arr.54.a = a;
> #pragma omp target oacc_parallel if(0) firstprivate(b) firstprivate(a) [child fn: main._omp_fn.1 (.omp_data_arr.54, .omp_data_sizes.55, .omp_data_kinds.56)]
>
> ..., but the "receiving" side stays the same, so we got a mismatch.
>
> If something like that, then the 'POINTER_TYPE' conditional should
> probably be inside the '!omp_is_reference' conditional, just guarding the
> 'build_fold_addr_expr'?
>
>
> Anyway, I had a look at this now, and seem to have gotten it work.
>
> I will admit, though, that I'm somewhat lost especially with all the
> 'omp_is_reference' usage ("should privatize what this DECL points to
> rather than the DECL itself"). Using that on 'OMP_CLAUSE_DECL ([...])'
> (the common case) makes sense given that's in context of the originating
> source language, but what exactly does it mean when 'omp_is_reference' is
> used on 'new_var = lookup_decl (var, ctx)', or on 'var =
> lookup_decl_in_outer_ctx (ovar, ctx)', where the things looked up by
> these (that is, stored in 'ctx->cb.decl_map') are "arbitrary"/"synthetic"
> items? (Jakub?) Or is it actually improper to use 'omp_is_reference' on
> these, but it just happens to do the expected things in the (several)
> existing cases?
>
> Anyway, for an 'integer, allocatable :: a' that is used 'firstprivate'
> inside an OpenACC offloading region, we now get the following 'omplower'
> changes:
>
> [...]
> integer(kind=4) * a;
> [...]
> a = __builtin_malloc (4);
> [...]
> - a.16 = a;
> - .omp_data_arr.13.a = &a.16;
> + .omp_data_arr.13.a = a;
> #pragma omp target oacc_parallel map(tofrom:b [len: 400]) firstprivate(a) [child fn: MAIN__._omp_fn.0 (.omp_data_arr.13, .omp_data_sizes.14, .omp_data_kinds.15)]
> {
> .omp_data_i = (const struct .omp_data_t.10 & restrict) &.omp_data_arr.13;
> - D.3981 = .omp_data_i->a;
> - a = *D.3981;
> + a = .omp_data_i->a;
> [...]
>
> ..., and that seems to work fine. (..., and no testsuite regressions.)
>
> (I have not yet looked into the related OpenMP changes required.)
>
> Jakub, is the following about right, do you have any comments? If
> approving this patch, please respond with "Reviewed-by: NAME <EMAIL>" so
> that your effort will be recorded in the commit log, see
> <https://gcc.gnu.org/wiki/Reviewed-by>.
>
> The code changes seem very ad-hoc, but that's the common impression that
> I got from looking at/working though a lot of all that OMP code... :-(
>
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -9685,7 +9685,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> {
> gcc_assert (is_gimple_omp_oacc (ctx->stmt));
> if (omp_is_reference (new_var)
> - && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE)
> + && TREE_CODE (TREE_TYPE (new_var)) == POINTER_TYPE)
> + {
> + /* Special handling for Fortran 'allocatable' scalars:
> + avoid indirection. */
> + x = build_receiver_ref (var, false, ctx);
> + }
> + else if (omp_is_reference (new_var))
> {
> /* Create a local object to hold the instance
> value. */
> @@ -9920,7 +9926,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
> {
> gcc_assert (is_gimple_omp_oacc (ctx->stmt));
> - if (!omp_is_reference (var))
> + if (omp_is_reference (lookup_decl (ovar, ctx))
> + && TREE_CODE (TREE_TYPE (ovar)) == POINTER_TYPE)
> + {
> + /* Special handling for Fortran 'allocatable' scalars:
> + avoid indirection. */
> + }
> + else if (!omp_is_reference (var))
> {
> if (is_gimple_reg (var)
> && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c))
Grüße
Thomas
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 658 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20191007/67ebfcd3/attachment.sig>
More information about the Gcc-patches
mailing list