| Message ID | 87mw1bkr7j.fsf@schwinge.name |
|---|---|
| State | New |
| Headers | show |
Hi! On Mon, 11 May 2015 18:35:12 +0200, I wrote: > On Tue, 05 May 2015 10:54:02 +0200, I wrote: > > In follow-up messages, I'll be posting the separated parts (for easier > > review) of a next set of OpenACC changes that we'd like to commit. > > ChangeLog updates not yet written; will do that before commit, obviously. > > In order for us to be able to make progress with staging our other > OpenACC changes in gomp-4_0-branch, I have now committed to > gomp-4_0-branch r223007, which is these patches as posted plus a tiny > last-minute typo fix (see below), and we shall then work on addressing > the review comments already provided (thanks!) (as well as those which I > found myself, upon reviewing our changes), before later re-submitting for > trunk. In a similar vein, I have now committed the following to gomp-4_0-branch in r223178. This is not meant to be integrated into trunk as-is: there are incompatible libgomp ABI changes, for example. We'd still appreciate any review comments, of course. To avoid running into mailing list size limits, the patch is attached as gcc-gomp-4_0-branch-r223178.patch.gz, so here's just the commit log: commit 631438a29e8d275570ba2e881ffcdea2dfe7b5e1 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Wed May 13 20:25:48 2015 +0000 Assorted OpenACC changes gcc/ada/ * gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. gcc/c-family/ * c-common.c (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. gcc/c/ * c-parser.c (c_parser_oacc_data_clause, oacc_split_loop_clauses) (c_parser_oacc_parallel): Handle PRAGMA_OACC_CLAUSE_FIRSTPRIVATE. (c_parser_oacc_all_clauses): Update handling of PRAGMA_OACC_CLAUSE_FIRSTPRIVATE. * c-typeck.c (c_finish_omp_clauses): Add error checking for GOMP_MAP_FORCE_TO_GANGLOCAL. gcc/cp/ * parser.c (cp_parser_oacc_data_clause) (cp_parser_oacc_all_clauses, oacc_split_loop_clauses) (cp_parser_oacc_parallel): Handle PRAGMA_OACC_CLAUSE_FIRSTPRIVATE. * semantics.c (finish_omp_clauses): Add error checking for GOMP_MAP_FORCE_TO_GANGLOCAL. gcc/fortran/ * f95-lang.c (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. * gfortran.h (gfc_omp_map_op): Add OMP_MAP_GANGLOCAL, OMP_MAP_FORCE_TO_GANGLOCAL. * trans-openmp.c (gfc_trans_omp_clauses): Handle them. * openmp.c (gfc_match_omp_clauses): Update handling of OMP_CLAUSE_FIRSTPRIVATE. * types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR): Remove. (BT_FN_INT_INT_INT_INT) (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR): New function types. gcc/jit/ * jit-builtins.c (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. * jit-builtins.h (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. gcc/lto/ * lto-lang.c (DEF_FUNCTION_TYPE_VAR_11): Remove. (DEF_FUNCTION_TYPE_VAR_12): New macro. gcc/testsuite/ * c-c++-common/goacc/dtype-1.c: Update. * c-c++-common/goacc/dtype-2.c: Likewise. * c-c++-common/goacc/host_data-1.c: Likewise. * c-c++-common/goacc/host_data-2.c: Likewise. * c-c++-common/goacc/host_data-3.c: Likewise. * c-c++-common/goacc/host_data-4.c: Likewise. * c-c++-common/goacc/sb-3.c: Likewise. * c-c++-common/goacc/tile.c: Likewise. * g++.dg/goacc/template-reduction.C: Likewise. * g++.dg/goacc/template.C: Likewise. * gfortran.dg/goacc/coarray.f95: Likewise. * gfortran.dg/goacc/dtype-1.f95: Likewise. * gfortran.dg/goacc/host_data-tree.f95: Likewise. * gfortran.dg/goacc/list.f95: Likewise. * gfortran.dg/goacc/loop-tree-1.f90: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * c-c++-common/goacc/executeables-1.c: New file. * c-c++-common/goacc/firstprivate.c: Likewise. * c-c++-common/goacc/private-reduction-1.c: Likewise. * g++.dg/goacc/loop-1.c: Likewise. * g++.dg/goacc/loop-2.c: Likewise. * g++.dg/goacc/loop-3.c: Likewise. * gcc.dg/goacc/sb-1.c: Likewise. * gcc.dg/goacc/sb-2.c: Likewise. * gcc.dg/goacc/sb-3.c: Likewise. * gfortran.dg/goacc/firstprivate-1.f95: Likewise. gcc/ * builtin-types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR): Remove. (BT_FN_INT_INT_INT_INT) (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR): New function types. * builtins.c (expand_oacc_builtin, expand_oacc_ganglocal_ptr): New functions. (expand_builtin, is_simple_builtin): Use them. * config/nvptx/mkoffload.c (process): Protect GOMP_offload_register prototype from C++ name mangling. * config/nvptx/nvptx.c (nvptx_file_start): Print declaration of sdata. * config/nvptx/nvptx.md (UNSPEC_NCTAID, UNSPEC_CTAID) (UNSPEC_SHARED_DATA): New constants. (oacc_nctaid_insn, oacc_nctaid, oacc_ctaid_insn, oacc_ctaid) (ganglocal_ptr<mode>, ganglocal_ptr): New patterns. * doc/md.texi (oacc_ntid, oacc_tid): Document. * gimple.h (gimple_statement_omp_parallel_layout): Add ganglocal_size member. (gimple_omp_target_ganglocal_size) (gimple_omp_target_set_ganglocal_size): New functions. * gimplify.c (gimplify_omp_var_data): Add GOVD_USE_DEVICE, GOVD_FORCE_MAP, GOVD_GANGLOCAL. (omp_region_type): Add ORT_HOST_DATA. (omp_region_kind, acc_region_kind): New enum types. (gimplify_omp_ctx): Add region_kind, acc_region_kind members. (new_omp_context, omp_add_variable, omp_notice_variable) (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses_1): Update for OpenACC. (gimplify_scan_omp_clauses): Add region_kind formal parameter. Adjust all users. (gimplify_oacc_host_data_1, gimplify_oacc_host_data): New functions. (gimplify_expr): Update handling of OACC_HOST_DATA. * omp-builtins.def (BUILT_IN_GOACC_KERNELS_INTERNAL) (BUILT_IN_GOACC_KERNELS, BUILT_IN_GOACC_PARALLEL): Change type from BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR to BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR. Adjust all users. (BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS): Change type from from BT_FN_INT to BT_FN_INT_INT_INT_INT. Adjust all users. (BUILT_IN_GOACC_NTID, BUILT_IN_GOACC_TID, BUILT_IN_GOACC_NCTAID) (BUILT_IN_GOACC_CTAID, BUILT_IN_GOACC_GET_GANGLOCAL_PTR) (BUILT_IN_GOACC_DEVICEPTR): New builtins. * omp-low.c (omp_context): Add oacc_reduction_set, ganglocal_init, ganglocal_ptr, ganglocal_size, ganglocal_size_host members. (new_omp_context, delete_omp_context): Initialize/deinitialize these, respectively. (omp_for_data): Add gang, worker, vector members. (extract_omp_for_data): Populate these. (oacc_max_threads, oacc_finalize_reduction_data): Rewrite functions. (is_oacc_parallel, oacc_parallel_max_reduction_array_size) (align_and_expand, alloc_var_ganglocal, install_var_ganglocal) (install_array_var_ganglocal) (oacc_outermost_parallel_kernels_context, oacc_inside_routine) (is_oacc_multithreaded, oacc_needs_global_memory) (is_atomic_compatible_reduction, oacc_serial_reduction) (oacc_process_reduction_data_helper): New functions. (build_outer_var_ref, fixup_remapped_decl, scan_sharing_clauses) (check_omp_nesting_restrictions, lower_rec_input_clauses) (lower_reduction_clauses, oacc_initialize_reduction_data) (oacc_process_reduction_data, lower_omp_target) (lower_omp_regimplify_p): Update for OpenACC. * tree-parloops.c (create_parallel_loop): For OpenACC, switch from vector to gang parallelism. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_FORCE_TO_GANGLOCAL. include/ * gomp-constants.h (GOMP_MAP_FLAG_GANGLOCAL): New macro. (gomp_map_kind): Add GOMP_MAP_GANGLOCAL, GOMP_MAP_FORCE_TO_GANGLOCAL. libgomp/ * libgomp.h (splay_tree_key_s): Add dealloc_host member. Adjust all users. * libgomp.map (GOACC_2.0.GOMP_4_BRANCH): Add GOACC_deviceptr, GOACC_get_ganglocal_ptr. * libgomp_g.h (GOACC_get_ganglocal_ptr): New prototype. * oacc-mem.c (GOACC_deviceptr): New function. * oacc-parallel.c (__goacc_host_ganglocal_ptr): New static variable. (GOACC_get_ganglocal_ptr, alloc_host_shared_mem) (free_host_shared_mem, alloc_ganglocal_addrs): New functions. (GOACC_parallel, GOACC_kernels): Use them. Add shared_size formal parameter. Adjust all users. (GOACC_parallel): Remove num_workers check. (GOACC_enter_exit_data, GOACC_update): Handle more mapping kinds. (GOACC_get_num_threads, GOACC_get_thread_num): Add gang, worker, vector formal parameters. Adjust all users. * plugin/plugin-host.c (GOMP_OFFLOAD_openacc_parallel): Add shared_size formal parameter. Adjust all users. * plugin/plugin-nvptx.c (nvptx_exec) (GOMP_OFFLOAD_openacc_parallel): Add shared_size formal parameter. Adjust all users. * target.c (gomp_map_vars, gomp_unmap_vars) (gomp_offload_image_to_device): Update for OpenACC. * testsuite/libgomp.oacc-c++/template-reduction.C: Update. * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise. * testsuite/libgomp.oacc-fortran/collapse-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-7.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-8.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/enter-data.c: New file. * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-1.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c: Likewise. * testsuite/libgomp.oacc-c/gwv.c: Likewise. * testsuite/libgomp.oacc-fortran/firstprivate-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/host_data-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/implicit-firstprivate-ref.f90: Likewise. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c: Remove file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@223178 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 81 ++ gcc/ada/ChangeLog.gomp | 12 + gcc/ada/gcc-interface/utils.c | 17 +- gcc/builtin-types.def | 5 +- gcc/builtins.c | 90 ++ gcc/c-family/ChangeLog.gomp | 12 + gcc/c-family/c-common.c | 17 +- gcc/c/ChangeLog.gomp | 16 + gcc/c/c-parser.c | 7 +- gcc/c/c-typeck.c | 4 + gcc/config/nvptx/mkoffload.c | 8 + gcc/config/nvptx/nvptx.c | 1 + gcc/config/nvptx/nvptx.md | 52 + gcc/cp/ChangeLog.gomp | 15 + gcc/cp/parser.c | 10 + gcc/cp/semantics.c | 4 + gcc/doc/md.texi | 18 + gcc/fortran/ChangeLog.gomp | 23 + gcc/fortran/f95-lang.c | 14 +- gcc/fortran/gfortran.h | 4 +- gcc/fortran/openmp.c | 20 +- gcc/fortran/trans-openmp.c | 6 + gcc/fortran/types.def | 5 +- gcc/gimple.h | 23 + gcc/gimplify.c | 347 +++++- gcc/jit/ChangeLog.gomp | 19 + gcc/jit/jit-builtins.c | 10 +- gcc/jit/jit-builtins.h | 7 +- gcc/lto/ChangeLog.gomp | 12 + gcc/lto/lto-lang.c | 17 +- gcc/omp-builtins.def | 22 +- gcc/omp-low.c | 1145 ++++++++++++++++---- gcc/testsuite/ChangeLog.gomp | 36 + gcc/testsuite/c-c++-common/goacc/dtype-1.c | 93 +- gcc/testsuite/c-c++-common/goacc/dtype-2.c | 2 +- gcc/testsuite/c-c++-common/goacc/executeables-1.c | 74 ++ gcc/testsuite/c-c++-common/goacc/firstprivate.c | 9 + gcc/testsuite/c-c++-common/goacc/host_data-1.c | 1 - gcc/testsuite/c-c++-common/goacc/host_data-2.c | 1 - gcc/testsuite/c-c++-common/goacc/host_data-3.c | 6 +- gcc/testsuite/c-c++-common/goacc/host_data-4.c | 1 - .../c-c++-common/goacc/private-reduction-1.c | 10 + gcc/testsuite/c-c++-common/goacc/sb-3.c | 2 +- gcc/testsuite/c-c++-common/goacc/tile.c | 3 - gcc/testsuite/g++.dg/goacc/loop-1.c | 23 + gcc/testsuite/g++.dg/goacc/loop-2.c | 70 ++ gcc/testsuite/g++.dg/goacc/loop-3.c | 43 + gcc/testsuite/g++.dg/goacc/template-reduction.C | 6 +- gcc/testsuite/g++.dg/goacc/template.C | 15 + gcc/testsuite/gcc.dg/goacc/sb-1.c | 73 ++ gcc/testsuite/gcc.dg/goacc/sb-2.c | 20 + .../{c-c++-common => gcc.dg}/goacc/sb-3.c | 4 +- gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 105 +- gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 | 11 + gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/list.f95 | 6 +- gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +- gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 | 2 +- gcc/tree-parloops.c | 18 +- gcc/tree-pretty-print.c | 12 + include/ChangeLog.gomp | 13 + include/gomp-constants.h | 7 +- libgomp/ChangeLog.gomp | 71 ++ libgomp/libgomp.h | 4 +- libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 9 +- libgomp/oacc-mem.c | 38 +- libgomp/oacc-parallel.c | 140 ++- libgomp/plugin/plugin-host.c | 1 + libgomp/plugin/plugin-nvptx.c | 35 +- libgomp/target.c | 11 + .../libgomp.oacc-c++/template-reduction.C | 14 +- .../libgomp.oacc-c-c++-common/collapse-2.c | 2 +- .../testsuite/libgomp.oacc-c-c++-common/data-2.c | 161 ++- .../libgomp.oacc-c-c++-common/enter-data.c | 23 + .../libgomp.oacc-c-c++-common/firstprivate-1.c | 32 + .../libgomp.oacc-c-c++-common/firstprivate-2.c | 55 + .../libgomp.oacc-c-c++-common/host_data-1.c | 125 +++ .../libgomp.oacc-c-c++-common/host_data-2.c | 50 + .../libgomp.oacc-c-c++-common/parallel-loop-1.c | 37 + .../libgomp.oacc-c-c++-common/parallel-loop-1.h | 20 + .../libgomp.oacc-c-c++-common/parallel-loop-2.h | 282 +++++ .../libgomp.oacc-c-c++-common/parallel-reduction.c | 67 ++ .../libgomp.oacc-c-c++-common/reduction-1.c | 13 +- .../libgomp.oacc-c-c++-common/reduction-2.c | 93 +- .../libgomp.oacc-c-c++-common/reduction-3.c | 93 +- .../libgomp.oacc-c-c++-common/reduction-4.c | 69 +- .../libgomp.oacc-c-c++-common/reduction-5.c | 6 +- .../reduction-initial-1.c | 25 - libgomp/testsuite/libgomp.oacc-c/gwv.c | 34 + .../testsuite/libgomp.oacc-fortran/collapse-5.f90 | 2 +- .../testsuite/libgomp.oacc-fortran/collapse-6.f90 | 2 +- .../testsuite/libgomp.oacc-fortran/collapse-7.f90 | 2 +- .../testsuite/libgomp.oacc-fortran/collapse-8.f90 | 2 +- libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 | 126 ++- .../libgomp.oacc-fortran/firstprivate-1.f90 | 42 + .../testsuite/libgomp.oacc-fortran/host_data-1.f90 | 28 + libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 | 453 ++++++++ .../implicit-firstprivate-ref.f90 | 42 + .../libgomp.oacc-fortran/parallel-reduction.f90 | 38 + .../testsuite/libgomp.oacc-fortran/reduction-1.f90 | 54 +- .../testsuite/libgomp.oacc-fortran/reduction-2.f90 | 46 +- .../testsuite/libgomp.oacc-fortran/reduction-3.f90 | 46 +- .../testsuite/libgomp.oacc-fortran/reduction-4.f90 | 36 +- .../testsuite/libgomp.oacc-fortran/reduction-5.f90 | 4 +- .../testsuite/libgomp.oacc-fortran/reduction-6.f90 | 4 +- 107 files changed, 4430 insertions(+), 722 deletions(-) Grüße, Thomas
On Wed, May 13, 2015 at 10:52:08PM +0200, Thomas Schwinge wrote: > In a similar vein, I have now committed the following to gomp-4_0-branch > in r223178. This is not meant to be integrated into trunk as-is: there > are incompatible libgomp ABI changes, for example. We'd still appreciate ABI incompatible libgomp changes are absolute no-no. You need to provide backwards compatibility... Jakub
--- a/gcc/testsuite/c-c++-common/goacc/host_data-4.c +++ b/gcc/testsuite/c-c++-common/goacc/host_data-4.c @@ -6,7 +6,7 @@ int main (int argc, char* argv[]) #pragma acc enter data copyin (x) /* Specifying an array index is not valid for host_data/use_device. */ - #pragma acc host_data use_device (x[4]) /* { dg-error "expected \\\')' before '\\\[' token" } */ + #pragma acc host_data use_device (x[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ ; #pragma acc exit data delete (x)