From: Thomas Schwinge <thomas@codesourcery.com>
To: Chung-Lin Tang <cltang@codesourcery.com>
Cc: <gcc-patches@gcc.gnu.org>, Catherine Moore <clm@codesourcery.com>
Subject: Re: [PATCH, OpenACC 2.7, v2] Implement host_data must have use_device clause requirement
Date: Thu, 20 Jul 2023 12:00:21 +0200 [thread overview]
Message-ID: <878rbac37u.fsf@euler.schwinge.homeip.net> (raw)
In-Reply-To: <aa222b94-ad9e-8c44-f99b-fbd8b8dc9d82@siemens.com>
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
prev parent reply other threads:[~2023-07-20 10:00 UTC|newest]
Thread overview: 4+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-06-06 15:10 [PATCH, OpenACC 2.7] " Chung-Lin Tang
2023-06-16 9:13 ` Thomas Schwinge
2023-07-13 10:54 ` [PATCH, OpenACC 2.7, v2] " Chung-Lin Tang
2023-07-20 10:00 ` Thomas Schwinge [this message]
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=878rbac37u.fsf@euler.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=clm@codesourcery.com \
--cc=cltang@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).