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 <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
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

Attachment: pgpvli97O2hHB.pgp
Description: PGP signature

Reply via email to