[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