From: Chung-Lin Tang <chunglin.tang@siemens.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>,
Thomas Schwinge <thomas@codesourcery.com>,
Catherine Moore <clm@codesourcery.com>
Subject: [PATCH, OpenACC 2.7] Implement host_data must have use_device clause requirement
Date: Tue, 6 Jun 2023 23:10:37 +0800 [thread overview]
Message-ID: <3be2222f-48ae-12a1-a83b-415360e0a506@siemens.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 1238 bytes --]
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.
[-- Attachment #2: 0001-OpenACC-2.7-host_data-must-have-use_device-clause-re.patch --]
[-- Type: text/plain, Size: 5694 bytes --]
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
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
--
2.27.0
next reply other threads:[~2023-06-06 15:10 UTC|newest]
Thread overview: 4+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-06-06 15:10 Chung-Lin Tang [this message]
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
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=3be2222f-48ae-12a1-a83b-415360e0a506@siemens.com \
--to=chunglin.tang@siemens.com \
--cc=clm@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=thomas@codesourcery.com \
/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).