[PR90742] OpenACC/OpenMP target offloading: Fortran 'allocatable' scalars in 'firstprivate' clauses

[Thomas Schwinge <thomas@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 9322 bytes --]

Hi!

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

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]