[PATCH] Fix PR64748

[PATCH] Fix PR64748

[James Norris]

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

Hi,

The attached patch resolves c/PR64748. The patch
adds the use of parm's with the deviceptr clause.

Question....

As there is VAR_P (), could there be a PARM_P ()?
Or would that obscure something I'm not aware of?

Regtested and bootstrapped on x86_64.

Thanks,
Jim

==== ChangeLog entries...

         gcc/c/
         PR c/64748
         * c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.

         gcc/cp/
         PR c/64748
         * parser.c (cp_parser_oacc_data_clause_deviceptr): Allow parms.

         gcc/testsuite/
         PR c/64748
         * c-c++-common/goacc/deviceptr-1.c: Add tests.


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

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 5341f04..f2d114c 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.
+
 2016-01-27  Jakub Jelinek  <jakub@redhat.com>
 
 	PR debug/66869
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index eede3a7..f61f559 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10760,7 +10760,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 (!VAR_P (v))
+      if (!VAR_P (v) && !(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/ChangeLog b/gcc/cp/ChangeLog
index 3b5c9d5..b11b859 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* parser.c (cp_parser_oacc_data_clause_deviceptr): Allow parms.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR debug/66869
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d03b0c9..de96b44 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -30087,7 +30087,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 (!VAR_P (v))
+      if (!VAR_P (v) && !(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/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 150ebc8..db281cd 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* c-c++-common/goacc/deviceptr-1.c: Add tests.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/69551
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 546fa82..6edbdb1 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -84,3 +84,21 @@ fun4 (void)
 #pragma acc parallel deviceptr(s2_p)
   s2_p = 0;
 }
+
+void
+func5 (float *fp)
+{
+
+#pragma acc data deviceptr (fp)
+{ }
+
+}
+
+void
+func6 (float fp)
+{
+
+#pragma acc data deviceptr (fp)	/* { dg-error "is not a pointer variable" } */
+{ }
+
+}

Re: [PATCH] Fix PR64748

[Jakub Jelinek]

On Mon, Feb 01, 2016 at 01:41:50PM -0600, James Norris wrote:
> The attached patch resolves c/PR64748. The patch
> adds the use of parm's with the deviceptr clause.
> 
> Question....
> 
> As there is VAR_P (), could there be a PARM_P ()?

Not for GCC 6.x, for 7 it is possible.

> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -10760,7 +10760,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 (!VAR_P (v))
> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))

Please don't write !(x == y) but x != y.

> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> @@ -30087,7 +30087,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 (!VAR_P (v))
> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))
>  	error_at (loc, "%qD is not a variable", v);
>        else if (TREE_TYPE (v) == error_mark_node)
>  	;

For C++, all this diagnostics is premature, if processing_template_decl
you really often don't know what the type will be, not sure if you always
know at least if it is a VAR_DECL, PARM_DECL or something else.  I bet you
can easily ICE with the current POINTER_TYPE_P (TREE_TYPE (v)) check as
in templates the type can be NULL, or it could be some lang type and only
later on become POINTER_TYPE, etc.
For C++ the diagnostics need to be done during finish_omp_clauses or so, not
earlier.

	Jakub

Re: [PATCH] Fix PR64748

[James Norris]

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

Hi!

On 02/01/2016 02:03 PM, Jakub Jelinek wrote:
> On Mon, Feb 01, 2016 at 01:41:50PM -0600, James Norris wrote:
>> The attached patch resolves c/PR64748. The patch
>> adds the use of parm's with the deviceptr clause.
>>
>  [snip snip]
>> --- a/gcc/c/c-parser.c
>> +++ b/gcc/c/c-parser.c
>> @@ -10760,7 +10760,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 (!VAR_P (v))
>> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))
>
> Please don't write !(x == y) but x != y.

Fixed.

>
>> --- a/gcc/cp/parser.c
>> +++ b/gcc/cp/parser.c
>> @@ -30087,7 +30087,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 (!VAR_P (v))
>> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))
>>   	error_at (loc, "%qD is not a variable", v);
>>         else if (TREE_TYPE (v) == error_mark_node)
>>   	;
>
> For C++, all this diagnostics is premature, if processing_template_decl
> you really often don't know what the type will be, not sure if you always
> know at least if it is a VAR_DECL, PARM_DECL or something else.  I bet you
> can easily ICE with the current POINTER_TYPE_P (TREE_TYPE (v)) check as
> in templates the type can be NULL, or it could be some lang type and only
> later on become POINTER_TYPE, etc.
> For C++ the diagnostics need to be done during finish_omp_clauses or so, not
> earlier.

