[gomp4] Allow parameter declarations with deviceptr

[gomp4] Allow parameter declarations with deviceptr

[James Norris]

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

Hi,

This patch allows parameter declarations to be used as
arguments to deviceptr for C and C++.

Committed to gomp-4_0-branch.

Jim

[-- Attachment #2: bugfix.patch --]
[-- Type: text/x-patch, Size: 1060 bytes --]

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 88e68ae..dc244ce 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10749,7 +10749,7 @@ c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
 	 c_parser_omp_var_list_parens() should construct a list of
 	 locations to go along with the var list.  */
 
-      if (TREE_CODE (v) != VAR_DECL)
+      if (TREE_CODE (v) != VAR_DECL && TREE_CODE (v) != PARM_DECL)
 	error_at (loc, "%qD is not a variable", v);
       else if (TREE_TYPE (v) == error_mark_node)
 	;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 41fb35e..c233595 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -28122,7 +28122,7 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
 	 c_parser_omp_var_list_parens should construct a list of
 	 locations to go along with the var list.  */
 
-      if (TREE_CODE (v) != VAR_DECL)
+      if (TREE_CODE (v) != VAR_DECL && TREE_CODE (v) != PARM_DECL)
 	error_at (loc, "%qD is not a variable", v);
       else if (TREE_TYPE (v) == error_mark_node)
 	;

Re: [gomp4] Allow parameter declarations with deviceptr

[Cesar Philippidis]

On 07/01/2015 02:25 PM, James Norris wrote:

> This patch allows parameter declarations to be used as
> arguments to deviceptr for C and C++.

Does this fix an existing failure? If not, can you please add a new test
case?

Thanks,
Cesar

[PATCH] Fix PR64748

[James Norris]

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


This fixes the validation of the argument to the deviceptr clause.

Bootstrapped and regtested on x86_64-unknown-linux-gnu.

OK to commit to trunk?

Jim


[-- Attachment #2: deviceptr.patch --]
[-- Type: text/x-patch, Size: 6109 bytes --]

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ceb9e1a..9f0d7af 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10334,11 +10334,11 @@ c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
 	 c_parser_omp_var_list_parens() should construct a list of
 	 locations to go along with the var list.  */
 
-      if (TREE_CODE (v) != VAR_DECL)
-	error_at (loc, "%qD is not a variable", v);
-      else if (TREE_TYPE (v) == error_mark_node)
+      if (TREE_TYPE (v) == error_mark_node)
 	;
-      else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+      else if ((TREE_CODE (v) != VAR_DECL ||
+	       TREE_CODE (v) != PARM_DECL) &&
+	       !POINTER_TYPE_P (TREE_TYPE (v)))
 	error_at (loc, "%qD is not a pointer variable", v);
 
       tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 57dfbcc..37b4712 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -27988,11 +27988,11 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
 	 c_parser_omp_var_list_parens should construct a list of
 	 locations to go along with the var list.  */
 
-      if (TREE_CODE (v) != VAR_DECL)
-	error_at (loc, "%qD is not a variable", v);
-      else if (TREE_TYPE (v) == error_mark_node)
+      if (TREE_TYPE (v) == error_mark_node)
 	;
-      else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+      else if ((TREE_CODE (v) != VAR_DECL ||
+	       TREE_CODE (v) != PARM_DECL) &&
+	       !POINTER_TYPE_P (TREE_TYPE (v)))
 	error_at (loc, "%qD is not a pointer variable", v);
 
       tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 546fa82..5ec7540 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -8,27 +8,29 @@ fun1 (void)
 #pragma acc kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
   ;
 
-#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
+#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a pointer variable" } */
+  /* { dg-error "'fun1' is not a variable in 'map' clause" "fun1 is not a varialbe in map clause" { target *-*-* } 11 } */
   ;
 #pragma acc parallel deviceptr(fun1[2:5])
-  /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 13 } */
-  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 13 } */
+  /* { dg-error "'fun1' is not a pointer variable" "not a pointer variable" { target *-*-* } 14 } */
+  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 14 } */
+  /* { dg-error "'fun1' is not a variable in 'map' clause" "fun1 is not a varialbe in map clause" { target *-*-* } 14 } */
   ;
 
   int i;
 #pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
   ;
 #pragma acc data deviceptr(i[0:4])
