[PATCH][OpenACC] Add support for firstprivate Fortran allocatable scalars

[PATCH][OpenACC] Add support for firstprivate Fortran allocatable scalars

[Cesar Philippidis]

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

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

Is this OK for trunk? I bootstrapped and regtested it for x86_64 with
nvptx offloading.

Thanks,
Cesar

[-- Attachment #2: 0001-OpenACC-Add-support-for-firstprivate-Fortran-allocat.patch --]
[-- Type: text/x-patch, Size: 4271 bytes --]

From b8fb83b36d0f96b12af9a1f5596f31b3c6b72ef0 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Mon, 6 Aug 2018 09:19:28 -0700
Subject: [PATCH] [OpenACC] Add support for firstprivate Fortran allocatable
 scalars

This patch updates the way that lower_omp_target uses firstprivate
pointers in OpenACC offloaded regions. On host side, when preparing
pointer type firstprivate data mapping, not to be confused with
GOMP_MAP_FIRSTPRIVATE_POINTER, the compiler passes passes the address
of the value being pointed to, not the address of the pointer
itself. Correspondingly, on the device side, the compiler generates to
deference the remapped pointer once and copy the data to a local
buffer.

While this behavior like it would break things, it will not affect C
or C++ data mappings, because those languages transfer pointers via
GOMP_MAP_FIRSTPRIVATE_POINTER. 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.

2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	omp-low.c (lower_omp_target): Update OpenACC handling of
	pointer variables with GOMP_MAP_FIRSTPRIVATE mappings.

	libgomp/
	testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New
	test.
---
 gcc/omp-low.c                                      | 18 ++++++++----
 .../libgomp.oacc-fortran/allocatable-scalar.f90    | 33 ++++++++++++++++++++++
 2 files changed, 46 insertions(+), 5 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/allocatable-scalar.f90

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 843c66f..47603c4 100644
--- 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);
@@ -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))
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/allocatable-scalar.f90 b/libgomp/testsuite/libgomp.oacc-fortran/allocatable-scalar.f90
new file mode 100644
index 0000000..be86d14
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/allocatable-scalar.f90
@@ -0,0 +1,33 @@
+! Test non-declared allocatable scalars in OpenACC data clauses.
+
+! { dg-do run }
+
+program main
+  implicit none
+  integer, parameter :: n = 100
+  integer, allocatable :: a, c
+  integer :: i, b(n)
+
+  allocate (a)
+
+  a = 50
+
+  !$acc parallel loop
+  do i = 1, n;
+     b(i) = a
+  end do
+
+  do i = 1, n
+     if (b(i) /= a) call abort
+  end do
+
+  allocate (c)
+
+  !$acc parallel copyout(c) num_gangs(1)
+  c = a
+  !$acc end parallel
+
+  if (c /= a) call abort
+
+  deallocate (a, c)
+end program main
-- 
2.7.4

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

[Thomas Schwinge]

[-- 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 --]

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

[Thomas Schwinge]

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

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

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