The check has been moved to finish_omp_clause (). I put the check at
the tail end of the checking, as I wasn't able to determine if there
was a checking precedence done by the if-else-if sequence.

Thanks for the review!

Jim


===== ChangeLog entries...

         gcc/testsuite/

         PR c/64748
         * c-c++-common/goacc/deviceptr-1.c: Add tests.
         * g++.dg/goacc/deviceptr-1.c: New file.


         gcc/cp/

         PR c/64748
         * parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking.
         * semantics.c (finish_omp_clauses): Add deviceptr checking.


         gcc/c/

         PR c/64748
         * c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.




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

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 5341f04..f2d114c 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.
+
 2016-01-27  Jakub Jelinek  <jakub@redhat.com>
 
 	PR debug/66869
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index eede3a7..229fd6e 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10760,7 +10760,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 (!VAR_P (v))
+      if (!VAR_P (v) && 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/ChangeLog b/gcc/cp/ChangeLog
index 3b5c9d5..76cf5b1 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking.
+	* semantics.c (finish_omp_clauses): Add deviceptr checking.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR debug/66869
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d03b0c9..10f3627 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -30080,20 +30080,6 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
   for (t = vars; t; t = TREE_CHAIN (t))
     {
       tree v = TREE_PURPOSE (t);
-
-      /* FIXME diagnostics: Ideally we should keep individual
-	 locations for all the variables in the var list to make the
-	 following errors more precise.  Perhaps
-	 c_parser_omp_var_list_parens should construct a list of
-	 locations to go along with the var list.  */
-
-      if (!VAR_P (v))
-	error_at (loc, "%qD is not a variable", v);
-      else if (TREE_TYPE (v) == error_mark_node)
-	;
-      else if (!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);
       OMP_CLAUSE_SET_MAP_KIND (u, GOMP_MAP_FORCE_DEVICEPTR);
       OMP_CLAUSE_DECL (u) = v;
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 95c4f19..1e376b1 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6683,6 +6683,14 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	      error ("%qD appears both in data and map clauses", t);
 	      remove = true;
 	    }
+	  else if (!processing_template_decl
+		   && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR
+		   && !POINTER_TYPE_P (TREE_TYPE (t)))
+	    {
+	      error ("%qD is not a pointer variable", t);
+	      remove = true;
+	    }
 	  else
 	    {
 	      bitmap_set_bit (&map_head, DECL_UID (t));
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 150ebc8..ce07496 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2016-02-XX  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* c-c++-common/goacc/deviceptr-1.c: Add tests.
+	* g++.dg/goacc/deviceptr-1.c: New file.
+
 2016-01-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/69551
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 546fa82..6edbdb1 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -84,3 +84,21 @@ fun4 (void)
 #pragma acc parallel deviceptr(s2_p)
   s2_p = 0;
 }
+
+void
+func5 (float *fp)
+{
+
+#pragma acc data deviceptr (fp)
+{ }
+
+}
+
+void
+func6 (float fp)
+{
+
+#pragma acc data deviceptr (fp)	/* { dg-error "is not a pointer variable" } */
+{ }
+
+}
diff --git a/gcc/testsuite/g++.dg/goacc/deviceptr-1.C b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
new file mode 100644
index 0000000..7879790
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
@@ -0,0 +1,28 @@
+// { dg-do compile }
+
+template <typename P>
+
+void
+func1 (P p)
+{
+
+#pragma acc data deviceptr (p)	// { dg-error "is not a pointer" }
+{ }
+
+}
+
+void
+func2 (void)
+{
+  int *p;
+
+  func1 <int*>(p);
+}
+
+void
+func3 (void)
+{
+  int p;
+
+  func1 <int>(p);
+}

Re: [PATCH] Fix PR64748

[James Norris]

Hi,


Ping!

Thanks,
Jim


On 02/02/2016 08:51 AM, James Norris wrote:
> Hi!
>
> On 02/01/2016 02:03 PM, Jakub Jelinek wrote:
>> On Mon, Feb 01, 2016 at 01:41:50PM -0600, James Norris wrote:
>>> The attached patch resolves c/PR64748. The patch
>>> adds the use of parm's with the deviceptr clause.
>>>
>>  [snip snip]
>>> --- a/gcc/c/c-parser.c
>>> +++ b/gcc/c/c-parser.c
>>> @@ -10760,7 +10760,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 (!VAR_P (v))
>>> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))
>>
>> Please don't write !(x == y) but x != y.
>
> Fixed.
>
>>
>>> --- a/gcc/cp/parser.c
>>> +++ b/gcc/cp/parser.c
>>> @@ -30087,7 +30087,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 (!VAR_P (v))
>>> +      if (!VAR_P (v) && !(TREE_CODE (v) == PARM_DECL))
>>>       error_at (loc, "%qD is not a variable", v);
>>>         else if (TREE_TYPE (v) == error_mark_node)
>>>       ;
>>
>> For C++, all this diagnostics is premature, if processing_template_decl
>> you really often don't know what the type will be, not sure if you always
>> know at least if it is a VAR_DECL, PARM_DECL or something else.  I bet you
>> can easily ICE with the current POINTER_TYPE_P (TREE_TYPE (v)) check as
>> in templates the type can be NULL, or it could be some lang type and only
>> later on become POINTER_TYPE, etc.
>> For C++ the diagnostics need to be done during finish_omp_clauses or so, not
>> earlier.
>
> The check has been moved to finish_omp_clause (). I put the check at
> the tail end of the checking, as I wasn't able to determine if there
> was a checking precedence done by the if-else-if sequence.
>
> Thanks for the review!
>
> Jim
>
>
> ===== ChangeLog entries...
>
>          gcc/testsuite/
>
>          PR c/64748
>          * c-c++-common/goacc/deviceptr-1.c: Add tests.
>          * g++.dg/goacc/deviceptr-1.c: New file.
>
>
>          gcc/cp/
>
>          PR c/64748
>          * parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking.
>          * semantics.c (finish_omp_clauses): Add deviceptr checking.
>
>
>          gcc/c/
>
>          PR c/64748
>          * c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.
>
>
>