-  /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 21 } */
-  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 21 } */
+  /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 23 } */
+  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 23 } */
   ;
 
   float fa[10];
 #pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
   ;
 #pragma acc kernels deviceptr(fa[1:5])
-  /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 29 } */
-  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 29 } */
+  /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 31 } */
+  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 31 } */
   ;
 
   float *fp;
@@ -44,10 +46,11 @@ fun2 (void)
   int i;
   float *fp;
 #pragma acc kernels deviceptr(fp,u,fun2,i,fp)
-  /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */
-  /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */
-  /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */
-  /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */
+  /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 48 } */
+  /* { dg-error "'fun2' is not a pointer variable" "fun2 not a pointer variable" { target *-*-* } 48 } */
+  /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 48 } */
+  /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 48 } */
+  /* { dg-error "'fun2' is not a variable in 'map' clause" "fun2 is not a variable in map clause" { target *-*-* } 48 } */
   ;
 }
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
index e271a37..5443433 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
@@ -1,10 +1,26 @@
 /* { dg-do run } */
 
+#include <string.h>
 #include <stdlib.h>
+#include <openacc.h>
+
+void foo (int *b, const int N)
+{
+  int i;
+
+#pragma acc parallel loop deviceptr (b)
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+  return;
+}
 
 int main (void)
 {
   void *a, *a_1, *a_2;
+  int *b;
+  const int N = 32;
+  int i;
 
 #define A (void *) 0x123
   a = A;
@@ -28,5 +44,35 @@ int main (void)
     abort ();
 #endif
 
+#if !ACC_MEM_SHARED
+  b = (int *) malloc (N * sizeof (int));
+
+  memset (b, 0, (N * sizeof (int)));
+
+#pragma acc data copy (b[0:N])
+  {
+#pragma acc parallel loop deviceptr (b)
+  for (i = 0; i < N; i++)
+    b[i] = i;
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (b[i] != i)
+        abort ();
+    }
+
+#pragma acc data copy (b[0:N])
+  {
+    foo ((int *) acc_deviceptr (b), N);
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (b[i] != i)
+        abort ();
+    }
+#endif
+
   return 0;
 }

Re: [PATCH] Fix PR64748

[James Norris]

Hi!

Ping.

Thanks!


On 02/16/2015 12:26 PM, James Norris wrote:
>
> This fixes the validation of the argument to the deviceptr clause.
>
> Bootstrapped and regtested on x86_64-unknown-linux-gnu.
>
> OK to commit to trunk?
>
> Jim
>

Re: [PATCH] Fix PR64748

[Jeff Law]

On 03/10/15 07:36, James Norris wrote:
> Hi!
>
> Ping.
Note that the GCC trunk is in regression bugfix stage, so this patch may 
(is likely?) be deferred until the next stage1 development cycle.

jeff

Re: [PATCH] Fix PR64748

[Thomas Schwinge]

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

Hi Jeff!

On Tue, 10 Mar 2015 12:48:34 -0600, Jeff Law <law@redhat.com> wrote:
> On 03/10/15 07:36, James Norris wrote:
> > Ping.
> Note that the GCC trunk is in regression bugfix stage, so this patch may 
> (is likely?) be deferred until the next stage1 development cycle.

