Re: [PATCH] Fix use of declare'd vars by routine procedures.

[James Norris <jnorris at codesourcery dot com>]

Hi!

On 01/29/2016 02:31 AM, Jakub Jelinek wrote:
> On Thu, Jan 28, 2016 at 12:26:38PM -0600, James Norris wrote:
> > I think the attached change is what you had in mind with
> > regard to doing the check at gimplification time.
>
> Nope, this is still a wrong location for that.
> If you look at the next line after the block you've added, you'll see
> if (gimplify_omp_ctxp && omp_notice_variable (gimplify_omp_ctxp, decl, true))
> And that function fairly early calls is_global_var (decl).  So you know
> already gimplify_omp_ctxp && is_global_var (decl), just put the rest into
> that block.

When said you said "rest into that block", I assumed you meant
the block in omp_notice_variable () that has the function call
to is_global_var () as the if condition.

> The TREE_CODE (decl) == VAR_DECL check is VAR_P (decl).

Fixed.

> What do you want to achieve with
>
> > +      && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
> > +       || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))
>
> ?  is_global_var already guarantees you that it is either TREE_STATIC
> or DECL_EXTERNAL, why is that not good enough?

As you pointed out above, the check is not necessary.

> > [snip snip]
> >
> > +  /* Validate variable for use within routine function.  */
> > +  if (gimplify_omp_ctxp && gimplify_omp_ctxp->region_type == ORT_TARGET
> > +      && get_oacc_fn_attrib (current_function_decl)
>
> If you only care about the implicit target region of acc routine, I think
> you also want to check that gimplify_omp_ctxp->outer_context == NULL.

Check added.

> > [snip snip]
>
> And I'm really confused by this error message.  If you are complaining that
> the variable is not listed in acc declare clauses, why don't you say that?
> What does the error have to do with its storage class?
> Also, splitting one error into two is weird, usually there would be one
> error message and perhaps inform after it.

Error messages re-written to make better sense.

Thanks!

Jim

===== ChangeLog entries....

        gcc/
        * gimplify.c (omp_notice_variable): Add variable usage check.
        gcc/testsuite/
        * c-c++-common/goacc/routine-5.c: Add tests.



diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9f99842..41cd7a7 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,7 @@
+2016-01-XX  James Norris  <jnorris@codesourcery.com>
+
+	* gimplify.c (omp_notice_variable): Add usage check.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/69551
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 32bc1fd..5833605 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6087,9 +6087,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   if (ctx->region_type == ORT_NONE)
     return lang_hooks.decls.omp_disregard_value_expr (decl, false);
 
-  /* Threadprivate variables are predetermined.  */
   if (is_global_var (decl))
     {
+      /* Threadprivate variables are predetermined.  */
       if (DECL_THREAD_LOCAL_P (decl))
 	return omp_notice_threadprivate_variable (ctx, decl, NULL_TREE);
 
@@ -6100,6 +6100,30 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	  if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
 	    return omp_notice_threadprivate_variable (ctx, decl, value);
 	}
+
+      if (gimplify_omp_ctxp->outer_context == NULL
+	  && VAR_P (decl)
+	  && get_oacc_fn_attrib (current_function_decl))
+	{
+	  location_t loc = DECL_SOURCE_LOCATION (decl);
+
+	  if (lookup_attribute ("omp declare target link",
+				DECL_ATTRIBUTES (decl)))
+	    {
+	      error_at (loc,
+			"%qE with %<link%> clause used in %<routine%> function",
+			DECL_NAME (decl));
+	      return false;
+	    }
+	  else if (!lookup_attribute ("omp declare target",
+				      DECL_ATTRIBUTES (decl)))
+	    {
+	      error_at (loc,
+			"%qE requires a %<declare%> directive for use "
+			"in a %<routine%> function", DECL_NAME (decl));
+	      return false;
+	    }
+	}
     }
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
@@ -8223,7 +8247,7 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
 	      if (oacc_declare_returns == NULL)
 		oacc_declare_returns = new hash_map<tree, tree>;
 
-	      oacc_declare_returns->put (decl, c);
+		oacc_declare_returns->put (decl, c);
 	    }
 	}
 
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 150ebc8..61003c2 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,7 @@
+2016-01-XX  James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/routine-5.c: Add tests.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/69551
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index ccda097..c34838f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -45,3 +45,97 @@ using namespace g;
 #pragma acc routine (a) /* { dg-error "does not refer to" } */
   
 #pragma acc routine (c) /* { dg-error "does not refer to" } */
+
+int vb1;		/* { dg-error "directive for use" } */
+extern int vb2;		/* { dg-error "directive for use" } */
+static int vb3;		/* { dg-error "directive for use" } */
+
+#pragma acc routine
+int
+func1 (int a)
+{
+  vb1 = a + 1;
+  vb2 = vb1 + 1;
+  vb3 = vb2 + 1;
+
+  return vb3;
+}
+
+#pragma acc routine
+int
+func2 (int a)
+{
+  extern int vb4;	/* { dg-error "directive for use" } */
+  static int vb5;	/* { dg-error "directive for use" } */
+
+  vb4 = a + 1;
+  vb5 = vb4 + 1;
+
+  return vb5;
+}
+
+extern int vb6;			/* { dg-error "clause used in" } */
+#pragma acc declare link (vb6)
+static int vb7;			/* { dg-error "clause used in" } */
+#pragma acc declare link (vb7)
+
+#pragma acc routine
+int
+func3 (int a)
+{
+  vb6 = a + 1;
+  vb7 = vb6 + 1;
+
+  return vb7;
+}
+
+int vb8;
+#pragma acc declare create (vb8)
+extern int vb9;
+#pragma acc declare create (vb9)
+static int vb10;
+#pragma acc declare create (vb10)
+
+#pragma acc routine
+int
+func4 (int a)
+{
+  vb8 = a + 1;
+  vb9 = vb8 + 1;
+  vb10 = vb9 + 1;
+
+  return vb10;
+}
+
+int vb11;
+#pragma acc declare device_resident (vb11)
+extern int vb12;
+#pragma acc declare device_resident (vb12)
+extern int vb13;
+#pragma acc declare device_resident (vb13)
+
+#pragma acc routine
+int
+func5 (int a)
+{
+  vb11 = a + 1;
+  vb12 = vb11 + 1;
+  vb13 = vb12 + 1;
+
+  return vb13;
+}
+
+#pragma acc routine
+int
+func6 (int a)
+{
+  extern int vb14;
+#pragma acc declare create (vb14)
+  static int vb15;
+#pragma acc declare create (vb15)
+
+  vb14 = a + 1;
+  vb15 = vb14 + 1;
+
+  return vb15;
+}