re PR c/64748 (OpenACC: "is not a variable" error with deviceptr())

	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.

From-SVN: r233458

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 0f3e7560056a72a048a6501a2a14505e19fb9cce..b51957ee87e3df652d4787db09af61a1ad8a2afe 100644 (file)
--- 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 7a272447900a9a0305a644bd0ca2f0e16738e7e6..23853be205fc045adbb50019ffb7f35f31121340 100644 (file)
--- 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 51c99e1759889624895e1d29eb5ec90ead8a0413..743e059956a60c6c4a365381b6582d6837d13849 100644 (file)
--- 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 07d182156d6375230fbeca6f5a6023fbd68f87de..b8d823787eca4b0c210c36d6f9feb027e0e31a10 100644 (file)
--- 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 0f6a6b5c539f752fff38a1a3c3b457877d08f889..70a7aa5bdf517bb32918ad198189f09bd0b8f9ee 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6634,6 +6634,14 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
                        omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
              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)
            {

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 242ed9717e33cf29feb691706bcef2e6265310c0..e78c3c1b5ae3808f676659aefbebca0a138feb89 100644 (file)
--- 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 546fa82958e6ca15114d97081ef3d345ea392e95..08ddb1072ff8e4b6c8c966cf5682309e16964928 100644 (file)
--- 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 (file)
index 0000000..d6d0483
--- /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" }
+  ;
+}