diff mbox

Fix use of declare'd vars by routine procedures.

Message ID 56AA5D5E.2030903@codesourcery.com
State New
Headers show

Commit Message

James Norris Jan. 28, 2016, 6:26 p.m. UTC 
Jakub,

On 01/22/2016 04:39 AM, Jakub Jelinek wrote:
>
> As for the extra error, I think it would be better to diagnose this during
> gimplification, that way you do it for all FEs, plus you have the omp
> context in there, etc.  The error wording is weird, the diagnostic
> should make it clear use of what is invalid in what and why.
>
> 	Jakub
>

I think the attached change is what you had in mind with
regard to doing the check at gimplification time.

OK?

Thanks for taking the time to review,
Jim


ChangeLog entries....

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

Comments

Jakub Jelinek Jan. 29, 2016, 8:31 a.m. UTC | #1 
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.
The TREE_CODE (decl) == VAR_DECL check is VAR_P (decl).
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?

> diff --git a/trunk/gcc/gimplify.c b/trunk/gcc/gimplify.c
> --- a/trunk/gcc/gimplify.c	(revision 232802)
> +++ b/trunk/gcc/gimplify.c	(working copy)
> @@ -1841,6 +1841,33 @@
>        return GS_ERROR;
>      }
>  
> +  /* 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.

> +      && TREE_CODE (decl) == VAR_DECL
> +      && is_global_var (decl)
> +      && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
> +	  || (!TREE_STATIC (decl) && DECL_EXTERNAL (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 GS_ERROR;
> +	}
> +      else if (!lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
> +	{
> +	  error_at (loc,
> +		    "storage class of %qE cannot be ", DECL_NAME (decl));
> +	  error_at (gimplify_omp_ctxp->location,
> +		    "used in enclosing %<routine%> function");

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.

	Jakub
diff mbox

Patch

Index: gcc/ChangeLog
===================================================================
diff --git a/trunk/gcc/ChangeLog b/trunk/gcc/ChangeLog
--- a/trunk/gcc/ChangeLog	(revision 232802)
+++ b/trunk/gcc/ChangeLog	(working copy)
@@ -1,3 +1,7 @@ 
+2016-01-XX  James Norris  <jnorris@codesourcery.com>
+
+	* gimplify.c (gimplify_var_or_parm_decl): Add usage check.
+
 2016-01-25  Jeff Law  <law@redhat.com>
 
 	PR tree-optimization/69196
Index: gcc/gimplify.c
===================================================================
diff --git a/trunk/gcc/gimplify.c b/trunk/gcc/gimplify.c
--- a/trunk/gcc/gimplify.c	(revision 232802)
+++ b/trunk/gcc/gimplify.c	(working copy)
@@ -1841,6 +1841,33 @@ 
       return GS_ERROR;
     }
 
+  /* 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)
+      && TREE_CODE (decl) == VAR_DECL
+      && is_global_var (decl)
+      && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
+	  || (!TREE_STATIC (decl) && DECL_EXTERNAL (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 GS_ERROR;
+	}
+      else if (!lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
+	{
+	  error_at (loc,
+		    "storage class of %qE cannot be ", DECL_NAME (decl));
+	  error_at (gimplify_omp_ctxp->location,
+		    "used in enclosing %<routine%> function");
+	  return GS_ERROR;
+	}
+    }
+
   /* When within an OMP context, notice uses of variables.  */
   if (gimplify_omp_ctxp && omp_notice_variable (gimplify_omp_ctxp, decl, true))
     return GS_ALL_DONE;
Index: gcc/testsuite/ChangeLog
===================================================================
diff --git a/trunk/gcc/testsuite/ChangeLog b/trunk/gcc/testsuite/ChangeLog
--- a/trunk/gcc/testsuite/ChangeLog	(revision 232802)
+++ b/trunk/gcc/testsuite/ChangeLog	(working copy)
@@ -1,3 +1,7 @@ 
+2016-01-XX  James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/routine-5.c: Add tests.
+
 2016-01-25  Jeff Law  <law@redhat.com>
 
 	PR tree-optimization/69196
Index: gcc/testsuite/c-c++-common/goacc/routine-5.c
===================================================================
diff --git a/trunk/gcc/testsuite/c-c++-common/goacc/routine-5.c b/trunk/gcc/testsuite/c-c++-common/goacc/routine-5.c
--- a/trunk/gcc/testsuite/c-c++-common/goacc/routine-5.c	(revision 232802)
+++ b/trunk/gcc/testsuite/c-c++-common/goacc/routine-5.c	(working copy)
@@ -45,3 +45,70 @@ 
 #pragma acc routine (a) /* { dg-error "does not refer to" } */
   
 #pragma acc routine (c) /* { dg-error "does not refer to" } */
+
+float vb1; /* { dg-error "storage class of" } */
+
+#pragma acc routine
+int
+func1 (int a)		/* { dg-error "used in enclosing" } */
+{
+  vb1 = a + 1;
+
+  return vb1;
+}
+
+#pragma acc routine
+int
+func2 (int a)		/* { dg-error "used in enclosing" } */
+{
+  static int vb2;	/* { dg-error "storage class of" } */
+
+  vb2 = a + 1;
+
+  return vb2;
+}
+
+float vb3;		/* { dg-error "'link' clause used" } */
+#pragma acc declare link (vb3)
+
+#pragma acc routine
+int
+func3 (int a)
+{
+  vb3 = a + 1;
+
+  return vb3;
+}
+
+float vb4;
+#pragma acc declare create (vb4)
+
+#pragma acc routine
+int
+func4 (int a)
+{
+  vb4 = a + 1;
+
+  return vb4;
+}
+
+extern float vb5;	/* { dg-error "storage class of" } */
+
+#pragma acc routine
+int
+func5 (int a)		/* { dg-error "used in enclosing" } */
+{
+  vb5 = a + 1;
+
+  return vb5;
+}
+
+#pragma acc routine
+int
+func6 (int a)		/* { dg-error "used in enclosing" } */
+{
+  extern float vb6;	/* { dg-error "storage class of" } */
+  vb6 = a + 1;
+
+  return vb6;
+}