Message ID | 3be2222f-48ae-12a1-a83b-415360e0a506@siemens.com |
---|---|
State | New |
Headers | show |
Series | [OpenACC,2.7] Implement host_data must have use_device clause requirement | expand |
Hi Chung-Lin! On 2023-06-06T23:10:37+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote: > this patch implements the OpenACC 2.7 change requiring the host_data construct > to have at least one use_device clause. Thanks! > This patch started out with a simple check during gimplify (much smaller patch), Heh, thanks for the explanation -- would've been my first question otherweise. ;-) > but turned out that front-ends removed use_device clauses when they have error, > and the gimplify check started to echo a "no use_device clause" message in such > cases, which seem confusing for the user. So ended up adding the check in each > front-end instead. I presume that's also the reason why you're doing this check before 'c_finish_omp_clauses' etc.? I'll clarify with the OpenACC Technical Committee whether really those diagnostics are intended as "error" or should instead just be "warning". After all, there's no actual problem with an OpenACC 'host_data' without 'use_device' clause (or no data clause on OpenACC 'data', 'enter data', 'exit data', 'update', etc.) -- it's just likely that the user missed something. That is, the OpenACC 2.7: "At least one 'use_device' clause must appear" is addressing the user, not at the implementation (in my current interpretation). Depending on the outcome of that, we can easily adjust GCC. Note for later, independently of your work here: 'c_parser_oacc_enter_exit_data' etc. for its corresponding "has no data movement clause" diagnostic actually does 'c_finish_omp_clauses' etc. first -- maybe that should be changed accordingly. (Actually, I note that it's only OpenACC 3.0 that "Required at least one data clause on a 'data' construct, an 'enter data' directive, or an 'exit data' directive", heh... Per his internal 2014-10-17 email, Cesar implemented the code of 'c_parser_oacc_enter_exit_data' etc. "similar to that of acc update", which indeed already back then did require "At least one 'self', 'host', or 'device' clause". Fortran does have the diagnostic for OpenACC 'update', but it's missing for OpenACC 'enter data', 'exit data' without data clause (have not checked other constructs with similar requirements).) > Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect > no surprises). Is this okay for trunk? OK with one small change, please -- unless there's a reason for doing it this way: > --- a/gcc/fortran/trans-openmp.cc > +++ b/gcc/fortran/trans-openmp.cc > @@ -4677,6 +4677,12 @@ gfc_trans_oacc_construct (gfc_code *code) > break; > case EXEC_OACC_HOST_DATA: > construct_code = OACC_HOST_DATA; > + if (code->ext.omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) > + { > + error_at (gfc_get_location (&code->loc), > + "%<host_data%> construct requires %<use_device%> clause"); > + return NULL_TREE; > + } > break; > default: > gcc_unreachable (); The OpenMP "must contain at least one [...] clause" checks are done in 'gcc/fortran/openmp.cc:resolve_omp_clauses'. For consistency (or, to let 'gcc/fortran/trans-openmp.cc' continue to just deal with "directive translation"), do similar for OpenACC 'host_data'? (..., and we later accordingly adjust 'gcc/fortran/openmp.cc:gfc_match_oacc_update', too?) Grüße Thomas ----------------- 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 5baa501dbee..b61aef8b1a2 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18398,8 +18398,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 1c9aa671851..dd7638f1c93 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45798,8 +45798,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/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 42b608f3d36..5e0079cce76 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4677,6 +4677,12 @@ gfc_trans_oacc_construct (gfc_code *code) break; case EXEC_OACC_HOST_DATA: construct_code = OACC_HOST_DATA; + if (code->ext.omp_clauses->lists[OMP_LIST_USE_DEVICE] == NULL) + { + error_at (gfc_get_location (&code->loc), + "%<host_data%> construct requires %<use_device%> clause"); + return NULL_TREE; + } break; default: gcc_unreachable (); 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..747a2b201e7 --- /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 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
Hi Thomas, this patch implements the OpenACC 2.7 change requiring the host_data construct to have at least one use_device clause. This patch started out with a simple check during gimplify (much smaller patch), but turned out that front-ends removed use_device clauses when they have error, and the gimplify check started to echo a "no use_device clause" message in such cases, which seem confusing for the user. So ended up adding the check in each front-end instead. Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect no surprises). Is this okay for trunk? Thanks, Chung-Lin 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: * trans-openmp.cc (gfc_trans_oacc_construct): 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. From 0d17b8d24fa6079d6c289305e9644c3fecd429f1 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang <cltang@codesourcery.com> Date: Tue, 6 Jun 2023 03:19:33 -0700 Subject: [PATCH 1/2] OpenACC 2.7: host_data must have use_device clause requirement This patch implements the OpenACC 2.7 change requiring the host_data construct to have at least one use_device clause. 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: * trans-openmp.cc (gfc_trans_oacc_construct): 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. --- gcc/c/c-parser.cc | 9 +++++++-- gcc/cp/parser.cc | 11 +++++++++-- gcc/fortran/trans-openmp.cc | 6 ++++++ gcc/testsuite/c-c++-common/goacc/host_data-2.c | 7 ++++++- gcc/testsuite/gfortran.dg/goacc/host_data-error.f90 | 6 ++++++ gcc/testsuite/gfortran.dg/goacc/pr71704.f90 | 5 +++-- 6 files changed, 37 insertions(+), 7 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/host_data-error.f90