From ba539195610794575f76e8dd2e16f548c4f89dbb Mon Sep 17 00:00:00 2001 From: James Norris Date: Tue, 16 Feb 2016 15:29:04 +0000 Subject: [PATCH] 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 --- gcc/c/ChangeLog | 5 +++ gcc/c/c-parser.c | 2 +- gcc/cp/ChangeLog | 5 +++ gcc/cp/parser.c | 14 ------- gcc/cp/semantics.c | 8 ++++ gcc/testsuite/ChangeLog | 5 +++ .../c-c++-common/goacc/deviceptr-1.c | 14 +++++++ gcc/testsuite/g++.dg/goacc/deviceptr-1.C | 38 +++++++++++++++++++ 8 files changed, 76 insertions(+), 15 deletions(-) create mode 100644 gcc/testsuite/g++.dg/goacc/deviceptr-1.C 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 + + PR c/64748 + * c-parser.c (c_parser_oacc_data_clause_deviceptr): Allow parms. + 2016-02-12 Bernd Schmidt * 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 + + * parser.c (cp_parser_oacc_data_clause_deviceptr): Remove checking. + * semantics.c (finish_omp_clauses): Add deviceptr checking. + 2016-02-15 Jakub Jelinek 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 @@ -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 242ed9717e3..e78c3c1b5ae 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2016-02-16 James Norris + + * c-c++-common/goacc/deviceptr-1.c: Add tests. + * g++.dg/goacc/deviceptr-1.c: New file. + 2016-02-16 Jakub Jelinek 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 +void +func1 (P p) +{ +#pragma acc data deviceptr (p)// { dg-bogus "is not a pointer" } + ; +} + +void +func2 (int *p) +{ + func1 (p); +} + +template +void +func3 (P p) +{ +#pragma acc data deviceptr (p)// { dg-error "is not a pointer" } + ; +} +void +func4 (int p) +{ + func3 (p); +} + +template +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" } + ; +} -- 2.30.2