Make static vars inside of target regions or declare target routines implicitly declare target to (PR middle-end/90779)

Hi!

The OpenMP specification isn't clear on this, I'll work on getting that
clarified for 5.1, but the agreement on omp-lang has been that it should
work the way the patch implements it, static block scope variables inside of
#pragma omp target or #pragma omp declare target routines are handled as if
they have #pragma omp declare target to (variable).

Bootstrapped/regtested on x86_64-linux and i686-linux, unfortunately it
regresses:
+FAIL: c-c++-common/goacc/routine-5.c  (test for errors, line 204)

Thus, I'm not committing it right now and want to ask what should be done
for OpenACC.  The patch uses & ORT_TARGET tests, so it affects both OpenMP
target region, and OpenACC parallel/kernels and both OpenMP and OpenACC
target routines.  Is it ok to do it that way and just adjust the routine-5.c
test, or shall it test (ctx->region_type & (ORT_TARGET | ORT_ACC)) ==
ORT_TARGET, i.e. only OpenMP and not OpenACC?  If so, there is still the
problem that gimplify_body.c does:
  if (flag_openacc || flag_openmp)
    {
      gcc_assert (gimplify_omp_ctxp == NULL);
      if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
        gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
    }
We'd need different attribute (or additional attribute) for OpenACC routines
and would need to use new_omp_context (cond ? ORT_TARGET : ORT_ACC_PARALLEL)
or similar to express OpenACC routines.

2019-06-12  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/90779
	* gimplify.c (gimplify_bind_expr): Add "omp declare target" attributes
	to static block scope variables inside of target region or target
	functions.

	* testsuite/libgomp.c/pr90779.c: New test.
	* testsuite/libgomp.fortran/pr90779.f90: New test.


	Jakub

On 13-06-19 10:34, Jakub Jelinek wrote:
> Hi!
> 
> The OpenMP specification isn't clear on this, I'll work on getting that
> clarified for 5.1, but the agreement on omp-lang has been that it should
> work the way the patch implements it, static block scope variables inside of
> #pragma omp target or #pragma omp declare target routines are handled as if
> they have #pragma omp declare target to (variable).
> 
> Bootstrapped/regtested on x86_64-linux and i686-linux, unfortunately it
> regresses:
> +FAIL: c-c++-common/goacc/routine-5.c  (test for errors, line 204)
> 
> Thus, I'm not committing it right now and want to ask what should be done
> for OpenACC.

OpenACC 2.6 - 2.15.1. Routine Directive - Restrictions:
...
In C and C++, function static variables are not supported in functions
to which a routine directive applies.
...
[ And text is still the same for 2.7. ]

Thanks,
- Tom

> The patch uses & ORT_TARGET tests, so it affects both OpenMP
> target region, and OpenACC parallel/kernels and both OpenMP and OpenACC
> target routines.  Is it ok to do it that way and just adjust the routine-5.c
> test, or shall it test (ctx->region_type & (ORT_TARGET | ORT_ACC)) ==
> ORT_TARGET, i.e. only OpenMP and not OpenACC?  If so, there is still the
> problem that gimplify_body.c does:
>   if (flag_openacc || flag_openmp)
>     {
>       gcc_assert (gimplify_omp_ctxp == NULL);
>       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
>         gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
>     }
> We'd need different attribute (or additional attribute) for OpenACC routines
> and would need to use new_omp_context (cond ? ORT_TARGET : ORT_ACC_PARALLEL)
> or similar to express OpenACC routines.
> 
> 2019-06-12  Jakub Jelinek  <jakub@redhat.com>
> 
> 	PR middle-end/90779
> 	* gimplify.c (gimplify_bind_expr): Add "omp declare target" attributes
> 	to static block scope variables inside of target region or target
> 	functions.
> 
> 	* testsuite/libgomp.c/pr90779.c: New test.
> 	* testsuite/libgomp.fortran/pr90779.f90: New test.
> 
> --- gcc/gimplify.c.jj	2019-06-10 19:42:03.868959986 +0200
> +++ gcc/gimplify.c	2019-06-12 13:00:18.765167777 +0200
> @@ -1323,17 +1323,37 @@ gimplify_bind_expr (tree *expr_p, gimple
>  	  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
>  
>  	  /* Mark variable as local.  */
> -	  if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t)
> -	      && (! DECL_SEEN_IN_BIND_EXPR_P (t)
> -		  || splay_tree_lookup (ctx->variables,
> -					(splay_tree_key) t) == NULL))
> +	  if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t))
>  	    {
> -	      if (ctx->region_type == ORT_SIMD
> -		  && TREE_ADDRESSABLE (t)
> -		  && !TREE_STATIC (t))
> -		omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
> -	      else
> -		omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
> +	      if (! DECL_SEEN_IN_BIND_EXPR_P (t)
> +		  || splay_tree_lookup (ctx->variables,
> +					(splay_tree_key) t) == NULL)
> +		{
> +		  if (ctx->region_type == ORT_SIMD
> +		      && TREE_ADDRESSABLE (t)
> +		      && !TREE_STATIC (t))
> +		    omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
> +		  else
> +		    omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
> +		}
> +	      /* Static locals inside of target construct or offloaded
> +		 routines need to be "omp declare target".  */
> +	      if (TREE_STATIC (t))
> +		for (; ctx; ctx = ctx->outer_context)
> +		  if ((ctx->region_type & ORT_TARGET) != 0)
> +		    {
> +		      if (!lookup_attribute ("omp declare target",
> +					     DECL_ATTRIBUTES (t)))
> +			{
> +			  tree id = get_identifier ("omp declare target");
> +			  DECL_ATTRIBUTES (t)
> +			    = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
> +			  varpool_node *node = varpool_node::get (t);
> +			  if (node)
> +			    node->offloadable = 1;
> +			}
> +		      break;
> +		    }
>  	    }
>  
>  	  DECL_SEEN_IN_BIND_EXPR_P (t) = 1;
> --- libgomp/testsuite/libgomp.c/pr90779.c.jj	2019-06-12 13:01:57.081667587 +0200
> +++ libgomp/testsuite/libgomp.c/pr90779.c	2019-06-12 12:41:15.637730797 +0200
> @@ -0,0 +1,18 @@
> +/* PR middle-end/90779 */
> +
> +extern void abort (void);
> +
> +int
> +main ()
> +{
> +  int i, j;
> +  for (i = 0; i < 2; ++i)
> +    #pragma omp target map(from: j)
> +    {
> +      static int k = 5;
> +      j = ++k;
> +    }
> +  if (j != 7)
> +    abort ();
> +  return 0;
> +}
> --- libgomp/testsuite/libgomp.fortran/pr90779.f90.jj	2019-06-12 12:43:17.891825811 +0200
> +++ libgomp/testsuite/libgomp.fortran/pr90779.f90	2019-06-12 12:43:08.421973375 +0200
> @@ -0,0 +1,12 @@
> +! PR middle-end/90779
> +
> +program pr90779
> +  implicit none
> +  integer :: v(4), i
> +
> +  !$omp target map(from:v)
> +    v(:) = (/ (i, i=1,4) /)
> +  !$omp end target
> +
> +  if (any (v .ne. (/ (i, i=1,4) /))) stop 1
> +end program
> 
> 	Jakub
>

