aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
authorChung-Lin Tang <cltang@codesourcery.com>2023-07-31 07:53:24 -0700
committerChung-Lin Tang <cltang@codesourcery.com>2023-07-31 07:56:19 -0700
commita104e9ac0ae9a7e78ec2edd0b81074946646a87d (patch)
tree3af10f8605d928435401ca81286f4b95c7d10c01 /gcc
parent7cdd0860949c6c3232e6cff1d7ca37bb5234074c (diff)
downloadgcc-a104e9ac0ae9a7e78ec2edd0b81074946646a87d.zip
gcc-a104e9ac0ae9a7e78ec2edd0b81074946646a87d.tar.gz
gcc-a104e9ac0ae9a7e78ec2edd0b81074946646a87d.tar.bz2
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: * 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.
Diffstat (limited to 'gcc')
-rw-r--r--gcc/c/c-parser.cc9
-rw-r--r--gcc/cp/parser.cc11
-rw-r--r--gcc/fortran/openmp.cc6
-rw-r--r--gcc/testsuite/c-c++-common/goacc/host_data-2.c7
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/host_data-error.f906
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/pr71704.f905
6 files changed, 37 insertions, 7 deletions
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 24a6eb6..80920b3 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 d7ef5b3..b1d2e14 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45899,8 +45899,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 2952cd3..234d896 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -8926,6 +8926,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 b3093e5..862a764 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 0000000..bd26298
--- /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 0235e85..31724c8 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