Hi! 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. commit cd00a35cd24a3ac05cac0061639ce631e52f2f49 Author: tschwinge Date: Mon May 11 16:29:03 2015 +0000 Next set of OpenACC changes gcc/c-family/ * c-common.c (c_common_attribute_table): Set min_len to -1 for "omp declare target". Add "oacc declare". * c-common.h (oacc_extract_device_id, oacc_filter_device_types): New prototypes. * c-omp.c (oacc_extract_device_id, oacc_filter_device_types): New functions. * c-pragma.c (oacc_pragmas): Add "atomic", "declare", "host_data", "routine". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ATOMIC, PRAGMA_OACC_DECLARE, PRAGMA_OACC_HOST_DATA, PRAGMA_OACC_ROUTINE. (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_LINK, PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_DEFAULT, remove PRAGMA_OACC_CLAUSE_SELF. gcc/c/ * c-parser.c (c_parser): Add oacc_routines member. (c_parse_file): Initialize it. (c_parser_declaration_or_fndef): Add oacc_routine_clauses, and oacc_routine_named formal parameters. Adjust all users. Support OpenACC routines. (c_parser_pragma): Handle PRAGMA_OACC_DECLARE, PRAGMA_OACC_ROUTINE, PRAGMA_OACC_WAIT. Add pragma context checking for PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA. (c_parser_omp_clause_name): Add consume_token formal parameter. Handle "bind", "device_resident", "device_type", "dtype", "independent", "link", "nohost", "tile", "use_device". (c_parser_oacc_wait_list): Change an error message. (c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_LINK. Don't handle PRAGMA_OACC_CLAUSE_SELF. (c_parser_omp_clause_default): Add only_none formal parameter. (c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_threads) (c_parser_omp_clause_num_workers) (c_parser_omp_clause_vector_length): Replace functions by... (require_positive_expr, c_parser_omp_positive_int_clause): ... these new functions. Adjust all users. (c_parser_omp_clause_untied, c_parser_omp_clause_branch): Replace functions by... (c_parser_omp_simple_clause): ... this new function. Adjust all users. (c_parser_oacc_shape_clause, c_parser_oacc_clause_bind) (c_parser_oacc_clause_device_type, c_parser_oacc_clause_tile) (c_parser_oacc_clause_use_device): New functions. (c_parser_oacc_all_clauses): Add dtype_mask, and scan_dtype formal parameters. Adjust all users. Handle PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_LINK, PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_WORKER. Don't handle PRAGMA_OACC_CLAUSE_SELF. (c_parser_oacc_declare, oacc_split_loop_clauses) (c_parser_oacc_host_data, c_parser_oacc_routine) (c_finish_oacc_routine): New functions. (c_parser_oacc_enter_exit_data): Change error reporting. (c_parser_oacc_loop): Add mask, and cclauses formal parameters. Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_WORKER, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_PRIVATE. (c_parser_oacc_kernels): Support combined directives. Handle PRAGMA_OACC_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE. (c_parser_oacc_parallel): Likewise. Handle PRAGMA_OACC_CLAUSE_PRIVATE. (c_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC, PRAGMA_OACC_HOST_DATA. (c_parser_oacc_update): Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_WAIT. Don't handle PRAGMA_OACC_CLAUSE_SELF. * c-tree.h (c_finish_oacc_host_data): New prototype. * c-typeck.c (c_finish_oacc_host_data): New function. (c_finish_omp_clauses): Add oacc formal parameter. Adjust all users. Handle OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE. gcc/cp/ * cp-gimplify.c (cxx_omp_clause_copy_ctor): Handle OMP_CLAUSE_MAP. * cp-tree.h (finish_oacc_host_data): New prototype. * parser.h (cp_parser): Add oacc_routine, named_oacc_routines members. * parser.c (cp_parser_new): Initialize them. (cp_ensure_no_omp_declare_simd, cp_parser_init_declarator) (cp_parser_late_return_type_opt, cp_parser_member_declaration) (cp_parser_function_definition_from_specifiers_and_declarator) (cp_parser_save_member_function_body, cp_parser_omp_declare_simd) (cp_parser_omp_declare, cp_parser_pragma): Extend for OpenACC routines. (cp_finalize_oacc_routine) (cp_parser_oacc_routine_check_parallelism, cp_parser_oacc_routine) (cp_parser_late_parsing_oacc_routine): New functions. (cp_parser_omp_clause_name): Add consume_token formal parameter. Handle "auto", "bind", "device_resident", "device_type", "dtype", "gang", "independent", "link", "nohost", "seq", "tile", "use_device", "vector", "worker". (cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_LINK. Don't handle PRAGMA_OACC_CLAUSE_SELF. (cp_parser_oacc_clause_vector_length) (cp_parser_omp_clause_num_gangs, cp_parser_omp_clause_num_threads) (cp_parser_omp_clause_num_workers): Replace functions by... (require_positive_expr, cp_parser_omp_positive_int_clause): ... these new functions. Adjust all users. (cp_parser_oacc_shape_clause, cp_parser_oacc_clause_device_type) (cp_parser_oacc_clause_tile, cp_parser_oacc_clause_bind): New functions. (cp_parser_oacc_wait_list): Change an error message. (cp_parser_omp_clause_default): Add is_omp formal parameter. (cp_parser_omp_clause_untied, cp_parser_omp_clause_branch): Replace functions by... (cp_parser_omp_simple_clause): ... this new function. Adjust all users. (cp_parser_oacc_all_clauses): Add dtype_mask, and scan_dtype formal parameters. Adjust all users. Handle PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_LINK, PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_PRIVATE, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_WORKER. Don't handle PRAGMA_OACC_CLAUSE_SELF. (cp_parser_oacc_host_data, oacc_split_loop_clauses): (cp_parser_oacc_enter_exit_data): Change error reporting. (cp_parser_oacc_loop): Add p_name, mask, and cclauses formal parameters. Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_PRIVATE PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_WORKER, PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE. Support combined directives. (cp_parser_oacc_kernels, cp_parser_oacc_parallel): Replace functions by... (cp_parser_oacc_parallel_kernels): ... this new function. Adjust all users. Support combined directives. For "kernels", handle PRAGMA_OACC_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE. For "parallel", handle PRAGMA_OACC_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_PRIVATE. (cp_parser_oacc_update): Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE. Don't handle PRAGMA_OACC_CLAUSE_SELF. (cp_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC, PRAGMA_OACC_HOST_DATA. (cp_parser_pragma): Handle PRAGMA_OACC_ATOMIC, PRAGMA_OACC_HOST_DATA, PRAGMA_OACC_ROUTINE. Add pragma context checking for PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_TILE. (tsubst_expr): Handle OACC_PARALLEL, OACC_KERNELS, OACC_LOOP, OACC_DATA, OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE. * semantics.c (finish_omp_clauses): Add oacc formal parameter. Adjust all users. Handle OMP_CLAUSE_GANG, OMP_CLAUSE_VECTOR, OMP_CLAUSE_WORKER, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_AUTO, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_SEQ, OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE. (finish_oacc_host_data): New function. gcc/fortran/ * dump-parse-tree.c (show_namespace): Rewrite handling of OpenACC declare. * gfortran.h (gfc_statement): Add ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (gfc_omp_clauses): Add routine_bind, dtype, dtype_clauses, nohost, acc_collapse, bind, num_gangs, num_workers, vector_length, tile members. (gfc_oacc_declare, gfc_oacc_routine_name): New typedefs. (gfc_get_oacc_declare, gfc_get_oacc_routine_name): New macros. (gfc_namespace): Add oacc_declare, oacc_routine_clauses, oacc_routine_names, oacc_routine members, remove oacc_declare_clauses member. (gfc_exec_op): Add EXEC_OACC_ROUTINE, EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE. (gfc_code): Add oacc_declare member. (gfc_free_oacc_declares, insert_oacc_declare): New prototypes. * match.h (gfc_match_oacc_atomic): New prototype. * openmp.c (OMP_CLAUSE_HOST_SELF): Rename to... (OMP_CLAUSE_HOST): ... this. Adjust all users. (OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_DEVICE_TYPE): New macros. (gfc_free_oacc_declares): New function. (gfc_match_omp_map_clause): Add allow_sections formal parameter. Adjust all users. (gfc_match_omp_clauses): Add dtype_mask formal parameter. Adjust all users. Change handling of OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_TILE, OMP_CLAUSE_DEFAULT, OMP_CLAUSE_COLLAPSE. Handle OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_DEVICE_TYPE. (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_LOOP_CLAUSES) (OACC_UPDATE_CLAUSES): Add OMP_CLAUSE_DEVICE_TYPE. (OACC_ROUTINE_CLAUSES, OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK) (OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK) (OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK) (OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK) (OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK): New macros. (gfc_match_oacc_declare, gfc_match_oacc_routine): Rewrite functions. (gfc_match_oacc_update): Add error reporting. (gfc_match_omp_atomic, gfc_match_oacc_atomic): New wrapper functions around... (gfc_match_omp_oacc_atomic): ... this new function. (check_array_not_assumed): Remove pointer check. (oacc_code_to_statement): Handle EXEC_OACC_ATOMIC. (resolve_oacc_loop_blocks): Don't error out for combined OpenACC gang, worker, and vector clauses. (resolve_oacc_cache): Remove function. (gfc_resolve_oacc_declare): Rewrite function. (gfc_resolve_oacc_directive): Handle EXEC_OACC_ATOMIC. Don't handle EXEC_OACC_CACHE. * parse.c (decode_oacc_directive): Handle "atomic", "end atomic". (case_exec_markers): Add ST_OACC_ATOMIC. (case_decl): Add ST_OACC_DECLARE. (gfc_ascii_statement): Handle ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (verify_st_order, parse_spec): Remove handling of ST_OACC_DECLARE. (parse_omp_atomic): Rename to... (parse_omp_oacc_atomic): ... this new function. Add omp_p formal parameter. Adjust all users. (parse_executable): Handle ST_OACC_ATOMIC. (parse_progunit): Remove handling of OpenACC declare. (is_oacc): Handle EXEC_OACC_ROUTINE. * parse.h (gfc_state_data): Add ext.oacc_declare member. Remove ext.oacc_declare_clauses member. * resolve.c (gfc_resolve_blocks): Handle EXEC_OACC_ATOMIC, EXEC_OACC_ROUTINE, EXEC_OACC_DECLARE. (gfc_resolve_code): Handle EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE. * st.c (gfc_free_statement): Handle EXEC_OACC_DECLARE, EXEC_OACC_ROUTINE, EXEC_OACC_ATOMIC. * trans-decl.c (find_end, insert_oacc_declare): New functions. (gfc_generate_function_code): Change handling of OpenACC declare. * trans-openmp.c (gfc_omp_clause_copy_ctor): Handle OMP_CLAUSE_REDUCTION. (gfc_trans_omp_clauses): Add appropriate, generate OMP_CLAUSE_SEQ (instead of OMP_CLAUSE_ORDERED), OMP_CLAUSE_AUTO, or OMP_CLAUSE_TILE. (gfc_trans_oacc_combined_directive): Don't set OACC_KERNELS_COMBINED, and OACC_PARALLEL_COMBINED. (gfc_trans_oacc_declare): Rewrite function. (gfc_trans_oacc_directive): Handle EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE. * trans-stmt.c (gfc_trans_block_construct): Change handling of OpenACC declare. * trans.c (trans_code): Handle EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE. gcc/ * gimplify.c (gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_TILE. (gimplify_expr): Don't verify OACC_KERNELS_COMBINED, and OACC_PARALLEL_COMBINED. * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE. (check_omp_nesting_restrictions): Support GIMPLE_OMP_ATOMIC_LOAD, GIMPLE_OMP_ATOMIC_STORE inside OpenACC contexts. * tree-core.h (omp_clause_code): Add OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE, OMP_CLAUSE_DEVICE_TYPE. * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1): Update for these. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE. * tree.h (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): Remove macros. (OMP_CLAUSE_BIND_NAME, OMP_CLAUSE_TILE_LIST) (OMP_CLAUSE_DEVICE_TYPE_DEVICES, OMP_CLAUSE_DEVICE_TYPE_CLAUSES): Add macros. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-1.c: Update. * c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise. * c-c++-common/goacc/asyncwait-1.c: Likewise. * c-c++-common/goacc/data-2.c: Likewise. * c-c++-common/goacc/reduction-1.c: Likewise. * c-c++-common/goacc/reduction-2.c: Likewise. * c-c++-common/goacc/reduction-3.c: Likewise. * c-c++-common/goacc/reduction-4.c: Likewise. * gfortran.dg/goacc/cache-1.f95: Likewise. * gfortran.dg/goacc/coarray.f95: Likewise. * gfortran.dg/goacc/coarray_2.f90: Likewise. * gfortran.dg/goacc/combined_loop.f90: Likewise. * gfortran.dg/goacc/cray.f95: Likewise. * gfortran.dg/goacc/declare-1.f95: Likewise. * gfortran.dg/goacc/host_data-tree.f95: Likewise. * gfortran.dg/goacc/loop-1.f95: Likewise. * gfortran.dg/goacc/loop-2.f95: Likewise. * gfortran.dg/goacc/parameter.f95: Likewise. * c-c++-common/goacc/loop-1.c: Enable for C++. * c-c++-common/goacc/kernels-1.c: Rename to... * c-c++-common/goacc/kernels-empty.c: ... this new file. * c-c++-common/goacc/parallel-1.c: Rename to... * c-c++-common/goacc/parallel-empty.c: ... this new file. * c-c++-common/goacc/declare-1.c: New file. * c-c++-common/goacc/declare-2.c: Likewise. * c-c++-common/goacc/dtype-1.c: Likewise. * 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/kernels-eternal.c: Likewise. * c-c++-common/goacc/kernels-noreturn.c: Likewise. * c-c++-common/goacc/parallel-eternal.c: Likewise. * c-c++-common/goacc/parallel-noreturn.c: Likewise. * c-c++-common/goacc/routine-1.c: Likewise. * c-c++-common/goacc/routine-2.c: Likewise. * c-c++-common/goacc/routine-3.c: Likewise. * c-c++-common/goacc/routine-4.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/declare-2.f95: Likewise. * gfortran.dg/goacc/default.f95: Likewise. * gfortran.dg/goacc/dtype-1.f95: Likewise. * gfortran.dg/goacc/dtype-2.f95: Likewise. * gfortran.dg/goacc/modules.f95: Likewise. * gfortran.dg/goacc/update.f95: Likewise. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT, GOMP_MAP_LINK. libgomp/ * oacc-mem.c (update_dev_host): Add missing initialization. * oacc-ptx.h (GOMP_ATOMIC_PTX): New macro. * plugin/plugin-nvptx.c (link_ptx): Link it in. * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-69.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-71.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-72.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-74.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-77.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-83.c: Likewise. * testsuite/libgomp.oacc-fortran/data-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-4.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/subr.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/subr.ptx: Likewise. * testsuite/libgomp.oacc-c-c++-common/timer.h: Remove file. * testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Move common code from here... * testsuite/libgomp.oacc-c-c++-common/parallel-1.c: ..., and here... * testsuite/libgomp.oacc-c-c++-common/data-clauses.h: ... into this new file. * testsuite/libgomp.oacc-c++/template-reduction.C: New test. * testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-2.c: Likewise. * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/atomic_update-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/cache-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/clauses-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-13.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-14.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-15.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-5.f90: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@223007 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 28 + gcc/c-family/ChangeLog.gomp | 26 + gcc/c-family/c-common.c | 3 +- gcc/c-family/c-common.h | 2 + gcc/c-family/c-omp.c | 105 ++ gcc/c-family/c-pragma.c | 4 + gcc/c-family/c-pragma.h | 14 +- gcc/c/ChangeLog.gomp | 71 + gcc/c/c-parser.c | 1353 ++++++++++++---- gcc/c/c-tree.h | 3 +- gcc/c/c-typeck.c | 112 +- gcc/cp/ChangeLog.gomp | 93 ++ gcc/cp/cp-gimplify.c | 3 +- gcc/cp/cp-tree.h | 3 +- gcc/cp/parser.c | 1382 +++++++++++++---- gcc/cp/parser.h | 4 + gcc/cp/pt.c | 43 +- gcc/cp/semantics.c | 151 +- gcc/fortran/ChangeLog.gomp | 94 ++ gcc/fortran/dump-parse-tree.c | 12 +- gcc/fortran/gfortran.h | 50 +- gcc/fortran/match.h | 1 + gcc/fortran/openmp.c | 581 +++++-- gcc/fortran/parse.c | 65 +- gcc/fortran/parse.h | 2 +- gcc/fortran/resolve.c | 5 + gcc/fortran/st.c | 7 + gcc/fortran/trans-decl.c | 62 +- gcc/fortran/trans-openmp.c | 66 +- gcc/fortran/trans-stmt.c | 7 +- gcc/fortran/trans-stmt.h | 2 +- gcc/fortran/trans.c | 2 + gcc/gimplify.c | 16 +- gcc/omp-low.c | 11 +- gcc/testsuite/ChangeLog.gomp | 58 + gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c | 46 + .../c-c++-common/goacc-gomp/nesting-fail-1.c | 25 - gcc/testsuite/c-c++-common/goacc/asyncwait-1.c | 4 +- gcc/testsuite/c-c++-common/goacc/data-2.c | 12 +- gcc/testsuite/c-c++-common/goacc/declare-1.c | 84 + gcc/testsuite/c-c++-common/goacc/declare-2.c | 67 + gcc/testsuite/c-c++-common/goacc/dtype-1.c | 113 ++ gcc/testsuite/c-c++-common/goacc/dtype-2.c | 31 + gcc/testsuite/c-c++-common/goacc/host_data-1.c | 14 + gcc/testsuite/c-c++-common/goacc/host_data-2.c | 14 + gcc/testsuite/c-c++-common/goacc/host_data-3.c | 16 + gcc/testsuite/c-c++-common/goacc/host_data-4.c | 15 + .../goacc/{kernels-1.c => kernels-empty.c} | 0 gcc/testsuite/c-c++-common/goacc/kernels-eternal.c | 11 + .../c-c++-common/goacc/kernels-noreturn.c | 12 + gcc/testsuite/c-c++-common/goacc/loop-1.c | 2 - .../goacc/{parallel-1.c => parallel-empty.c} | 0 .../c-c++-common/goacc/parallel-eternal.c | 11 + .../c-c++-common/goacc/parallel-noreturn.c | 12 + gcc/testsuite/c-c++-common/goacc/reduction-1.c | 25 +- gcc/testsuite/c-c++-common/goacc/reduction-2.c | 22 +- gcc/testsuite/c-c++-common/goacc/reduction-3.c | 22 +- gcc/testsuite/c-c++-common/goacc/reduction-4.c | 40 +- gcc/testsuite/c-c++-common/goacc/routine-1.c | 35 + gcc/testsuite/c-c++-common/goacc/routine-2.c | 36 + gcc/testsuite/c-c++-common/goacc/routine-3.c | 52 + gcc/testsuite/c-c++-common/goacc/routine-4.c | 87 ++ gcc/testsuite/c-c++-common/goacc/tile.c | 26 + gcc/testsuite/g++.dg/goacc/template-reduction.C | 100 ++ gcc/testsuite/g++.dg/goacc/template.C | 131 ++ gcc/testsuite/gfortran.dg/goacc/cache-1.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 2 +- gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 | 1 + gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 | 2 +- gcc/testsuite/gfortran.dg/goacc/cray.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/declare-1.f95 | 3 +- gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 44 + gcc/testsuite/gfortran.dg/goacc/default.f95 | 17 + gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 161 ++ gcc/testsuite/gfortran.dg/goacc/dtype-2.f95 | 39 + gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 2 +- gcc/testsuite/gfortran.dg/goacc/loop-1.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/loop-2.f95 | 26 +- gcc/testsuite/gfortran.dg/goacc/modules.f95 | 55 + gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/update.f95 | 5 + gcc/tree-core.h | 14 +- gcc/tree-pretty-print.c | 6 + gcc/tree.c | 13 +- gcc/tree.h | 21 +- include/ChangeLog.gomp | 6 + include/gomp-constants.h | 4 + libgomp/ChangeLog.gomp | 70 + libgomp/oacc-mem.c | 3 + libgomp/oacc-ptx.h | 28 + libgomp/plugin/plugin-nvptx.c | 10 + .../libgomp.oacc-c++/template-reduction.C | 102 ++ .../libgomp.oacc-c-c++-common/atomic_capture-1.c | 866 +++++++++++ .../libgomp.oacc-c-c++-common/atomic_capture-2.c | 1626 ++++++++++++++++++++ .../libgomp.oacc-c-c++-common/atomic_update-1.c | 760 +++++++++ .../libgomp.oacc-c-c++-common/clauses-1.c | 26 + .../testsuite/libgomp.oacc-c-c++-common/data-2.c | 44 +- .../testsuite/libgomp.oacc-c-c++-common/data-3.c | 18 +- .../{parallel-1.c => data-clauses.h} | 32 +- .../libgomp.oacc-c-c++-common/kernels-1.c | 182 +-- .../testsuite/libgomp.oacc-c-c++-common/lib-69.c | 70 +- .../testsuite/libgomp.oacc-c-c++-common/lib-70.c | 79 +- .../testsuite/libgomp.oacc-c-c++-common/lib-71.c | 55 +- .../testsuite/libgomp.oacc-c-c++-common/lib-72.c | 60 +- .../testsuite/libgomp.oacc-c-c++-common/lib-73.c | 64 +- .../testsuite/libgomp.oacc-c-c++-common/lib-74.c | 91 +- .../testsuite/libgomp.oacc-c-c++-common/lib-75.c | 89 +- .../testsuite/libgomp.oacc-c-c++-common/lib-76.c | 88 +- .../testsuite/libgomp.oacc-c-c++-common/lib-77.c | 91 +- .../testsuite/libgomp.oacc-c-c++-common/lib-78.c | 91 +- .../testsuite/libgomp.oacc-c-c++-common/lib-79.c | 91 +- .../testsuite/libgomp.oacc-c-c++-common/lib-80.c | 95 +- .../testsuite/libgomp.oacc-c-c++-common/lib-81.c | 106 +- .../testsuite/libgomp.oacc-c-c++-common/lib-82.c | 43 +- .../testsuite/libgomp.oacc-c-c++-common/lib-83.c | 22 +- .../libgomp.oacc-c-c++-common/parallel-1.c | 204 +-- .../libgomp.oacc-c-c++-common/routine-1.c | 40 + .../libgomp.oacc-c-c++-common/routine-2.c | 41 + libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h | 44 +- .../testsuite/libgomp.oacc-c-c++-common/subr.ptx | 222 +-- .../testsuite/libgomp.oacc-c-c++-common/timer.h | 103 -- .../libgomp.oacc-fortran/atomic_capture-1.f90 | 784 ++++++++++ .../libgomp.oacc-fortran/atomic_update-1.f90 | 338 ++++ libgomp/testsuite/libgomp.oacc-fortran/cache-1.f90 | 26 + .../testsuite/libgomp.oacc-fortran/clauses-1.f90 | 290 ++++ libgomp/testsuite/libgomp.oacc-fortran/data-1.f90 | 231 ++- libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 | 50 + libgomp/testsuite/libgomp.oacc-fortran/data-3.f90 | 34 +- .../testsuite/libgomp.oacc-fortran/data-4-2.f90 | 19 +- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 | 19 +- .../testsuite/libgomp.oacc-fortran/declare-1.f90 | 229 +++ libgomp/testsuite/libgomp.oacc-fortran/lib-12.f90 | 24 + libgomp/testsuite/libgomp.oacc-fortran/lib-13.f90 | 28 + libgomp/testsuite/libgomp.oacc-fortran/lib-14.f90 | 79 + libgomp/testsuite/libgomp.oacc-fortran/lib-15.f90 | 52 + .../testsuite/libgomp.oacc-fortran/routine-5.f90 | 27 + 136 files changed, 11216 insertions(+), 2501 deletions(-) This is the patches as posted, including the following last-minute typo fix: --- 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) Grüße, Thomas