diff options
author | Cesar Philippidis <cesar@codesourcery.com> | 2015-11-05 18:03:48 -0800 |
---|---|---|
committer | Cesar Philippidis <cesar@gcc.gnu.org> | 2015-11-05 18:03:48 -0800 |
commit | 7a5e4956cc026cba54159d5c764486ac4151db85 (patch) | |
tree | 2f50525608c6a6a2a595620ad902d2efc85ab031 /gcc/testsuite/g++.dg/goacc | |
parent | 906f9ad995368b4891186c07ed3c5c2fd5f53bea (diff) | |
download | gcc-7a5e4956cc026cba54159d5c764486ac4151db85.zip gcc-7a5e4956cc026cba54159d5c764486ac4151db85.tar.gz gcc-7a5e4956cc026cba54159d5c764486ac4151db85.tar.bz2 |
gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Add support for
OMP_CLAUSE_TILE. Update handling of OMP_CLAUSE_INDEPENDENT.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE.
* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE.
* tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE.
(omp_clause_code_name): Likewise.
(walk_tree_1): Handle OMP_CLAUSE_TILE.
* tree.h (OMP_TILE_LIST): New macro.
gcc/c-family/
* c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR,
AUTO, SEQ, INDEPENDENT and PRIVATE loop clauses. Associate REDUCTION
clauses with parallel and kernels and loops.
* c-pragma.h (enum pragma_omp_clause): Add entries for
PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.
* pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS,
NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO,
INDEPENDENT,SEQ}.
(tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add support for
PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
(c_parser_omp_clause_default): Add is_oacc argument. Handle
default(none) in OpenACC.
(c_parser_oacc_shape_clause): Allow pointer variables as gang static
arguments.
(c_parser_oacc_clause_tile): New function.
(c_parser_oacc_all_clauses): Add support for OMP_CLAUSE_DEFAULT,
OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
TILE}.
(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
FIRSTPRIVATE}.
(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default.
(c_parser_oacc_update): Update the error message for missing clauses.
* c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE
and OMP_CLAUSE_INDEPENDENT.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Add support for
PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
(cp_parser_oacc_shape_clause): Allow pointer variables as gang static
arguments.
(cp_parser_oacc_clause_tile): New function.
(cp_parser_omp_clause_default): Add is_oacc argument. Handle
default(none) in OpenACC.
(cp_parser_oacc_all_clauses): Add support for
(cp_parser_omp_all_clauses): Update call to
cp_parser_omp_clause_default.
PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}.
(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
TILE}.
(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
FIRSTPRIVATE}.
(cp_parser_oacc_update): Update the error message for missing clauses.
* semantics.c (finish_omp_clauses): Add support for
OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
gcc/fortran/
* openmp.c (gfc_match_omp_clauses): Update support for the tile
and default clauses in OpenACC.
(gfc_match_oacc_update): Error when data clauses are supplied.
(oacc_compatible_clauses): Delete.
(resolve_omp_clauses): Give special care for OpenACC reductions.
Also update error reporting for the tile clause.
(resolve_oacc_loop_blocks): Update error reporting for the tile clause.
* trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add
OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument.
(gfc_trans_oacc_combined_directive): Update the list of clauses which
are split to acc loops.
gcc/testsuite/
* c-c++-common/goacc/combined-directives.c: New test.
* c-c++-common/goacc/loop-clauses.c: New test.
* c-c++-common/goacc/tile.c: New test.
* c-c++-common/goacc/loop-shape.c: Add test for pointer variable
as gang static arguments.
* c-c++-common/goacc/update-1.c: Adjust expected error message.
* g++.dg/goacc/template.C: New test.
* gfortran.dg/goacc/combined-directives.f90: New test.
* gfortran.dg/goacc/default.f95: New test.
* gfortran.dg/goacc/default_none.f95: New test.
* gfortran.dg/goacc/firstprivate-1.f95: New test.
* gfortran.dg/goacc/gang-static.f95: New test.
* gfortran.dg/goacc/kernels-loop-inner.f95: New test.
* gfortran.dg/goacc/kernels-loops-adjacent.f95: New test.
* gfortran.dg/goacc/list.f95: Update test.
* gfortran.dg/goacc/loop-2.f95: Likewise.
* gfortran.dg/goacc/loop-4.f95: New test.
* gfortran.dg/goacc/loop-5.f95: New test.
* gfortran.dg/goacc/loop-6.f95: New test.
* gfortran.dg/goacc/loop-tree-1.f90: Update test.
* gfortran.dg/goacc/multi-clause.f90: New test.
* gfortran.dg/goacc/parallel-tree.f95: Update test.
* gfortran.dg/goacc/update.f95: New test.
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>
From-SVN: r229832
Diffstat (limited to 'gcc/testsuite/g++.dg/goacc')
-rw-r--r-- | gcc/testsuite/g++.dg/goacc/template.C | 141 |
1 files changed, 141 insertions, 0 deletions
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C new file mode 100644 index 0000000..f7a717b --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -0,0 +1,141 @@ +// This error is temporary. Remove when support is added for these clauses +// in the middle end. Also remove the comments from the reduction test +// after the FE learns that reduction variables may appear in data clauses too. +// { dg-prune-output "sorry, unimplemented" } + +#pragma acc routine +template <typename T> T +accDouble(int val) +{ + return val * 2; +} + +template<typename T> T +oacc_parallel_copy (T a) +{ + T b = 0; + char w = 1; + int x = 2; + float y = 3; + double z = 4; + +#pragma acc parallel num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a) + { + b = a; + } + +#pragma acc parallel num_gangs (a) copy (w, x, y, z) + { + w = accDouble<char>(w); + x = accDouble<int>(x); + y = accDouble<float>(y); + z = accDouble<double>(z); + } + +#pragma acc parallel num_gangs (a) if (1) + { +#pragma acc loop auto tile (a, 3) + for (int i = 0; i < a; i++) + for (int j = 0; j < 5; j++) + b = a; + +#pragma acc loop seq + for (int i = 0; i < a; i++) + b = a; + } + + T c; + +#pragma acc parallel num_workers (10) + { +#pragma acc atomic capture + c = b++; + +#pragma atomic update + c++; + +#pragma acc atomic read + b = a; + +#pragma acc atomic write + b = a; + } + +//#pragma acc parallel reduction (+:c) +// { +// c = 1; +// } + +#pragma acc data if (1) copy (b) + { + #pragma acc parallel + { + b = a; + } + } + +#pragma acc enter data copyin (b) +#pragma acc parallel present (b) + { + b = a; + } + +#pragma acc update host (b) +#pragma acc update self (b) +#pragma acc update device (b) +#pragma acc exit data delete (b) + + return b; +} + +template<typename T> T +oacc_kernels_copy (T a) +{ + T b = 0; + T c = 0; + char w = 1; + int x = 2; + float y = 3; + double z = 4; + +#pragma acc kernels copy (w, x, y, z) + { + w = accDouble<char>(w); + x = accDouble<int>(x); + y = accDouble<float>(y); + z = accDouble<double>(z); + } + +#pragma acc kernels copyout (b) copyin (a) + b = a; + +//#pragma acc kernels loop reduction (+:c) +// for (int i = 0; i < 10; i++) +// { +// c = 1; +// } + +#pragma acc data if (1) copy (b) + { + #pragma acc kernels + { + b = a; + } + } + +#pragma acc enter data copyin (b) +#pragma acc kernels present (b) + { + b = a; + } + return b; +} + +int +main () +{ + int b = oacc_parallel_copy<int> (5); + int c = oacc_kernels_copy<int> (5); + + return b + c; +} |