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.
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 <jnor...@codesourcery.com>
+
+ PR c/64748
+ * c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms.
+
2016-01-27 Jakub Jelinek <ja...@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 <jnor...@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 <ja...@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 <jnor...@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 <ja...@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);
+}
