| Message ID | 56B0C26B.7020200@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
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. > > >
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
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); +}