Hi!

So I recently had reason to verify how 'static' variables behave in
OpenACC compute construct regions as well as OpenACC 'routine'.  Two
weeks ago I started writing a few testcases -- and today then wondered if
maybe there's something in the GCC archives about this.  And there is:
directly related are <https://gcc.gnu.org/PR84991> "[openacc] Misleading
error message for function static var in routine",
<https://gcc.gnu.org/PR84992> "[openacc] function static var in
parallel", and more generally <https://gcc.gnu.org/PR90779> "Fortran
array initialization in offload regions" that is discussed here.  (I had
taken part in at least some of these discussions, yet didn't directly
remember these now...  8-| Too much going on?)

Anyway:

On 2019-06-15T10:19:23+0200, Tom de Vries <tdevries@suse.de> wrote:
> On 13-06-19 10:34, Jakub Jelinek wrote:
>> The OpenMP specification isn't clear on this, I'll work on getting that
>> clarified for 5.1, but the agreement on omp-lang has been that it should
>> work the way the patch implements it, static block scope variables inside of
>> #pragma omp target or #pragma omp declare target routines are handled as if
>> they have #pragma omp declare target to (variable).
>>
>> Bootstrapped/regtested on x86_64-linux and i686-linux, unfortunately it
>> regresses:
>> +FAIL: c-c++-common/goacc/routine-5.c  (test for errors, line 204)
>>
>> Thus, I'm not committing it right now and want to ask what should be done
>> for OpenACC.
>
> OpenACC 2.6 - 2.15.1. Routine Directive - Restrictions:
> ...
> In C and C++, function static variables are not supported in functions
> to which a routine directive applies.
> ...
> [ And text is still the same for 2.7. ]

..., and still in OpenACC 3.1.

But yes, that seems somewhat incomplete and/or inconsistent.  I've filed
<https://github.com/OpenACC/openacc-spec/issues/372> "C/C++ 'static'
variables" (only visible to members of the GitHub OpenACC organization).

I do agree that even if OpenACC ultimately doesn't want to support
certain cases of (?) 'static' variables, we still have to (and evidently
can) support 'static' for compiler-synthesized variables, per our own
desired semantics, which happen to match OpenMP's (as I understand this).

>> The patch uses & ORT_TARGET tests, so it affects both OpenMP
>> target region, and OpenACC parallel/kernels and both OpenMP and OpenACC
>> target routines.  Is it ok to do it that way and just adjust the routine-5.c
>> test, or shall it test (ctx->region_type & (ORT_TARGET | ORT_ACC)) ==
>> ORT_TARGET, i.e. only OpenMP and not OpenACC?  If so, there is still the
>> problem that gimplify_body.c does:
>>   if (flag_openacc || flag_openmp)
>>     {
>>       gcc_assert (gimplify_omp_ctxp == NULL);
>>       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
>>         gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
>>     }
>> We'd need different attribute (or additional attribute) for OpenACC routines
>> and would need to use new_omp_context (cond ? ORT_TARGET : ORT_ACC_PARALLEL)
>> or similar to express OpenACC routines.