Re: [PATCH] Fix PR64748

[Jakub Jelinek]

On Tue, Feb 02, 2016 at 08:51:23AM -0600, James Norris wrote:
> --- a/gcc/cp/semantics.c
> +++ b/gcc/cp/semantics.c
> @@ -6683,6 +6683,14 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
>  	      error ("%qD appears both in data and map clauses", t);
>  	      remove = true;
>  	    }
> +	  else if (!processing_template_decl
> +		   && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR
> +		   && !POINTER_TYPE_P (TREE_TYPE (t)))
> +	    {
> +	      error ("%qD is not a pointer variable", t);
> +	      remove = true;
> +	    }

Please move this a few lines up, before the first duplicate check, thus
above
          else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
                   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
Also, testing it only for !processing_template_decl is undesirable, then you
can't diagnose obvious issues in non-instantiated templates.  Better use:

else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
	 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR
	 && !type_dependent_expression_p (t)
	 && !POINTER_TYPE_P (TREE_TYPE (t)))

> --- /dev/null
> +++ b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
> @@ -0,0 +1,28 @@
> +// { dg-do compile }
> +
> +template <typename P>
> +
> +void
> +func1 (P p)
> +{
> +

Please avoid the unnecessary empty lines above (both of them).

> +#pragma acc data deviceptr (p)	// { dg-error "is not a pointer" }
> +{ }
> +

And here too.  Perhaps use "  ;" instead of "{ }"?  And, more importantly,
by using a single template and instantiating it with both arguments, you are
not testing that you are not diagnosing it for the pointer case.

> +}
> +
> +void
> +func2 (void)
> +{
> +  int *p;
> +
> +  func1 <int*>(p);
> +}
> +
> +void
> +func3 (void)
> +{
> +  int p;
> +
> +  func1 <int>(p);
> +}

Also, I don't like the uses of uninitialized vars.
So better

template <typename P>
void
func1 (P p)
{
#pragma acc data deviceptr (p)	// { dg-bogus "is not a pointer" }
  ;
}

void
func2 (int *p)
{
  func1 (p);
}

template <typename P>
void
func3 (P p)
{
#pragma acc data deviceptr (p)	// { dg-error "is not a pointer" }
  ;
}

void
func4 (int p)
{
  func3 (p);
}

template <int N>
void
func5 (int *p, int q)
{
#pragma acc data deviceptr (p)	// { dg-bogus "is not a pointer" }
  ;
#pragma acc data deviceptr (q)	// { dg-error "is not a pointer" }
  ;
}

func5 added so to test that you diagnose even uninstantiated templates
if the vars/parameters are not type dependent.

Ok for trunk with those changes.

	Jakub

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

[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

[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

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

[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;
 }