PR c/64748

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

	gcc/cp/
	* parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking.
	* semantics.c (finish_omp_clauses): Add deviceptr checking.

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


git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@233458 138bc75d-0d04-0410-961f-82ee72b054a4

8 files changed, 76 insertions, 15 deletions

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 0f3e7560056..b51957ee87e 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-16  James Norris  <jnorris@codesourcery.com>
+
+	PR c/64748
+	* c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.
+
 2016-02-12  Bernd Schmidt  <bschmidt@redhat.com>
 
 	* c-decl.c (build_null_declspecs): Zero the entire struct.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 7a272447900..23853be205f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10766,7 +10766,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 51c99e17598..743e059956a 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-16  James Norris  <jnorris@codesourcery.com>
+
+	* parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking.
+	* semantics.c (finish_omp_clauses): Add deviceptr checking.
+
 2016-02-15  Jakub Jelinek  <jakub@redhat.com>
 
 	PR c++/69658
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 07d182156d6..b8d823787ec 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -30105,20 +30105,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 0f6a6b5c539..70a7aa5bdf5 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6635,6 +6635,14 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	      remove = true;
 	    }
 	  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)))
+	    {
+	      error ("%qD is not a pointer variable", t);
+	      remove = true;
+	    }
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
 	      if (bitmap_bit_p (&generic_head, DECL_UID (t))
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 242ed9717e3..e78c3c1b5ae 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2016-02-16  James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/deviceptr-1.c: Add tests.
+	* g++.dg/goacc/deviceptr-1.c: New file.
+
 2016-02-16  Jakub Jelinek  <jakub@redhat.com>
 
 	PR tree-optimization/69820
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 546fa82958e..08ddb1072ff 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -84,3 +84,17 @@ 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 00000000000..d6d0483129b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
@@ -0,0 +1,38 @@
+// { dg-do compile }
+
+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" }
+  ;
+}