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

Message ID 20190613083430.GA19695@tucnak
State New
Headers show
Series
  • Make static vars inside of target regions or declare target routines implicitly declare target to (PR middle-end/90779)
Related show

Commit Message

Jakub Jelinek June 13, 2019, 8:34 a.m.
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

Comments

Tom de Vries June 15, 2019, 8:19 a.m. | #1
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

>

Patch

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