The decision is with you guys, but my understanding has been, that we can
still apply bug fixes to (new) OpenACC code (Jim's patch is a bug fix),
and possibly even extend the OpenACC code (which has not been part of a
GCC release before, so can't regress).


Grüße,
 Thomas

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

Re: [PATCH] Fix PR64748

[Jakub Jelinek]

On Fri, Mar 13, 2015 at 04:24:09PM +0100, Thomas Schwinge wrote:
> Hi Jeff!
> 
> On Tue, 10 Mar 2015 12:48:34 -0600, Jeff Law <law@redhat.com> wrote:
> > On 03/10/15 07:36, James Norris wrote:
> > > Ping.
> > Note that the GCC trunk is in regression bugfix stage, so this patch may 
> > (is likely?) be deferred until the next stage1 development cycle.
> 
> The decision is with you guys, but my understanding has been, that we can
> still apply bug fixes to (new) OpenACC code (Jim's patch is a bug fix),
> and possibly even extend the OpenACC code (which has not been part of a
> GCC release before, so can't regress).

Yeah, it is acceptable, but as you noted, formatting should be fixed and
also ensure that the same error isn't reported multiple times.
For OpenMP, generally most of the diagnostics that isn't related to parsing
clauses is deferred until *finish_omp_clauses, where after diagnosing
something on a clause the clause is removed and so not reported again later.

	Jakub

Re: [PATCH] Fix PR64748: OpenACC: "is not a variable" error with deviceptr()

[Thomas Schwinge]

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

Hi Jim!

Sorry for the delay.  Please be a little more descriptive in submissions:
for example, if you had used »OpenACC deviceptr clause« instead of
»deviceptr clause«, my mail client's search/filter would have pointed me
to this email earlier, or, by using a more descriptive subject line:
»[PATCH] Fix PR64748: OpenACC: "is not a variable" error with
deviceptr()« (copied from the PR) makes it easier to classify your
submission, compared to just »[PATCH] Fix PR64748«.  Also, it's a good
idea to CC the respective maintainers, that is, Jakub for anything
related to OpenMP (which we're building OpenACC upon).

On Mon, 16 Feb 2015 12:26:32 -0600, James Norris <jnorris@codesourcery.com> wrote:
> This fixes the validation of the argument to the deviceptr clause.
> 
> Bootstrapped and regtested on x86_64-unknown-linux-gnu.
> 
> OK to commit to trunk?

> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -10334,11 +10334,11 @@ c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
>  	 c_parser_omp_var_list_parens() should construct a list of
>  	 locations to go along with the var list.  */
>  
> -      if (TREE_CODE (v) != VAR_DECL)
> -	error_at (loc, "%qD is not a variable", v);
> -      else if (TREE_TYPE (v) == error_mark_node)
> +      if (TREE_TYPE (v) == error_mark_node)
>  	;
> -      else if (!POINTER_TYPE_P (TREE_TYPE (v)))
> +      else if ((TREE_CODE (v) != VAR_DECL ||
> +	       TREE_CODE (v) != PARM_DECL) &&
> +	       !POINTER_TYPE_P (TREE_TYPE (v)))
>  	error_at (loc, "%qD is not a pointer variable", v);

> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> @@ -27988,11 +27988,11 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
>  	 c_parser_omp_var_list_parens should construct a list of
>  	 locations to go along with the var list.  */
>  
> -      if (TREE_CODE (v) != VAR_DECL)
> -	error_at (loc, "%qD is not a variable", v);
> -      else if (TREE_TYPE (v) == error_mark_node)
> +      if (TREE_TYPE (v) == error_mark_node)
>  	;
> -      else if (!POINTER_TYPE_P (TREE_TYPE (v)))
> +      else if ((TREE_CODE (v) != VAR_DECL ||
> +	       TREE_CODE (v) != PARM_DECL) &&
> +	       !POINTER_TYPE_P (TREE_TYPE (v)))
>  	error_at (loc, "%qD is not a pointer variable", v);

I assume there isn't a convenience macro to check for VAR_DECL or
PARM_DECL?  (I don't see any.)

Formatting: »||« and »&&« don't end a line, but begin the next one.  See
<https://www.gnu.org/prep/standards/html_node/Formatting.html>, »When you
split an expression into multiple lines [...]«.

Given these test changes:

> --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
> @@ -8,27 +8,29 @@ fun1 (void)
>  #pragma acc kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
>    ;
>  
> -#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
> +#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a pointer variable" } */
> +  /* { dg-error "'fun1' is not a variable in 'map' clause" "fun1 is not a varialbe in map clause" { target *-*-* } 11 } */

..., and so on, I think we should avoid reporting several errors for one
wrong deviceptr usage.  Did the code's original structure have this
effect, to "let through" the erroneous v, without raising an additional
error message?

>    ;
>  #pragma acc parallel deviceptr(fun1[2:5])
> -  /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 13 } */
> -  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 13 } */
> +  /* { dg-error "'fun1' is not a pointer variable" "not a pointer variable" { target *-*-* } 14 } */
> +  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 14 } */
> +  /* { dg-error "'fun1' is not a variable in 'map' clause" "fun1 is not a varialbe in map clause" { target *-*-* } 14 } */
>    ;
>  
>    int i;
>  #pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
>    ;
>  #pragma acc data deviceptr(i[0:4])
> -  /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 21 } */
> -  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 21 } */
> +  /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 23 } */
> +  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 23 } */
>    ;
>  
>    float fa[10];
>  #pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
>    ;
>  #pragma acc kernels deviceptr(fa[1:5])
> -  /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 29 } */
> -  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 29 } */
> +  /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 31 } */
> +  /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 31 } */
>    ;
>  
>    float *fp;
> @@ -44,10 +46,11 @@ fun2 (void)
>    int i;
>    float *fp;
>  #pragma acc kernels deviceptr(fp,u,fun2,i,fp)
> -  /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */
> -  /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */
> -  /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */
> -  /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */
> +  /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 48 } */
> +  /* { dg-error "'fun2' is not a pointer variable" "fun2 not a pointer variable" { target *-*-* } 48 } */
> +  /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 48 } */
> +  /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 48 } */
> +  /* { dg-error "'fun2' is not a variable in 'map' clause" "fun2 is not a variable in map clause" { target *-*-* } 48 } */
>    ;
>  }
>  
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
> index e271a37..5443433 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c
> @@ -1,10 +1,26 @@
>  /* { dg-do run } */
>  
> +#include <string.h>
>  #include <stdlib.h>
> +#include <openacc.h>
> +
> +void foo (int *b, const int N)
> +{
> +  int i;
> +
> +#pragma acc parallel loop deviceptr (b)
> +  for (i = 0; i < N; i++)
> +    b[i] = i;
> +
> +  return;
> +}
>  
>  int main (void)
>  {
>    void *a, *a_1, *a_2;
> +  int *b;
> +  const int N = 32;
> +  int i;
>  
>  #define A (void *) 0x123
>    a = A;
> @@ -28,5 +44,35 @@ int main (void)
>      abort ();
>  #endif
>  
> +#if !ACC_MEM_SHARED
> +  b = (int *) malloc (N * sizeof (int));
> +
> +  memset (b, 0, (N * sizeof (int)));
> +
> +#pragma acc data copy (b[0:N])
> +  {
> +#pragma acc parallel loop deviceptr (b)
> +  for (i = 0; i < N; i++)
> +    b[i] = i;
> +  }
> +
> +  for (i = 0; i < N; i++)
> +    {
> +      if (b[i] != i)
> +        abort ();
> +    }
> +
> +#pragma acc data copy (b[0:N])
> +  {
> +    foo ((int *) acc_deviceptr (b), N);
> +  }
> +
> +  for (i = 0; i < N; i++)
> +    {
> +      if (b[i] != i)
> +        abort ();
> +    }
> +#endif
> +
>    return 0;
>  }


Grüße,
 Thomas

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

Re: [gomp4] Allow parameter declarations with deviceptr

[Thomas Schwinge]

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

Hi!

On Wed, 1 Jul 2015 16:33:24 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 07/01/2015 02:25 PM, James Norris wrote:
> 
> > This patch allows parameter declarations to be used as
> > arguments to deviceptr for C and C++.

Thanks!  I suppose this does fix <http://gcc.gnu.org/PR64748>?

> Does this fix an existing failure? If not, can you please add a new test
> case?

An earlier submission,
<http://news.gmane.org/find-root.php?message_id=%3C54E23658.6060105%40codesourcery.com%3E>,
did include some testsuite changes -- but I had not seen any update of
this patch after Jakub's and my review comments.


Grüße,
 Thomas

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