Message ID | aa222b94-ad9e-8c44-f99b-fbd8b8dc9d82@siemens.com |
---|---|
State | New |
Headers | show |
Series | [OpenACC,2.7,v2] Implement host_data must have use_device clause requirement | expand |
Hi Chung-Lin! On 2023-07-13T18:54:00+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote: > On 2023/6/16 5:13 PM, Thomas Schwinge wrote: >> OK with one small change, please -- unless there's a reason for doing it >> this way: [...] > I've adjusted the Fortran implementation as you described. Yes, I agree this way > more fits current Fortran FE conventions. > > I've re-tested the attached v2 patch, will commit later this week if no major > objections. ACK, thanks. Grüße Thomas > gcc/c/ChangeLog: > > * c-parser.cc (c_parser_oacc_host_data): Add checking requiring OpenACC > host_data construct to have an use_device clause. > > gcc/cp/ChangeLog: > > * parser.cc (cp_parser_oacc_host_data): Add checking requiring OpenACC > host_data construct to have an use_device clause. > > gcc/fortran/ChangeLog: > > * openmp.cc (resolve_omp_clauses): Add checking requiring > OpenACC host_data construct to have an use_device clause. > > gcc/testsuite/ChangeLog: > > * c-c++-common/goacc/host_data-2.c: Adjust testcase. > * gfortran.dg/goacc/host_data-error.f90: New testcase. > * gfortran.dg/goacc/pr71704.f90: Adjust testcase. > diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc > index 24a6eb6e459..80920b31f83 100644 > --- a/gcc/c/c-parser.cc > +++ b/gcc/c/c-parser.cc > @@ -18461,8 +18461,13 @@ c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p) > tree stmt, clauses, block; > > clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, > - "#pragma acc host_data"); > - > + "#pragma acc host_data", false); > + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) > + { > + error_at (loc, "%<host_data%> construct requires %<use_device%> clause"); > + return error_mark_node; > + } > + clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); > block = c_begin_omp_parallel (); > add_stmt (c_parser_omp_structured_block (parser, if_p)); > stmt = c_finish_oacc_host_data (loc, clauses, block); > diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc > index 5e2b5cba57e..beb5b632e5e 100644 > --- a/gcc/cp/parser.cc > +++ b/gcc/cp/parser.cc > @@ -45895,8 +45895,15 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) > unsigned int save; > > clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, > - "#pragma acc host_data", pragma_tok); > - > + "#pragma acc host_data", pragma_tok, > + false); > + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) > + { > + error_at (pragma_tok->location, > + "%<host_data%> construct requires %<use_device%> clause"); > + return error_mark_node; > + } > + clauses = finish_omp_clauses (clauses, C_ORT_ACC); > block = begin_omp_parallel (); > save = cp_parser_begin_omp_structured_block (parser); > cp_parser_statement (parser, NULL_TREE, false, if_p); > diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc > index 8efc4b3ecfa..f7af02845de 100644 > --- a/gcc/fortran/openmp.cc > +++ b/gcc/fortran/openmp.cc > @@ -8764,6 +8764,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, > "%<MERGEABLE%> clause", &omp_clauses->detach->where); > } > > + if (openacc > + && code->op == EXEC_OACC_HOST_DATA > + && omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) > + gfc_error ("%<host_data%> construct at %L requires %<use_device%> clause", > + &code->loc); > + > if (omp_clauses->assume) > gfc_resolve_omp_assumptions (omp_clauses->assume); > } > diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-2.c b/gcc/testsuite/c-c++-common/goacc/host_data-2.c > index b3093e575ff..862a764eb3a 100644 > --- a/gcc/testsuite/c-c++-common/goacc/host_data-2.c > +++ b/gcc/testsuite/c-c++-common/goacc/host_data-2.c > @@ -8,7 +8,9 @@ void > f (void) > { > int v2 = 3; > -#pragma acc host_data copy(v2) /* { dg-error ".copy. is not valid for ..pragma acc host_data." } */ > +#pragma acc host_data copy(v2) > + /* { dg-error ".copy. is not valid for ..pragma acc host_data." "" { target *-*-* } .-1 } */ > + /* { dg-error ".host_data. construct requires .use_device. clause" "" { target *-*-* } .-2 } */ > ; > > #pragma acc host_data use_device(v2) > @@ -20,6 +22,9 @@ f (void) > /* { dg-error ".use_device_ptr. variable is neither a pointer nor an array" "" { target c } .-1 } */ > /* { dg-error ".use_device_ptr. variable is neither a pointer, nor an array nor reference to pointer or array" "" { target c++ } .-2 } */ > ; > + > +#pragma acc host_data /* { dg-error ".host_data. construct requires .use_device. clause" } */ > + ; > } > > > diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 > new file mode 100644 > index 00000000000..bd262989410 > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 > @@ -0,0 +1,6 @@ > +! { dg-do compile } > + > +subroutine foo () > +!$acc host_data ! { dg-error "'host_data' construct at .1. requires 'use_device' clause" } > +!$acc end host_data > +end > diff --git a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 > index 0235e85d42a..31724c8b046 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 > @@ -47,8 +47,9 @@ real function f8 () > f8 = 1 > end > > -real function f9 () > -!$acc host_data > +real function f9 (a) > + integer a(:) > +!$acc host_data use_device(a) > !$acc end host_data > f8 = 1 > end ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 24a6eb6e459..80920b31f83 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18461,8 +18461,13 @@ c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p) tree stmt, clauses, block; clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, - "#pragma acc host_data"); - + "#pragma acc host_data", false); + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) + { + error_at (loc, "%<host_data%> construct requires %<use_device%> clause"); + return error_mark_node; + } + clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser, if_p)); stmt = c_finish_oacc_host_data (loc, clauses, block); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 5e2b5cba57e..beb5b632e5e 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45895,8 +45895,15 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) unsigned int save; clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, - "#pragma acc host_data", pragma_tok); - + "#pragma acc host_data", pragma_tok, + false); + if (!omp_find_clause (clauses, OMP_CLAUSE_USE_DEVICE_PTR)) + { + error_at (pragma_tok->location, + "%<host_data%> construct requires %<use_device%> clause"); + return error_mark_node; + } + clauses = finish_omp_clauses (clauses, C_ORT_ACC); block = begin_omp_parallel (); save = cp_parser_begin_omp_structured_block (parser); cp_parser_statement (parser, NULL_TREE, false, if_p); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 8efc4b3ecfa..f7af02845de 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -8764,6 +8764,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "%<MERGEABLE%> clause", &omp_clauses->detach->where); } + if (openacc + && code->op == EXEC_OACC_HOST_DATA + && omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) + gfc_error ("%<host_data%> construct at %L requires %<use_device%> clause", + &code->loc); + if (omp_clauses->assume) gfc_resolve_omp_assumptions (omp_clauses->assume); } diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-2.c b/gcc/testsuite/c-c++-common/goacc/host_data-2.c index b3093e575ff..862a764eb3a 100644 --- a/gcc/testsuite/c-c++-common/goacc/host_data-2.c +++ b/gcc/testsuite/c-c++-common/goacc/host_data-2.c @@ -8,7 +8,9 @@ void f (void) { int v2 = 3; -#pragma acc host_data copy(v2) /* { dg-error ".copy. is not valid for ..pragma acc host_data." } */ +#pragma acc host_data copy(v2) + /* { dg-error ".copy. is not valid for ..pragma acc host_data." "" { target *-*-* } .-1 } */ + /* { dg-error ".host_data. construct requires .use_device. clause" "" { target *-*-* } .-2 } */ ; #pragma acc host_data use_device(v2) @@ -20,6 +22,9 @@ f (void) /* { dg-error ".use_device_ptr. variable is neither a pointer nor an array" "" { target c } .-1 } */ /* { dg-error ".use_device_ptr. variable is neither a pointer, nor an array nor reference to pointer or array" "" { target c++ } .-2 } */ ; + +#pragma acc host_data /* { dg-error ".host_data. construct requires .use_device. clause" } */ + ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 new file mode 100644 index 00000000000..bd262989410 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 @@ -0,0 +1,6 @@ +! { dg-do compile } + +subroutine foo () +!$acc host_data ! { dg-error "'host_data' construct at .1. requires 'use_device' clause" } +!$acc end host_data +end diff --git a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 index 0235e85d42a..31724c8b046 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 @@ -47,8 +47,9 @@ real function f8 () f8 = 1 end -real function f9 () -!$acc host_data +real function f9 (a) + integer a(:) +!$acc host_data use_device(a) !$acc end host_data f8 = 1 end