I'm fine to have this supported for GCC/OpenACC in the way that it
currently is, so no need to special-case that.  If OpenACC decides
otherwise, we'll then adjust.


(I have not reviewed the PR90779 code changes; it was sufficient for my
case to understand what I called GCC's observed behavior.)


I've now pushed "Add 'libgomp.oacc-c-c++-common/static-variable-1.c'
[PR84991, PR84992, PR90779]" to master branch in commit
ffa0ae6eeef3ad15d3f288283e4c477193052f1a, and releases/gcc-10 branch in
commit 60b589b5858fb8ad414583c6b493e0897f1bde5f, see attached.  (The
PR90779 code changes never got backported to GCC 9 and 8 release
branches.)

Also I've filed <https://gcc.gnu.org/PR100001> "[GCN offloading]
Occasional C++ 'libgomp.oacc-c-c++-common/static-variable-1.c' execution
failure".


Grüße
 Thomas


>> 2019-06-12  Jakub Jelinek  <jakub@redhat.com>
>>
>>      PR middle-end/90779
>>      * gimplify.c (gimplify_bind_expr): Add "omp declare target" attributes
>>      to static block scope variables inside of target region or target
>>      functions.
>>
>>      * testsuite/libgomp.c/pr90779.c: New test.
>>      * testsuite/libgomp.fortran/pr90779.f90: New test.
>>
>> --- gcc/gimplify.c.jj        2019-06-10 19:42:03.868959986 +0200
>> +++ gcc/gimplify.c   2019-06-12 13:00:18.765167777 +0200
>> @@ -1323,17 +1323,37 @@ gimplify_bind_expr (tree *expr_p, gimple
>>        struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
>>
>>        /* Mark variable as local.  */
>> -      if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t)
>> -          && (! DECL_SEEN_IN_BIND_EXPR_P (t)
>> -              || splay_tree_lookup (ctx->variables,
>> -                                    (splay_tree_key) t) == NULL))
>> +      if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t))
>>          {
>> -          if (ctx->region_type == ORT_SIMD
>> -              && TREE_ADDRESSABLE (t)
>> -              && !TREE_STATIC (t))
>> -            omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
>> -          else
>> -            omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
>> +          if (! DECL_SEEN_IN_BIND_EXPR_P (t)
>> +              || splay_tree_lookup (ctx->variables,
>> +                                    (splay_tree_key) t) == NULL)
>> +            {
>> +              if (ctx->region_type == ORT_SIMD
>> +                  && TREE_ADDRESSABLE (t)
>> +                  && !TREE_STATIC (t))
>> +                omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
>> +              else
>> +                omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
>> +            }
>> +          /* Static locals inside of target construct or offloaded
>> +             routines need to be "omp declare target".  */
>> +          if (TREE_STATIC (t))
>> +            for (; ctx; ctx = ctx->outer_context)
>> +              if ((ctx->region_type & ORT_TARGET) != 0)
>> +                {
>> +                  if (!lookup_attribute ("omp declare target",
>> +                                         DECL_ATTRIBUTES (t)))
>> +                    {
>> +                      tree id = get_identifier ("omp declare target");
>> +                      DECL_ATTRIBUTES (t)
>> +                        = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
>> +                      varpool_node *node = varpool_node::get (t);
>> +                      if (node)
>> +                        node->offloadable = 1;
>> +                    }
>> +                  break;
>> +                }
>>          }
>>
>>        DECL_SEEN_IN_BIND_EXPR_P (t) = 1;
>> --- libgomp/testsuite/libgomp.c/pr90779.c.jj 2019-06-12 13:01:57.081667587 +0200
>> +++ libgomp/testsuite/libgomp.c/pr90779.c    2019-06-12 12:41:15.637730797 +0200
>> @@ -0,0 +1,18 @@
>> +/* PR middle-end/90779 */
>> +
>> +extern void abort (void);
>> +
>> +int
>> +main ()
>> +{
>> +  int i, j;
>> +  for (i = 0; i < 2; ++i)
>> +    #pragma omp target map(from: j)
>> +    {
>> +      static int k = 5;
>> +      j = ++k;
>> +    }
>> +  if (j != 7)
>> +    abort ();
>> +  return 0;
>> +}
>> --- libgomp/testsuite/libgomp.fortran/pr90779.f90.jj 2019-06-12 12:43:17.891825811 +0200
>> +++ libgomp/testsuite/libgomp.fortran/pr90779.f90    2019-06-12 12:43:08.421973375 +0200
>> @@ -0,0 +1,12 @@
>> +! PR middle-end/90779
>> +
>> +program pr90779
>> +  implicit none
>> +  integer :: v(4), i
>> +
>> +  !$omp target map(from:v)
>> +    v(:) = (/ (i, i=1,4) /)
>> +  !$omp end target
>> +
>> +  if (any (v .ne. (/ (i, i=1,4) /))) stop 1
>> +end program
>>
>>      Jakub
>>


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf