From: Thomas Schwinge Date: Wed, 2 Dec 2015 15:53:34 +0000 (+0100) Subject: Some OpenACC host_data cleanup X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=ff7a55bf56a6f148b29c8fd44ca90c4e535fc2bd;p=gcc.git Some OpenACC host_data cleanup gcc/c/ * c-parser.c (c_parser_omp_clause_name) (c_parser_oacc_all_clauses): Alphabetical sorting. gcc/cp/ * parser.c (cp_parser_omp_clause_name) (cp_parser_oacc_all_clauses): Alphabetical sorting. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE. gcc/testsuite/ * c-c++-common/goacc/host_data-5.c: New file. * c-c++-common/goacc/host_data-6.c: Likewise. * gfortran.dg/goacc/coarray.f95: XFAIL. * gfortran.dg/goacc/coarray_2.f90: Adjust dg-excess-errors directive. * gfortran.dg/goacc/host_data-tree.f95: Remove dg-prune-output directive. libgomp/ * testsuite/libgomp.oacc-c-c++-common/host_data-2.c: Restrict to target openacc_nvidia_accel_selected. * testsuite/libgomp.oacc-c-c++-common/host_data-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-3.c: Remove file. * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: Remove file. Co-Authored-By: James Norris Co-Authored-By: Julian Brown From-SVN: r231184 --- diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index acb8ee4e6cc..e517467bd6a 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,8 @@ +2015-12-02 Thomas Schwinge + + * c-parser.c (c_parser_omp_clause_name) + (c_parser_oacc_all_clauses): Alphabetical sorting. + 2015-12-02 Jakub Jelinek PR c/68533 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index d4c512fee04..ee0a30517af 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10277,10 +10277,10 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; - else if (!strcmp ("use_device_ptr", p)) - result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; + else if (!strcmp ("use_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; case 'v': if (!strcmp ("vector", p)) @@ -12951,10 +12951,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; - case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = c_parser_oacc_clause_use_device (parser, clauses); - c_name = "use_device"; - break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses); @@ -12964,6 +12960,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_clause_tile (parser, clauses); c_name = "tile"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = c_parser_oacc_clause_use_device (parser, clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 385ba635c4b..d2a7e99c63f 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,9 @@ +2015-12-02 Thomas Schwinge + + * parser.c (cp_parser_omp_clause_name) + (cp_parser_oacc_all_clauses): Alphabetical sorting. + * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE. + 2015-12-02 Andreas Arnez PR gcov-profile/68603 diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index f78df028193..b4ecac7acbd 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29230,10 +29230,10 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; - else if (!strcmp ("use_device_ptr", p)) - result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; + else if (!strcmp ("use_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; case 'v': if (!strcmp ("vector", p)) @@ -31600,11 +31600,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; - case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, - clauses); - c_name = "use_device"; - break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses, here); @@ -31614,6 +31609,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_clause_tile (parser, here, clauses); c_name = "tile"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, @@ -34516,13 +34516,13 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) return stmt; } -#define OACC_HOST_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) - /* OpenACC 2.0: # pragma acc host_data new-line structured-block */ +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + static tree cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok) { diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 5befd644c09..d1d1e4e9a6c 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -14387,6 +14387,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: case OMP_CLAUSE_MAP: + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: OMP_CLAUSE_DECL (nc) @@ -14513,6 +14514,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: /* tsubst_expr on SCOPE_REF results in returning diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 5fe26bb9aa0..4f7af876162 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,18 @@ +2015-12-02 Thomas Schwinge + + * gfortran.dg/goacc/coarray.f95: XFAIL. + * gfortran.dg/goacc/coarray_2.f90: Adjust dg-excess-errors + directive. + * gfortran.dg/goacc/host_data-tree.f95: Remove dg-prune-output + directive. + +2015-12-02 Thomas Schwinge + Julian Brown + James Norris + + * c-c++-common/goacc/host_data-5.c: New file. + * c-c++-common/goacc/host_data-6.c: Likewise. + 2015-12-02 Tom de Vries * c-c++-common/goacc/kernels-default-2.c: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-5.c b/gcc/testsuite/c-c++-common/goacc/host_data-5.c new file mode 100644 index 00000000000..f372fbdc437 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/host_data-5.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N]; + +#pragma acc data copyin (x[0:N]) + { + int *xp; +#pragma acc host_data use_device (x) + { + /* This use of the present clause is undefined behaviour for OpenACC. */ +#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ + { + xp = x; + } + } + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-6.c b/gcc/testsuite/c-c++-common/goacc/host_data-6.c new file mode 100644 index 00000000000..8be7912e280 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/host_data-6.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N]; + +#pragma acc data copyin (x[0:N]) + { + int *xp; +#pragma acc host_data use_device (x) + { + /* Here 'x' being implicitly firstprivate for the parallel region + conflicts with it being declared as use_device in the enclosing + host_data region. */ +#pragma acc parallel copyout (xp) + { + xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ + } + } + } + + return 0; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/coarray.f95 b/gcc/testsuite/gfortran.dg/goacc/coarray.f95 index 130ffc3ce9d..d2f10d54db5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/coarray.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/coarray.f95 @@ -1,7 +1,9 @@ ! { dg-do compile } ! { dg-additional-options "-fcoarray=single" } - -! TODO: These cases must fail +! +! PR fortran/63861 +! { dg-xfail-if "" { *-*-* } } +! { dg-excess-errors "TODO" } module test contains @@ -9,7 +11,6 @@ contains implicit none integer :: i integer, codimension[*] :: a - ! { dg-excess-errors "sorry, unimplemented: directive not yet implemented" } !$acc declare device_resident (a) !$acc data copy (a) !$acc end data @@ -17,7 +18,6 @@ contains !$acc end data !$acc parallel private (a) !$acc end parallel - ! { dg-excess-errors "sorry, unimplemented: directive not yet implemented" } !$acc host_data use_device (a) !$acc end host_data !$acc parallel loop reduction(+:a) diff --git a/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 b/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 index f9cf9ac6ca0..87e04d58bee 100644 --- a/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 @@ -3,6 +3,7 @@ ! ! PR fortran/63861 ! { dg-xfail-if "" { *-*-* } } +! { dg-excess-errors "TODO" } module test contains @@ -106,4 +107,3 @@ contains !$acc update self (a) end subroutine oacc4 end module test -! { dg-excess-errors "sorry, unimplemented: directive not yet implemented" } diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 index e4c8205a591..7a5eea6d287 100644 --- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 @@ -8,5 +8,4 @@ program test !$acc host_data use_device(i) !$acc end host_data end program test -! { dg-prune-output "unimplemented" } ! { dg-final { scan-tree-dump-times "pragma acc host_data use_device\\(i\\)" 1 "original" } } diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ddf836a078b..cde0b5c55a3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,12 @@ +2015-12-02 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/host_data-2.c: Restrict to + target openacc_nvidia_accel_selected. + * testsuite/libgomp.oacc-c-c++-common/host_data-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/host_data-5.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/host_data-3.c: Remove file. + * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: Remove file. + 2015-12-01 Julian Brown James Norris diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c index 98202867b4d..614f14324d4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c @@ -1,4 +1,4 @@ -/* { dg-do run } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c deleted file mode 100644 index 7d9b5f72667..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c +++ /dev/null @@ -1,29 +0,0 @@ -/* { dg-do compile } */ - -#include -#include - -#define N 1024 - -int main (int argc, char* argv[]) -{ - int x[N]; - -#pragma acc data copyin (x[0:N]) - { - int *xp; -#pragma acc host_data use_device (x) - { - /* This use of the present clause is undefined behaviour for OpenACC. */ -#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ - { - xp = x; - } - } - - if (xp != acc_deviceptr (x)) - abort (); - } - - return 0; -} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c index 3504f2710fd..0ab5a356912 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c @@ -1,4 +1,4 @@ -/* { dg-do run } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c index 268e9194d35..a3737a7dae0 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c @@ -1,4 +1,4 @@ -/* { dg-do run } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c deleted file mode 100644 index a841488515e..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c +++ /dev/null @@ -1,31 +0,0 @@ -/* { dg-do compile } */ - -#include -#include - -#define N 1024 - -int main (int argc, char* argv[]) -{ - int x[N]; - -#pragma acc data copyin (x[0:N]) - { - int *xp; -#pragma acc host_data use_device (x) - { - /* Here 'x' being implicitly firstprivate for the parallel region - conflicts with it being declared as use_device in the enclosing - host_data region. */ -#pragma acc parallel copyout (xp) - { - xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ - } - } - - if (xp != acc_deviceptr (x)) - abort (); - } - - return 0; -}