Merge current set of OpenACC changes from gomp-4_0-branch.

	contrib/
	* gcc_update (files_and_dependencies): Update rules for new
	libgomp/plugin/Makefrag.am and libgomp/plugin/configfrag.ac files.
	gcc/
	* builtin-types.def (BT_FN_VOID_INT_INT_VAR)
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
	New function types.
	* builtins.c: Include "gomp-constants.h".
	(expand_builtin_acc_on_device): New function.
	(expand_builtin, is_inexpensive_builtin): Handle
	BUILT_IN_ACC_ON_DEVICE.
	* builtins.def (DEF_GOACC_BUILTIN, DEF_GOACC_BUILTIN_COMPILER):
	New macros.
	* cgraph.c (cgraph_node::create): Consider flag_openacc next to
	flag_openmp.
	* config.gcc <nvptx-*> (tm_file): Add nvptx/offload.h.
	<*-intelmic-* | *-intelmicemul-*> (tm_file): Add
	i386/intelmic-offload.h.
	* gcc.c (LINK_COMMAND_SPEC, GOMP_SELF_SPECS): For -fopenacc, link
	to libgomp and its dependencies.
	* config/arc/arc.h (LINK_COMMAND_SPEC): Likewise.
	* config/darwin.h (LINK_COMMAND_SPEC_A): Likewise.
	* config/i386/mingw32.h (GOMP_SELF_SPECS): Likewise.
	* config/ia64/hpux.h (LIB_SPEC): Likewise.
	* config/pa/pa-hpux11.h (LIB_SPEC): Likewise.
	* config/pa/pa64-hpux.h (LIB_SPEC): Likewise.
	* doc/generic.texi: Update for OpenACC changes.
	* doc/gimple.texi: Likewise.
	* doc/invoke.texi: Likewise.
	* doc/sourcebuild.texi: Likewise.
	* gimple-pretty-print.c (dump_gimple_omp_for): Handle
	GF_OMP_FOR_KIND_OACC_LOOP.
	(dump_gimple_omp_target): Handle GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_DATA,
	GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
	Dump more data.
	* gimple.c: Update comments for OpenACC changes.
	* gimple.def: Likewise.
	* gimple.h: Likewise.
	(enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP,
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_DATA, GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
	(gimple_omp_for_cond, gimple_omp_for_set_cond): Sort in the
	appropriate place.
	(is_gimple_omp_oacc, is_gimple_omp_offloaded): New functions.
	* gimplify.c: Include "gomp-constants.h".
	Update comments for OpenACC changes.
	(is_gimple_stmt): Handle OACC_PARALLEL, OACC_KERNELS, OACC_DATA,
	OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE, OACC_ENTER_DATA,
	OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
	(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
	OMP_CLAUSE__CACHE_, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER,
	OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
	OMP_CLAUSE_SEQ.
	(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(gimplify_oacc_cache): New function.
	(gimplify_omp_for): Handle OACC_LOOP.
	(gimplify_omp_workshare): Handle OACC_KERNELS, OACC_PARALLEL,
	OACC_DATA.
	(gimplify_omp_target_update): Handle OACC_ENTER_DATA,
	OACC_EXIT_DATA, OACC_UPDATE.
	(gimplify_expr): Handle OACC_LOOP, OACC_CACHE, OACC_HOST_DATA,
	OACC_DECLARE, OACC_KERNELS, OACC_PARALLEL, OACC_DATA,
	OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE.
	(gimplify_body): Consider flag_openacc next to flag_openmp.
	* lto-streamer-out.c: Include "gomp-constants.h".
	* omp-builtins.def (BUILT_IN_ACC_GET_DEVICE_TYPE)
	(BUILT_IN_GOACC_DATA_START, BUILT_IN_GOACC_DATA_END)
	(BUILT_IN_GOACC_ENTER_EXIT_DATA, BUILT_IN_GOACC_PARALLEL)
	(BUILT_IN_GOACC_UPDATE, BUILT_IN_GOACC_WAIT)
	(BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS)
	(BUILT_IN_ACC_ON_DEVICE): New builtins.
	* omp-low.c: Include "gomp-constants.h".
	Update comments for OpenACC changes.
	(struct omp_context): Add reduction_map, gwv_below, gwv_this
	members.
	(extract_omp_for_data, use_pointer_for_field, install_var_field)
	(new_omp_context, delete_omp_context, scan_sharing_clauses)
	(create_omp_child_function, scan_omp_for, scan_omp_target)
	(check_omp_nesting_restrictions, lower_reduction_clauses)
	(build_omp_regions_1, diagnose_sb_0, make_gimple_omp_edges):
	Update for OpenACC changes.
	(scan_sharing_clauses): Handle OMP_CLAUSE_NUM_GANGS:
	OMP_CLAUSE_NUM_WORKERS: OMP_CLAUSE_VECTOR_LENGTH,
	OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT, OMP_CLAUSE_GANG,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_INDEPENDENT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ.  Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.
	(expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
	Handle GF_OMP_FOR_KIND_OACC_LOOP.
	(expand_omp_target, lower_omp_target): Handle
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA,
	GF_OMP_TARGET_KIND_OACC_DATA.
	(pass_expand_omp::execute, execute_lower_omp)
	(pass_diagnose_omp_blocks::gate): Consider flag_openacc next to
	flag_openmp.
	(offload_symbol_decl): New variable.
	(oacc_get_reduction_array_id, oacc_max_threads)
	(get_offload_symbol_decl, get_base_type, lookup_oacc_reduction)
	(maybe_lookup_oacc_reduction, enclosing_target_ctx)
	(oacc_loop_or_target_p, oacc_lower_reduction_var_helper)
	(oacc_gimple_assign, oacc_initialize_reduction_data)
	(oacc_finalize_reduction_data, oacc_process_reduction_data): New
	functions.
	(is_targetreg_ctx): Remove function.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CACHE_,
	OMP_CLAUSE_DEVICE_RESIDENT, OMP_CLAUSE_USE_DEVICE,
	OMP_CLAUSE_GANG, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_NUM_GANGS,
	OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH.
	* tree.c (omp_clause_code_name, walk_tree_1): Update accordingly.
	* tree.h (OMP_CLAUSE_GANG_EXPR, OMP_CLAUSE_GANG_STATIC_EXPR)
	(OMP_CLAUSE_ASYNC_EXPR, OMP_CLAUSE_WAIT_EXPR)
	(OMP_CLAUSE_VECTOR_EXPR, OMP_CLAUSE_WORKER_EXPR)
	(OMP_CLAUSE_NUM_GANGS_EXPR, OMP_CLAUSE_NUM_WORKERS_EXPR)
	(OMP_CLAUSE_VECTOR_LENGTH_EXPR): New macros.
	* tree-core.h: Update comments for OpenACC changes.
	(enum omp_clause_map_kind): Remove.
	(struct tree_omp_clause): Change type of map_kind member from enum
	omp_clause_map_kind to unsigned char.
	* tree-inline.c: Update comments for OpenACC changes.
	* tree-nested.c: Likewise.  Include "gomp-constants.h".
	(convert_nonlocal_reference_stmt, convert_local_reference_stmt)
	(convert_tramp_reference_stmt, convert_gimple_call): Update for
	OpenACC changes.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	* tree-pretty-print.c: Include "gomp-constants.h".
	(dump_omp_clause): Handle OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_GANG,
	OMP_CLAUSE_ASYNC, OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ,
	OMP_CLAUSE_WAIT, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_INDEPENDENT.  Use GOMP_MAP_*
	instead of OMP_CLAUSE_MAP_*.
	(dump_generic_node): Handle OACC_PARALLEL, OACC_KERNELS,
	OACC_DATA, OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE,
	OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
	* tree-streamer-in.c: Include "gomp-constants.h".
	(unpack_ts_omp_clause_value_fields) Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.  Use OMP_CLAUSE_SET_MAP_KIND.
	* tree-streamer-out.c: Include "gomp-constants.h".
	(pack_ts_omp_clause_value_fields): Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.
	* tree.def (OACC_PARALLEL, OACC_KERNELS, OACC_DATA)
	(OACC_HOST_DATA, OACC_LOOP, OACC_CACHE, OACC_DECLARE)
	(OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE): New tree codes.
	* tree.c (omp_clause_num_ops): Update accordingly.
	* tree.h (OMP_BODY, OMP_CLAUSES, OMP_LOOP_CHECK, OMP_CLAUSE_SIZE):
	Likewise.
	(OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES, OACC_KERNELS_BODY)
	(OACC_KERNELS_CLAUSES, OACC_DATA_BODY, OACC_DATA_CLAUSES)
	(OACC_HOST_DATA_BODY, OACC_HOST_DATA_CLAUSES, OACC_CACHE_CLAUSES)
	(OACC_DECLARE_CLAUSES, OACC_ENTER_DATA_CLAUSES)
	(OACC_EXIT_DATA_CLAUSES, OACC_UPDATE_CLAUSES)
	(OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): New macros.
	* tree.h (OMP_CLAUSE_MAP_KIND): Cast it to enum gomp_map_kind.
	(OMP_CLAUSE_SET_MAP_KIND): New macro.
	* varpool.c (varpool_node::get_create): Consider flag_openacc next
	to flag_openmp.
	* config/i386/intelmic-offload.h: New file.
	* config/nvptx/offload.h: Likewise.
	gcc/ada/
	* gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_8)
	(DEF_FUNCTION_TYPE_VAR_12): New macros.
	gcc/c-family/
	* c.opt (fopenacc): New option.
	* c-cppbuiltin.c (c_cpp_builtins): Conditionally define _OPENACC.
	* c-common.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
	New macros.
	* c-common.h (c_finish_oacc_wait): New prototype.
	* c-omp.c: Include "omp-low.h" and "gomp-constants.h".
	(c_finish_oacc_wait): New function.
	* c-pragma.c (oacc_pragmas): New variable.
	(c_pp_lookup_pragma, init_pragma): Handle it.
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_CACHE,
	PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA,
	PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL,
	PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
	(enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ASYNC,
	PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_COLLAPSE,
	PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYIN,
	PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE,
	PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICE,
	PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE,
	PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST,
	PRAGMA_OACC_CLAUSE_IF, PRAGMA_OACC_CLAUSE_NUM_GANGS,
	PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_PRIVATE,
	PRAGMA_OACC_CLAUSE_REDUCTION, PRAGMA_OACC_CLAUSE_SELF,
	PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_VECTOR,
	PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT,
	PRAGMA_OACC_CLAUSE_WORKER.
	gcc/c/
	* c-parser.c: Include "gomp-constants.h".
	(c_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
	omp_clause_map_kind.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	Use OMP_CLAUSE_SET_MAP_KIND.
	(c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA,
	PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_UPDATE.
	(c_parser_omp_construct): Handle PRAGMA_OACC_CACHE,
	PRAGMA_OACC_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP,
	PRAGMA_OACC_PARALLEL, PRAGMA_OACC_WAIT.
	(c_parser_omp_clause_name): Handle "auto", "async", "copy",
	"copyout", "create", "delete", "deviceptr", "gang", "host",
	"num_gangs", "num_workers", "present", "present_or_copy", "pcopy",
	"present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
	"present_or_create", "pcreate", "seq", "self", "vector",
	"vector_length", "wait", "worker".
	(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
	(OACC_ENTER_DATA_CLAUSE_MASK, OACC_EXIT_DATA_CLAUSE_MASK)
	(OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
	(OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
	(c_parser_omp_variable_list): Handle OMP_CLAUSE__CACHE_.
	(c_parser_oacc_wait_list, c_parser_oacc_data_clause)
	(c_parser_oacc_data_clause_deviceptr)
	(c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_workers)
	(c_parser_oacc_clause_async, c_parser_oacc_clause_wait)
	(c_parser_omp_clause_vector_length, c_parser_oacc_all_clauses)
	(c_parser_oacc_cache, c_parser_oacc_data, c_parser_oacc_kernels)
	(c_parser_oacc_enter_exit_data, c_parser_oacc_loop)
	(c_parser_oacc_parallel, c_parser_oacc_update)
	(c_parser_oacc_wait): New functions.
	* c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels)
	(c_finish_oacc_data): New prototypes.
	* c-typeck.c: Include "gomp-constants.h".
	(handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR.  Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(c_finish_oacc_parallel, c_finish_oacc_kernels)
	(c_finish_oacc_data): New functions.
	(c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_GANG,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, and OMP_CLAUSE_MAP's
	GOMP_MAP_FORCE_DEVICEPTR.
	gcc/cp/
	* parser.c: Include "gomp-constants.h".
	(cp_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
	omp_clause_map_kind.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	Use OMP_CLAUSE_SET_MAP_KIND.
	(cp_parser_omp_construct, cp_parser_pragma): Handle
	PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA,
	PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_PARALLEL,
	PRAGMA_OACC_LOOP, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
	(cp_parser_omp_clause_name): Handle "async", "copy", "copyout",
	"create", "delete", "deviceptr", "host", "num_gangs",
	"num_workers", "present", "present_or_copy", "pcopy",
	"present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
	"present_or_create", "pcreate", "vector_length", "wait".
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
	(OACC_EXIT_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
	(OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
	(OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
	(cp_parser_omp_var_list_no_open): Handle OMP_CLAUSE__CACHE_.
	(cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr)
	(cp_parser_oacc_clause_vector_length, cp_parser_oacc_wait_list)
	(cp_parser_oacc_clause_wait, cp_parser_omp_clause_num_gangs)
	(cp_parser_omp_clause_num_workers, cp_parser_oacc_clause_async)
	(cp_parser_oacc_all_clauses, cp_parser_oacc_cache)
	(cp_parser_oacc_data, cp_parser_oacc_enter_exit_data)
	(cp_parser_oacc_kernels, cp_parser_oacc_loop)
	(cp_parser_oacc_parallel, cp_parser_oacc_update)
	(cp_parser_oacc_wait): New functions.
	* cp-tree.h (finish_oacc_data, finish_oacc_kernels)
	(finish_oacc_parallel): New prototypes.
	* semantics.c: Include "gomp-constants.h".
	(handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR.  Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT, OMP_CLAUSE__CACHE_.
	Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	(finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
	functions.
	gcc/fortran/
	* lang.opt (fopenacc): New option.
	* cpp.c (cpp_define_builtins): Conditionally define _OPENACC.
	* dump-parse-tree.c (show_omp_node): Split part of it into...
	(show_omp_clauses): ... this new function.
	(show_omp_node, show_code_node): Handle EXEC_OACC_PARALLEL_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	(show_namespace): Update for OpenACC.
	* f95-lang.c (DEF_FUNCTION_TYPE_VAR_2, DEF_FUNCTION_TYPE_VAR_8)
	(DEF_FUNCTION_TYPE_VAR_12, DEF_GOACC_BUILTIN)
	(DEF_GOACC_BUILTIN_COMPILER): New macros.
	* types.def (BT_FN_VOID_INT_INT_VAR)
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
	New function types.
	* gfortran.h (gfc_statement): Add ST_OACC_PARALLEL_LOOP,
	ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_DATA,
	ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA,
	ST_OACC_LOOP, ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE,
	ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP,
	ST_OACC_END_KERNELS_LOOP, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
	ST_OACC_ROUTINE.
	(struct gfc_expr_list): New data type.
	(gfc_get_expr_list): New macro.
	(gfc_omp_map_op): Add OMP_MAP_FORCE_ALLOC, OMP_MAP_FORCE_DEALLOC,
	OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM, OMP_MAP_FORCE_TOFROM,
	OMP_MAP_FORCE_PRESENT, OMP_MAP_FORCE_DEVICEPTR.
	(OMP_LIST_FIRST, OMP_LIST_DEVICE_RESIDENT, OMP_LIST_USE_DEVICE)
	(OMP_LIST_CACHE): New enumerators.
	(struct gfc_omp_clauses): Add async_expr, gang_expr, worker_expr,
	vector_expr, num_gangs_expr, num_workers_expr, vector_length_expr,
	wait_list, tile_list, async, gang, worker, vector, seq,
	independent, wait, par_auto, gang_static, and loc members.
	(struct gfc_namespace): Add oacc_declare_clauses member.
	(gfc_exec_op): Add EXEC_OACC_KERNELS_LOOP,
	EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	(gfc_free_expr_list, gfc_resolve_oacc_directive)
	(gfc_resolve_oacc_declare, gfc_resolve_oacc_parallel_loop_blocks)
	(gfc_resolve_oacc_blocks): New prototypes.
	* match.c (match_exit_cycle): Handle EXEC_OACC_LOOP and
	EXEC_OACC_PARALLEL_LOOP.
	* match.h (gfc_match_oacc_cache, gfc_match_oacc_wait)
	(gfc_match_oacc_update, gfc_match_oacc_declare)
	(gfc_match_oacc_loop, gfc_match_oacc_host_data)
	(gfc_match_oacc_data, gfc_match_oacc_kernels)
	(gfc_match_oacc_kernels_loop, gfc_match_oacc_parallel)
	(gfc_match_oacc_parallel_loop, gfc_match_oacc_enter_data)
	(gfc_match_oacc_exit_data, gfc_match_oacc_routine): New
	prototypes.
	* openmp.c: Include "diagnostic.h" and "gomp-constants.h".
	(gfc_free_omp_clauses): Update for members added to struct
	gfc_omp_clauses.
	(gfc_match_omp_clauses): Change mask paramter to uint64_t.  Add
	openacc parameter.
	(resolve_omp_clauses): Add openacc parameter.  Update for OpenACC.
	(struct fortran_omp_context): Add is_openmp member.
	(gfc_resolve_omp_parallel_blocks): Initialize it.
	(gfc_resolve_do_iterator): Update for OpenACC.
	(gfc_resolve_omp_directive): Call
	resolve_omp_directive_inside_oacc_region.
	(OMP_CLAUSE_PRIVATE, OMP_CLAUSE_FIRSTPRIVATE)
	(OMP_CLAUSE_LASTPRIVATE, OMP_CLAUSE_COPYPRIVATE)
	(OMP_CLAUSE_SHARED, OMP_CLAUSE_COPYIN, OMP_CLAUSE_REDUCTION)
	(OMP_CLAUSE_IF, OMP_CLAUSE_NUM_THREADS, OMP_CLAUSE_SCHEDULE)
	(OMP_CLAUSE_DEFAULT, OMP_CLAUSE_ORDERED, OMP_CLAUSE_COLLAPSE)
	(OMP_CLAUSE_UNTIED, OMP_CLAUSE_FINAL, OMP_CLAUSE_MERGEABLE)
	(OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND, OMP_CLAUSE_INBRANCH)
	(OMP_CLAUSE_LINEAR, OMP_CLAUSE_NOTINBRANCH, OMP_CLAUSE_PROC_BIND)
	(OMP_CLAUSE_SAFELEN, OMP_CLAUSE_SIMDLEN, OMP_CLAUSE_UNIFORM)
	(OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO)
	(OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT)
	(OMP_CLAUSE_DIST_SCHEDULE): Use uint64_t.
	(OMP_CLAUSE_ASYNC, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS)
	(OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_COPY, OMP_CLAUSE_COPYOUT)
	(OMP_CLAUSE_CREATE, OMP_CLAUSE_PRESENT)
	(OMP_CLAUSE_PRESENT_OR_COPY, OMP_CLAUSE_PRESENT_OR_COPYIN)
	(OMP_CLAUSE_PRESENT_OR_COPYOUT, OMP_CLAUSE_PRESENT_OR_CREATE)
	(OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER)
	(OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT)
	(OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_DEVICE_RESIDENT)
	(OMP_CLAUSE_HOST_SELF, OMP_CLAUSE_OACC_DEVICE, OMP_CLAUSE_WAIT)
	(OMP_CLAUSE_DELETE, OMP_CLAUSE_AUTO, OMP_CLAUSE_TILE): New macros.
	(gfc_match_omp_clauses): Handle those.
	(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
	(OACC_LOOP_CLAUSES, OACC_PARALLEL_LOOP_CLAUSES)
	(OACC_KERNELS_LOOP_CLAUSES, OACC_HOST_DATA_CLAUSES)
	(OACC_DECLARE_CLAUSES, OACC_UPDATE_CLAUSES)
	(OACC_ENTER_DATA_CLAUSES, OACC_EXIT_DATA_CLAUSES)
	(OACC_WAIT_CLAUSES): New macros.
	(gfc_free_expr_list, match_oacc_expr_list, match_oacc_clause_gang)
	(gfc_match_omp_map_clause, gfc_match_oacc_parallel_loop)
	(gfc_match_oacc_parallel, gfc_match_oacc_kernels_loop)
	(gfc_match_oacc_kernels, gfc_match_oacc_data)
	(gfc_match_oacc_host_data, gfc_match_oacc_loop)
	(gfc_match_oacc_declare, gfc_match_oacc_update)
	(gfc_match_oacc_enter_data, gfc_match_oacc_exit_data)
	(gfc_match_oacc_wait, gfc_match_oacc_cache)
	(gfc_match_oacc_routine, oacc_is_loop)
	(resolve_oacc_scalar_int_expr, resolve_oacc_positive_int_expr)
	(check_symbol_not_pointer, check_array_not_assumed)
	(resolve_oacc_data_clauses, resolve_oacc_deviceptr_clause)
	(oacc_compatible_clauses, oacc_is_parallel, oacc_is_kernels)
	(omp_code_to_statement, oacc_code_to_statement)
	(resolve_oacc_directive_inside_omp_region)
	(resolve_omp_directive_inside_oacc_region)
	(resolve_oacc_nested_loops, resolve_oacc_params_in_parallel)
	(resolve_oacc_loop_blocks, gfc_resolve_oacc_blocks)
	(resolve_oacc_loop, resolve_oacc_cache, gfc_resolve_oacc_declare)
	(gfc_resolve_oacc_directive): New functions.
	* parse.c (next_free): Update for OpenACC.  Move some code into...
	(verify_token_free): ... this new function.
	(next_fixed): Update for OpenACC.  Move some code into...
	(verify_token_fixed): ... this new function.
	(case_executable): Add ST_OACC_UPDATE, ST_OACC_WAIT,
	ST_OACC_CACHE, ST_OACC_ENTER_DATA, and ST_OACC_EXIT_DATA.
	(case_exec_markers): Add ST_OACC_PARALLEL_LOOP, ST_OACC_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA, ST_OACC_LOOP,
	ST_OACC_KERNELS_LOOP.
	(case_decl): Add ST_OACC_ROUTINE.
	(push_state, parse_critical_block, parse_progunit): Update for
	OpenACC.
	(gfc_ascii_statement): Handle ST_OACC_PARALLEL_LOOP,
	ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_KERNELS_LOOP,
	ST_OACC_END_KERNELS_LOOP, ST_OACC_DATA, ST_OACC_END_DATA,
	ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP,
	ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
	ST_OACC_CACHE, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
	ST_OACC_ROUTINE.
	(verify_st_order, parse_spec): Handle ST_OACC_DECLARE.
	(parse_executable): Handle ST_OACC_PARALLEL_LOOP,
	ST_OACC_KERNELS_LOOP, ST_OACC_LOOP, ST_OACC_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA.
	(decode_oacc_directive, parse_oacc_structured_block)
	(parse_oacc_loop, is_oacc): New functions.
	* parse.h (struct gfc_state_data): Add oacc_declare_clauses
	member.
	(is_oacc): New prototype.
	* resolve.c (gfc_resolve_blocks, gfc_resolve_code): Handle
	EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL,
	EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS, EXEC_OACC_DATA,
	EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE,
	EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA,
	EXEC_OACC_EXIT_DATA.
	(resolve_codes): Call gfc_resolve_oacc_declare.
	* scanner.c (openacc_flag, openacc_locus): New variables.
	(skip_free_comments): Update for OpenACC.  Move some code into...
	(skip_omp_attribute): ... this new function.
	(skip_oacc_attribute): New function.
	(skip_fixed_comments, gfc_next_char_literal): Update for OpenACC.
	* st.c (gfc_free_statement): Handle EXEC_OACC_PARALLEL_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	* trans-decl.c (gfc_generate_function_code): Update for OpenACC.
	* trans-openmp.c: Include "gomp-constants.h".
	(gfc_omp_finish_clause, gfc_trans_omp_clauses): Use GOMP_MAP_*
	instead of OMP_CLAUSE_MAP_*.  Use OMP_CLAUSE_SET_MAP_KIND.
	(gfc_trans_omp_clauses): Handle OMP_LIST_USE_DEVICE,
	OMP_LIST_DEVICE_RESIDENT, OMP_LIST_CACHE, and OMP_MAP_FORCE_ALLOC,
	OMP_MAP_FORCE_DEALLOC, OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM,
	OMP_MAP_FORCE_TOFROM, OMP_MAP_FORCE_PRESENT,
	OMP_MAP_FORCE_DEVICEPTR, and gfc_omp_clauses' async, seq,
	independent, wait_list, num_gangs_expr, num_workers_expr,
	vector_length_expr, vector, vector_expr, worker, worker_expr,
	gang, gang_expr members.
	(gfc_trans_omp_do): Handle EXEC_OACC_LOOP.
	(gfc_convert_expr_to_tree, gfc_trans_oacc_construct)
	(gfc_trans_oacc_executable_directive)
	(gfc_trans_oacc_wait_directive, gfc_trans_oacc_combined_directive)
	(gfc_trans_oacc_declare, gfc_trans_oacc_directive): New functions.
	* trans-stmt.c (gfc_trans_block_construct): Update for OpenACC.
	* trans-stmt.h (gfc_trans_oacc_directive, gfc_trans_oacc_declare):
	New prototypes.
	* trans.c (tranc_code): Handle EXEC_OACC_CACHE, EXEC_OACC_WAIT,
	EXEC_OACC_UPDATE, EXEC_OACC_LOOP, EXEC_OACC_HOST_DATA,
	EXEC_OACC_DATA, EXEC_OACC_KERNELS, EXEC_OACC_KERNELS_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ENTER_DATA,
	EXEC_OACC_EXIT_DATA.
	* gfortran.texi: Update for OpenACC.
	* intrinsic.texi: Likewise.
	* invoke.texi: Likewise.
	gcc/lto/
	* lto-lang.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
	New macros.
	* lto.c: Include "gomp-constants.h".
	gcc/testsuite/
	* lib/target-supports.exp (check_effective_target_fopenacc): New
	procedure.
	* g++.dg/goacc-gomp/goacc-gomp.exp: New file.
	* g++.dg/goacc/goacc.exp: Likewise.
	* gcc.dg/goacc-gomp/goacc-gomp.exp: Likewise.
	* gcc.dg/goacc/goacc.exp: Likewise.
	* gfortran.dg/goacc/goacc.exp: Likewise.
	* c-c++-common/cpp/openacc-define-1.c: New file.
	* c-c++-common/cpp/openacc-define-2.c: Likewise.
	* c-c++-common/cpp/openacc-define-3.c: Likewise.
	* c-c++-common/goacc-gomp/nesting-1.c: Likewise.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/acc_on_device-2-off.c: Likewise.
	* c-c++-common/goacc/acc_on_device-2.c: Likewise.
	* c-c++-common/goacc/asyncwait-1.c: Likewise.
	* c-c++-common/goacc/cache-1.c: Likewise.
	* c-c++-common/goacc/clauses-fail.c: Likewise.
	* c-c++-common/goacc/collapse-1.c: Likewise.
	* c-c++-common/goacc/data-1.c: Likewise.
	* c-c++-common/goacc/data-2.c: Likewise.
	* c-c++-common/goacc/data-clause-duplicate-1.c: Likewise.
	* c-c++-common/goacc/deviceptr-1.c: Likewise.
	* c-c++-common/goacc/deviceptr-2.c: Likewise.
	* c-c++-common/goacc/deviceptr-3.c: Likewise.
	* c-c++-common/goacc/if-clause-1.c: Likewise.
	* c-c++-common/goacc/if-clause-2.c: Likewise.
	* c-c++-common/goacc/kernels-1.c: Likewise.
	* c-c++-common/goacc/loop-1.c: Likewise.
	* c-c++-common/goacc/loop-private-1.c: Likewise.
	* c-c++-common/goacc/nesting-1.c: Likewise.
	* c-c++-common/goacc/nesting-data-1.c: Likewise.
	* c-c++-common/goacc/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/parallel-1.c: Likewise.
	* c-c++-common/goacc/pcopy.c: Likewise.
	* c-c++-common/goacc/pcopyin.c: Likewise.
	* c-c++-common/goacc/pcopyout.c: Likewise.
	* c-c++-common/goacc/pcreate.c: Likewise.
	* c-c++-common/goacc/pragma_context.c: Likewise.
	* c-c++-common/goacc/present-1.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.
	* c-c++-common/goacc/sb-1.c: Likewise.
	* c-c++-common/goacc/sb-2.c: Likewise.
	* c-c++-common/goacc/sb-3.c: Likewise.
	* c-c++-common/goacc/update-1.c: Likewise.
	* gcc.dg/goacc/acc_on_device-1.c: Likewise.
	* gfortran.dg/goacc/acc_on_device-1.f95: Likewise.
	* gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise.
	* gfortran.dg/goacc/acc_on_device-2.f95: Likewise.
	* gfortran.dg/goacc/assumed.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-1.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-2.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-3.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-4.f95: Likewise.
	* gfortran.dg/goacc/branch.f95: Likewise.
	* gfortran.dg/goacc/cache-1.f95: Likewise.
	* gfortran.dg/goacc/coarray.f95: Likewise.
	* gfortran.dg/goacc/continuation-free-form.f95: Likewise.
	* gfortran.dg/goacc/cray.f95: Likewise.
	* gfortran.dg/goacc/critical.f95: Likewise.
	* gfortran.dg/goacc/data-clauses.f95: Likewise.
	* gfortran.dg/goacc/data-tree.f95: Likewise.
	* gfortran.dg/goacc/declare-1.f95: Likewise.
	* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
	* gfortran.dg/goacc/fixed-1.f: Likewise.
	* gfortran.dg/goacc/fixed-2.f: Likewise.
	* gfortran.dg/goacc/fixed-3.f: Likewise.
	* gfortran.dg/goacc/fixed-4.f: Likewise.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.
	* gfortran.dg/goacc/if.f95: Likewise.
	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
	* gfortran.dg/goacc/list.f95: Likewise.
	* gfortran.dg/goacc/literal.f95: Likewise.
	* gfortran.dg/goacc/loop-1.f95: Likewise.
	* gfortran.dg/goacc/loop-2.f95: Likewise.
	* gfortran.dg/goacc/loop-3.f95: Likewise.
	* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
	* gfortran.dg/goacc/omp.f95: Likewise.
	* gfortran.dg/goacc/parallel-kernels-clauses.f95: Likewise.
	* gfortran.dg/goacc/parallel-kernels-regions.f95: Likewise.
	* gfortran.dg/goacc/parallel-tree.f95: Likewise.
	* gfortran.dg/goacc/parameter.f95: Likewise.
	* gfortran.dg/goacc/private-1.f95: Likewise.
	* gfortran.dg/goacc/private-2.f95: Likewise.
	* gfortran.dg/goacc/private-3.f95: Likewise.
	* gfortran.dg/goacc/pure-elemental-procedures.f95: Likewise.
	* gfortran.dg/goacc/reduction-2.f95: Likewise.
	* gfortran.dg/goacc/reduction.f95: Likewise.
	* gfortran.dg/goacc/routine-1.f90: Likewise.
	* gfortran.dg/goacc/routine-2.f90: Likewise.
	* gfortran.dg/goacc/sentinel-free-form.f95: Likewise.
	* gfortran.dg/goacc/several-directives.f95: Likewise.
	* gfortran.dg/goacc/sie.f95: Likewise.
	* gfortran.dg/goacc/subarrays.f95: Likewise.
	* gfortran.dg/gomp/map-1.f90: Likewise.
	* gfortran.dg/openacc-define-1.f90: Likewise.
	* gfortran.dg/openacc-define-2.f90: Likewise.
	* gfortran.dg/openacc-define-3.f90: Likewise.
	* g++.dg/gomp/block-1.C: Update for changed compiler output.
	* g++.dg/gomp/block-2.C: Likewise.
	* g++.dg/gomp/block-3.C: Likewise.
	* g++.dg/gomp/block-5.C: Likewise.
	* g++.dg/gomp/target-1.C: Likewise.
	* g++.dg/gomp/target-2.C: Likewise.
	* g++.dg/gomp/taskgroup-1.C: Likewise.
	* g++.dg/gomp/teams-1.C: Likewise.
	* gcc.dg/cilk-plus/jump-openmp.c: Likewise.
	* gcc.dg/cilk-plus/jump.c: Likewise.
	* gcc.dg/gomp/block-1.c: Likewise.
	* gcc.dg/gomp/block-10.c: Likewise.
	* gcc.dg/gomp/block-2.c: Likewise.
	* gcc.dg/gomp/block-3.c: Likewise.
	* gcc.dg/gomp/block-4.c: Likewise.
	* gcc.dg/gomp/block-5.c: Likewise.
	* gcc.dg/gomp/block-6.c: Likewise.
	* gcc.dg/gomp/block-7.c: Likewise.
	* gcc.dg/gomp/block-8.c: Likewise.
	* gcc.dg/gomp/block-9.c: Likewise.
	* gcc.dg/gomp/target-1.c: Likewise.
	* gcc.dg/gomp/target-2.c: Likewise.
	* gcc.dg/gomp/taskgroup-1.c: Likewise.
	* gcc.dg/gomp/teams-1.c: Likewise.
	include/
	* gomp-constants.h: New file.
	libgomp/
	* Makefile.am (search_path): Add $(top_srcdir)/../include.
	(libgomp_la_SOURCES): Add splay-tree.c, libgomp-plugin.c,
	oacc-parallel.c, oacc-host.c, oacc-init.c, oacc-mem.c,
	oacc-async.c, oacc-plugin.c, oacc-cuda.c.
	[USE_FORTRAN] (libgomp_la_SOURCES): Add openacc.f90.
	Include $(top_srcdir)/plugin/Makefrag.am.
	(nodist_libsubinclude_HEADERS): Add openacc.h.
	[USE_FORTRAN] (nodist_finclude_HEADERS): Add openacc_lib.h,
	openacc.f90, openacc.mod, openacc_kinds.mod.
	(omp_lib.mod): Generalize into...
	(%.mod): ... this new rule.
	(openacc_kinds.mod, openacc.mod): New rules.
	* plugin/configfrag.ac: New file.
	* configure.ac: Move plugin/offloading support into it.  Include
	it.  Instantiate testsuite/libgomp-test-support.pt.exp.
	* plugin/Makefrag.am: New file.
	* testsuite/Makefile.am (OFFLOAD_TARGETS)
	(OFFLOAD_ADDITIONAL_OPTIONS, OFFLOAD_ADDITIONAL_LIB_PATHS): Don't
	export.
	(libgomp-test-support.exp): New rule.
	(all-local): Depend on it.
	* Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.
	* config.h.in: Likewise.
	* configure: Likewise.
	* configure.tgt: Harden shell syntax.
	* env.c: Include "oacc-int.h".
	(parse_acc_device_type): New function.
	(gomp_debug_var, goacc_device_type, goacc_device_num): New
	variables.
	(initialize_env): Initialize those.  Call
	goacc_runtime_initialize.
	* error.c (gomp_vdebug, gomp_debug, gomp_vfatal): New functions.
	(gomp_fatal): Call gomp_vfatal.
	* libgomp.h: Include "libgomp-plugin.h" and <stdarg.h>.
	(gomp_debug_var, goacc_device_type, goacc_device_num, gomp_vdebug)
	(gomp_debug, gomp_verror, gomp_vfatal, gomp_init_targets_once)
	(splay_tree_node, splay_tree, splay_tree_key)
	(struct target_mem_desc, struct splay_tree_key_s)
	(struct gomp_memory_mapping, struct acc_dispatch_t)
	(struct gomp_device_descr, gomp_acc_insert_pointer)
	(gomp_acc_remove_pointer, target_mem_desc, gomp_copy_from_async)
	(gomp_unmap_vars, gomp_init_device, gomp_init_tables)
	(gomp_free_memmap, gomp_fini_device): New declarations.
	(gomp_vdebug, gomp_debug): New macros.
	Include "splay-tree.h".
	* libgomp.map (OACC_2.0): New symbol version.  Use for
	acc_get_num_devices, acc_get_num_devices_h_, acc_set_device_type,
	acc_set_device_type_h_, acc_get_device_type,
	acc_get_device_type_h_, acc_set_device_num, acc_set_device_num_h_,
	acc_get_device_num, acc_get_device_num_h_, acc_async_test,
	acc_async_test_h_, acc_async_test_all, acc_async_test_all_h_,
	acc_wait, acc_wait_h_, acc_wait_async, acc_wait_async_h_,
	acc_wait_all, acc_wait_all_h_, acc_wait_all_async,
	acc_wait_all_async_h_, acc_init, acc_init_h_, acc_shutdown,
	acc_shutdown_h_, acc_on_device, acc_on_device_h_, acc_malloc,
	acc_free, acc_copyin, acc_copyin_32_h_, acc_copyin_64_h_,
	acc_copyin_array_h_, acc_present_or_copyin,
	acc_present_or_copyin_32_h_, acc_present_or_copyin_64_h_,
	acc_present_or_copyin_array_h_, acc_create, acc_create_32_h_,
	acc_create_64_h_, acc_create_array_h_, acc_present_or_create,
	acc_present_or_create_32_h_, acc_present_or_create_64_h_,
	acc_present_or_create_array_h_, acc_copyout, acc_copyout_32_h_,
	acc_copyout_64_h_, acc_copyout_array_h_, acc_delete,
	acc_delete_32_h_, acc_delete_64_h_, acc_delete_array_h_,
	acc_update_device, acc_update_device_32_h_,
	acc_update_device_64_h_, acc_update_device_array_h_,
	acc_update_self, acc_update_self_32_h_, acc_update_self_64_h_,
	acc_update_self_array_h_, acc_map_data, acc_unmap_data,
	acc_deviceptr, acc_hostptr, acc_is_present, acc_is_present_32_h_,
	acc_is_present_64_h_, acc_is_present_array_h_,
	acc_memcpy_to_device, acc_memcpy_from_device,
	acc_get_current_cuda_device, acc_get_current_cuda_context,
	acc_get_cuda_stream, acc_set_cuda_stream.
	(GOACC_2.0): New symbol version.  Use for GOACC_data_end,
	GOACC_data_start, GOACC_enter_exit_data, GOACC_parallel,
	GOACC_update, GOACC_wait, GOACC_get_thread_num,
	GOACC_get_num_threads.
	(GOMP_PLUGIN_1.0): New symbol version.  Use for
	GOMP_PLUGIN_malloc, GOMP_PLUGIN_malloc_cleared,
	GOMP_PLUGIN_realloc, GOMP_PLUGIN_debug, GOMP_PLUGIN_error,
	GOMP_PLUGIN_fatal, GOMP_PLUGIN_async_unmap_vars,
	GOMP_PLUGIN_acc_thread.
	* libgomp.texi: Update for OpenACC changes, and GOMP_DEBUG
	environment variable.
	* libgomp_g.h (GOACC_data_start, GOACC_data_end)
	(GOACC_enter_exit_data, GOACC_parallel, GOACC_update, GOACC_wait)
	(GOACC_get_num_threads, GOACC_get_thread_num): New declarations.
	* splay-tree.h (splay_tree_lookup, splay_tree_insert)
	(splay_tree_remove): New declarations.
	(rotate_left, rotate_right, splay_tree_splay, splay_tree_insert)
	(splay_tree_remove, splay_tree_lookup): Move into...
	* splay-tree.c: ... this new file.
	* target.c: Include "oacc-plugin.h", "oacc-int.h", <assert.h>.
	(splay_tree_node, splay_tree, splay_tree_key)
	(struct target_mem_desc, struct splay_tree_key_s)
	(struct gomp_device_descr): Don't declare.
	(num_devices_openmp): New variable.
	(gomp_get_num_devices ): Use it.
	(gomp_init_targets_once): New function.
	(gomp_get_num_devices ): Use it.
	(get_kind, gomp_copy_from_async, gomp_free_memmap)
	(gomp_fini_device, gomp_register_image_for_device): New functions.
	(gomp_map_vars): Add devaddrs parameter.
	(gomp_update): Add mm parameter.
	(gomp_init_device): Move most of it into...
	(gomp_init_tables): ... this new function.
	(gomp_register_images_for_device): Remove function.
	(splay_compare, gomp_map_vars, gomp_unmap_vars, gomp_init_device):
	Make them hidden instead of static.
	(gomp_map_vars_existing, gomp_map_vars, gomp_unmap_vars)
	(gomp_update, gomp_init_device, GOMP_target, GOMP_target_data)
	(GOMP_target_end_data, GOMP_target_update)
	(gomp_load_plugin_for_device, gomp_target_init): Update for
	OpenACC changes.
	* oacc-async.c: New file.
	* oacc-cuda.c: Likewise.
	* oacc-host.c: Likewise.
	* oacc-init.c: Likewise.
	* oacc-int.h: Likewise.
	* oacc-mem.c: Likewise.
	* oacc-parallel.c: Likewise.
	* oacc-plugin.c: Likewise.
	* oacc-plugin.h: Likewise.
	* oacc-ptx.h: Likewise.
	* openacc.f90: Likewise.
	* openacc.h: Likewise.
	* openacc_lib.h: Likewise.
	* plugin/plugin-host.c: Likewise.
	* plugin/plugin-nvptx.c: Likewise.
	* libgomp-plugin.c: Likewise.
	* libgomp-plugin.h: Likewise.
	* libgomp_target.h: Remove file after merging content into the
	former file.  Update all users.
	* testsuite/lib/libgomp.exp: Load libgomp-test-support.exp.
	(offload_targets_s, offload_targets_s_openacc): New variables.
	(check_effective_target_openacc_nvidia_accel_present)
	(check_effective_target_openacc_nvidia_accel_selected): New
	procedures.
	(libgomp_init): Update for OpenACC changes.
	* testsuite/libgomp-test-support.exp.in: New file.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/cache-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-1.c: Likewise.
	* 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/data-already-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-10.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-11.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-12.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-19.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-26.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-27.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-31.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-33.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-35.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-36.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-37.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-39.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-40.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-41.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-45.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-46.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-49.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-50.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-51.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-55.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-56.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-57.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-58.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-59.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-6.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-60.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-61.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-62.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-63.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-64.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-65.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-66.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-67.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-68.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-7.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-c-c++-common/lib-84.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-86.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-87.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-88.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-89.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-9.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-90.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-92.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/offset-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-empty.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pointer-align-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-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-c-c++-common/reduction-initial-1.c:
	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: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Likewise.
	* testsuite/libgomp.oacc-fortran/abort-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/abort-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-4.f90: 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-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-fortran/data-already-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-10.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-5.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-6.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-7.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-8.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/map-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/openacc_version-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/pointer-align-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/pset-1.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-fortran/routine-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/subarrays-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/subarrays-2.f90: Likewise.
	liboffloadmic/
	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_name)
	(GOMP_OFFLOAD_get_caps, GOMP_OFFLOAD_fini_device): New functions.

Co-Authored-By: Bernd Schmidt <bernds@codesourcery.com>
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Dmitry Bocharnikov <dmitry.b@samsung.com>
Co-Authored-By: Evgeny Gavrin <e.gavrin@samsung.com>
Co-Authored-By: Ilmir Usmanov <i.usmanov@samsung.com>
Co-Authored-By: Jakub Jelinek <jakub@redhat.com>
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Tobias Burnus <burnus@net-b.de>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r219682

1 files changed, 1017 insertions, 19 deletions

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index d309465..665ee42 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -81,6 +81,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "plugin.h"
 #include "omp-low.h"
 #include "builtins.h"
+#include "gomp-constants.h"
 
 
 /* Initialization routine for this file.  */
@@ -1249,10 +1250,15 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
 					     vec<tree, va_gc> **, location_t *,
 					     tree *, vec<location_t> *,
 					     unsigned int * = NULL);
+static void c_parser_oacc_enter_exit_data (c_parser *, bool);
+static void c_parser_oacc_update (c_parser *);
+static tree c_parser_oacc_loop (location_t, c_parser *, char *);
 static void c_parser_omp_construct (c_parser *);
 static void c_parser_omp_threadprivate (c_parser *);
 static void c_parser_omp_barrier (c_parser *);
 static void c_parser_omp_flush (c_parser *);
+static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code,
+				   tree, tree *);
 static void c_parser_omp_taskwait (c_parser *);
 static void c_parser_omp_taskyield (c_parser *);
 static void c_parser_omp_cancel (c_parser *);
@@ -4492,6 +4498,14 @@ c_parser_initval (c_parser *parser, struct c_expr *after,
    Although they are erroneous if the labels declared aren't defined,
    is it useful for the syntax to be this way?
 
+   OpenACC:
+
+   block-item:
+     openacc-directive
+
+   openacc-directive:
+     update-directive
+
    OpenMP:
 
    block-item:
@@ -4838,6 +4852,29 @@ c_parser_label (c_parser *parser)
      @throw expression ;
      @throw ;
 
+   OpenACC:
+
+   statement:
+     openacc-construct
+
+   openacc-construct:
+     parallel-construct
+     kernels-construct
+     data-construct
+     loop-construct
+
+   parallel-construct:
+     parallel-directive structured-block
+
+   kernels-construct:
+     kernels-directive structured-block
+
+   data-construct:
+     data-directive structured-block
+
+   loop-construct:
+     loop-directive structured-block
+
    OpenMP:
 
    statement:
@@ -9570,6 +9607,25 @@ c_parser_pragma (c_parser *parser, enum pragma_context context)
 
   switch (id)
     {
+    case PRAGMA_OACC_ENTER_DATA:
+      c_parser_oacc_enter_exit_data (parser, true);
+      return false;
+
+    case PRAGMA_OACC_EXIT_DATA:
+      c_parser_oacc_enter_exit_data (parser, false);
+      return false;
+
+    case PRAGMA_OACC_UPDATE:
+      if (context != pragma_compound)
+	{
+	  if (context == pragma_stmt)
+	    c_parser_error (parser, "%<#pragma acc update%> may only be "
+			    "used in compound statements");
+	  goto bad_stmt;
+	}
+      c_parser_oacc_update (parser);
+      return false;
+
     case PRAGMA_OMP_BARRIER:
       if (context != pragma_compound)
 	{
@@ -9772,7 +9828,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser)
     c_common_pch_pragma (parse_in, TREE_STRING_POINTER (name));
 }
 
-/* OpenMP 2.5 / 3.0 / 3.1 / 4.0 parsing routines.  */
+/* OpenACC and OpenMP parsing routines.  */
 
 /* Returns name of the next clause.
    If the clause is not recognized PRAGMA_OMP_CLAUSE_NONE is returned and
@@ -9784,7 +9840,9 @@ c_parser_omp_clause_name (c_parser *parser)
 {
   pragma_omp_clause result = PRAGMA_OMP_CLAUSE_NONE;
 
-  if (c_parser_next_token_is_keyword (parser, RID_IF))
+  if (c_parser_next_token_is_keyword (parser, RID_AUTO))
+    result = PRAGMA_OACC_CLAUSE_AUTO;
+  else if (c_parser_next_token_is_keyword (parser, RID_IF))
     result = PRAGMA_OMP_CLAUSE_IF;
   else if (c_parser_next_token_is_keyword (parser, RID_DEFAULT))
     result = PRAGMA_OMP_CLAUSE_DEFAULT;
@@ -9799,20 +9857,32 @@ c_parser_omp_clause_name (c_parser *parser)
 	case 'a':
 	  if (!strcmp ("aligned", p))
 	    result = PRAGMA_OMP_CLAUSE_ALIGNED;
+	  else if (!strcmp ("async", p))
+	    result = PRAGMA_OACC_CLAUSE_ASYNC;
 	  break;
 	case 'c':
 	  if (!strcmp ("collapse", p))
 	    result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+	  else if (!strcmp ("copy", p))
+	    result = PRAGMA_OACC_CLAUSE_COPY;
 	  else if (!strcmp ("copyin", p))
 	    result = PRAGMA_OMP_CLAUSE_COPYIN;
+	  else if (!strcmp ("copyout", p))
+	    result = PRAGMA_OACC_CLAUSE_COPYOUT;
           else if (!strcmp ("copyprivate", p))
 	    result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+	  else if (!strcmp ("create", p))
+	    result = PRAGMA_OACC_CLAUSE_CREATE;
 	  break;
 	case 'd':
-	  if (!strcmp ("depend", p))
+	  if (!strcmp ("delete", p))
+	    result = PRAGMA_OACC_CLAUSE_DELETE;
+	  else if (!strcmp ("depend", p))
 	    result = PRAGMA_OMP_CLAUSE_DEPEND;
 	  else if (!strcmp ("device", p))
 	    result = PRAGMA_OMP_CLAUSE_DEVICE;
+	  else if (!strcmp ("deviceptr", p))
+	    result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
 	  else if (!strcmp ("dist_schedule", p))
 	    result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
 	  break;
@@ -9824,6 +9894,14 @@ c_parser_omp_clause_name (c_parser *parser)
 	  else if (!strcmp ("from", p))
 	    result = PRAGMA_OMP_CLAUSE_FROM;
 	  break;
+	case 'g':
+	  if (!strcmp ("gang", p))
+	    result = PRAGMA_OACC_CLAUSE_GANG;
+	  break;
+	case 'h':
+	  if (!strcmp ("host", p))
+	    result = PRAGMA_OACC_CLAUSE_HOST;
+	  break;
 	case 'i':
 	  if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -9847,10 +9925,14 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
 	  else if (!strcmp ("nowait", p))
 	    result = PRAGMA_OMP_CLAUSE_NOWAIT;
+	  else if (!strcmp ("num_gangs", p))
+	    result = PRAGMA_OACC_CLAUSE_NUM_GANGS;
 	  else if (!strcmp ("num_teams", p))
 	    result = PRAGMA_OMP_CLAUSE_NUM_TEAMS;
 	  else if (!strcmp ("num_threads", p))
 	    result = PRAGMA_OMP_CLAUSE_NUM_THREADS;
+	  else if (!strcmp ("num_workers", p))
+	    result = PRAGMA_OACC_CLAUSE_NUM_WORKERS;
 	  else if (flag_cilkplus && !strcmp ("nomask", p))
 	    result = PRAGMA_CILK_CLAUSE_NOMASK;
 	  break;
@@ -9861,6 +9943,20 @@ c_parser_omp_clause_name (c_parser *parser)
 	case 'p':
 	  if (!strcmp ("parallel", p))
 	    result = PRAGMA_OMP_CLAUSE_PARALLEL;
+	  else if (!strcmp ("present", p))
+	    result = PRAGMA_OACC_CLAUSE_PRESENT;
+	  else if (!strcmp ("present_or_copy", p)
+		   || !strcmp ("pcopy", p))
+	    result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+	  else if (!strcmp ("present_or_copyin", p)
+		   || !strcmp ("pcopyin", p))
+	    result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+	  else if (!strcmp ("present_or_copyout", p)
+		   || !strcmp ("pcopyout", p))
+	    result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+	  else if (!strcmp ("present_or_create", p)
+		   || !strcmp ("pcreate", p))
+	    result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
 	  else if (!strcmp ("private", p))
 	    result = PRAGMA_OMP_CLAUSE_PRIVATE;
 	  else if (!strcmp ("proc_bind", p))
@@ -9877,10 +9973,14 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_SCHEDULE;
 	  else if (!strcmp ("sections", p))
 	    result = PRAGMA_OMP_CLAUSE_SECTIONS;
+	  else if (!strcmp ("seq", p))
+	    result = PRAGMA_OACC_CLAUSE_SEQ;
 	  else if (!strcmp ("shared", p))
 	    result = PRAGMA_OMP_CLAUSE_SHARED;
 	  else if (!strcmp ("simdlen", p))
 	    result = PRAGMA_OMP_CLAUSE_SIMDLEN;
+	  else if (!strcmp ("self", p))
+	    result = PRAGMA_OACC_CLAUSE_SELF;
 	  break;
 	case 't':
 	  if (!strcmp ("taskgroup", p))
@@ -9897,9 +9997,19 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_UNTIED;
 	  break;
 	case 'v':
-	  if (flag_cilkplus && !strcmp ("vectorlength", p))
+	  if (!strcmp ("vector", p))
+	    result = PRAGMA_OACC_CLAUSE_VECTOR;
+	  else if (!strcmp ("vector_length", p))
+	    result = PRAGMA_OACC_CLAUSE_VECTOR_LENGTH;
+	  else if (flag_cilkplus && !strcmp ("vectorlength", p))
 	    result = PRAGMA_CILK_CLAUSE_VECTORLENGTH;
 	  break;
+	case 'w':
+	  if (!strcmp ("wait", p))
+	    result = PRAGMA_OACC_CLAUSE_WAIT;
+	  else if (!strcmp ("worker", p))
+	    result = PRAGMA_OACC_CLAUSE_WORKER;
+	  break;
 	}
     }
 
@@ -9926,7 +10036,57 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
       }
 }
 
-/* OpenMP 2.5:
+/* OpenACC 2.0
+   Parse wait clause or wait directive parameters.  */
+
+static tree
+c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
+{
+  vec<tree, va_gc> *args;
+  tree t, args_tree;
+
+  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    return list;
+
+  args = c_parser_expr_list (parser, false, true, NULL, NULL, NULL, NULL);
+
+  if (args->length () == 0)
+    {
+      c_parser_error (parser, "expected integer expression before ')'");
+      release_tree_vector (args);
+      return list;
+    }
+
+  args_tree = build_tree_list_vec (args);
+
+  for (t = args_tree; t; t = TREE_CHAIN (t))
+    {
+      tree targ = TREE_VALUE (t);
+
+      if (targ != error_mark_node)
+	{
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (targ)))
+	    {
+	      c_parser_error (parser, "expression must be integral");
+	      targ = error_mark_node;
+	    }
+	  else
+	    {
+	      tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
+
+	      OMP_CLAUSE_DECL (c) = targ;
+	      OMP_CLAUSE_CHAIN (c) = list;
+	      list = c;
+	    }
+	}
+    }
+
+  release_tree_vector (args);
+  c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+  return list;
+}
+
+/* OpenACC 2.0, OpenMP 2.5:
    variable-list:
      identifier
      variable-list , identifier
@@ -9967,6 +10127,14 @@ c_parser_omp_variable_list (c_parser *parser,
 	{
 	  switch (kind)
 	    {
+	    case OMP_CLAUSE__CACHE_:
+	      if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE)
+		{
+		  c_parser_error (parser, "expected %<[%>");
+		  t = error_mark_node;
+		  break;
+		}
+	      /* FALL THROUGH.  */
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
@@ -10005,6 +10173,26 @@ c_parser_omp_variable_list (c_parser *parser,
 		      t = error_mark_node;
 		      break;
 		    }
+
+		  if (kind == OMP_CLAUSE__CACHE_)
+		    {
+		      if (TREE_CODE (low_bound) != INTEGER_CST
+			  && !TREE_READONLY (low_bound))
+			{
+			  error_at (clause_loc,
+					"%qD is not a constant", low_bound);
+			  t = error_mark_node;
+			}
+
+		      if (TREE_CODE (length) != INTEGER_CST
+			  && !TREE_READONLY (length))
+			{
+			  error_at (clause_loc,
+					"%qD is not a constant", length);
+			  t = error_mark_node;
+			}
+		    }
+
 		  t = tree_cons (low_bound, length, t);
 		}
 	      break;
@@ -10033,7 +10221,7 @@ c_parser_omp_variable_list (c_parser *parser,
 }
 
 /* Similarly, but expect leading and trailing parenthesis.  This is a very
-   common case for omp clauses.  */
+   common case for OpenACC and OpenMP clauses.  */
 
 static tree
 c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
@@ -10050,7 +10238,119 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
   return list;
 }
 
-/* OpenMP 3.0:
+/* OpenACC 2.0:
+   copy ( variable-list )
+   copyin ( variable-list )
+   copyout ( variable-list )
+   create ( variable-list )
+   delete ( variable-list )
+   present ( variable-list )
+   present_or_copy ( variable-list )
+     pcopy ( variable-list )
+   present_or_copyin ( variable-list )
+     pcopyin ( variable-list )
+   present_or_copyout ( variable-list )
+     pcopyout ( variable-list )
+   present_or_create ( variable-list )
+     pcreate ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
+			   tree list)
+{
+  enum gomp_map_kind kind;
+  switch (c_kind)
+    {
+    case PRAGMA_OACC_CLAUSE_COPY:
+      kind = GOMP_MAP_FORCE_TOFROM;
+      break;
+    case PRAGMA_OACC_CLAUSE_COPYIN:
+      kind = GOMP_MAP_FORCE_TO;
+      break;
+    case PRAGMA_OACC_CLAUSE_COPYOUT:
+      kind = GOMP_MAP_FORCE_FROM;
+      break;
+    case PRAGMA_OACC_CLAUSE_CREATE:
+      kind = GOMP_MAP_FORCE_ALLOC;
+      break;
+    case PRAGMA_OACC_CLAUSE_DELETE:
+      kind = GOMP_MAP_FORCE_DEALLOC;
+      break;
+    case PRAGMA_OACC_CLAUSE_DEVICE:
+      kind = GOMP_MAP_FORCE_TO;
+      break;
+    case PRAGMA_OACC_CLAUSE_HOST:
+    case PRAGMA_OACC_CLAUSE_SELF:
+      kind = GOMP_MAP_FORCE_FROM;
+      break;
+    case PRAGMA_OACC_CLAUSE_PRESENT:
+      kind = GOMP_MAP_FORCE_PRESENT;
+      break;
+    case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
+      kind = GOMP_MAP_TOFROM;
+      break;
+    case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
+      kind = GOMP_MAP_TO;
+      break;
+    case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
+      kind = GOMP_MAP_FROM;
+      break;
+    case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
+      kind = GOMP_MAP_ALLOC;
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  tree nl, c;
+  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
+
+  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+
+  return nl;
+}
+
+/* OpenACC 2.0:
+   deviceptr ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+  tree vars, t;
+
+  /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic
+     c_parser_oacc_data_clause), as for PRAGMA_OACC_CLAUSE_DEVICEPTR,
+     variable-list must only allow for pointer variables.  */
+  vars = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_ERROR, NULL);
+  for (t = vars; t && t; t = TREE_CHAIN (t))
+    {
+      tree v = TREE_PURPOSE (t);
+
+      /* FIXME diagnostics: Ideally we should keep individual
+	 locations for all the variables in the var list to make the
+	 following errors more precise.  Perhaps
+	 c_parser_omp_var_list_parens() should construct a list of
+	 locations to go along with the var list.  */
+
+      if (TREE_CODE (v) != VAR_DECL)
+	error_at (loc, "%qD is not a variable", v);
+      else if (TREE_TYPE (v) == error_mark_node)
+	;
+      else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+	error_at (loc, "%qD is not a pointer variable", v);
+
+      tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (u, GOMP_MAP_FORCE_DEVICEPTR);
+      OMP_CLAUSE_DECL (u) = v;
+      OMP_CLAUSE_CHAIN (u) = list;
+      list = u;
+    }
+
+  return list;
+}
+
+/* OpenACC 2.0, OpenMP 3.0:
    collapse ( constant-expression ) */
 
 static tree
@@ -10193,7 +10493,7 @@ c_parser_omp_clause_final (c_parser *parser, tree list)
   return list;
 }
 
-/* OpenMP 2.5:
+/* OpenACC, OpenMP 2.5:
    if ( expression ) */
 
 static tree
@@ -10261,6 +10561,51 @@ c_parser_omp_clause_nowait (c_parser *parser ATTRIBUTE_UNUSED, tree list)
   return c;
 }
 
+/* OpenACC:
+   num_gangs ( expression ) */
+
+static tree
+c_parser_omp_clause_num_gangs (c_parser *parser, tree list)
+{
+  location_t num_gangs_loc = c_parser_peek_token (parser)->location;
+  if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    {
+      location_t expr_loc = c_parser_peek_token (parser)->location;
+      tree c, t = c_parser_expression (parser).value;
+      mark_exp_read (t);
+      t = c_fully_fold (t, false, NULL);
+
+      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  return list;
+	}
+
+      /* Attempt to statically determine when the number isn't positive.  */
+      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
+		       build_int_cst (TREE_TYPE (t), 0));
+      if (CAN_HAVE_LOCATION_P (c))
+	SET_EXPR_LOCATION (c, expr_loc);
+      if (c == boolean_true_node)
+	{
+	  warning_at (expr_loc, 0,
+		      "%<num_gangs%> value must be positive");
+	  t = integer_one_node;
+	}
+
+      check_no_duplicate_clause (list, OMP_CLAUSE_NUM_GANGS, "num_gangs");
+
+      c = build_omp_clause (num_gangs_loc, OMP_CLAUSE_NUM_GANGS);
+      OMP_CLAUSE_NUM_GANGS_EXPR (c) = t;
+      OMP_CLAUSE_CHAIN (c) = list;
+      list = c;
+    }
+
+  return list;
+}
+
 /* OpenMP 2.5:
    num_threads ( expression ) */
 
@@ -10306,6 +10651,100 @@ c_parser_omp_clause_num_threads (c_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC:
+   num_workers ( expression ) */
+
+static tree
+c_parser_omp_clause_num_workers (c_parser *parser, tree list)
+{
+  location_t num_workers_loc = c_parser_peek_token (parser)->location;
+  if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    {
+      location_t expr_loc = c_parser_peek_token (parser)->location;
+      tree c, t = c_parser_expression (parser).value;
+      mark_exp_read (t);
+      t = c_fully_fold (t, false, NULL);
+
+      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  return list;
+	}
+
+      /* Attempt to statically determine when the number isn't positive.  */
+      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
+		       build_int_cst (TREE_TYPE (t), 0));
+      if (CAN_HAVE_LOCATION_P (c))
+	SET_EXPR_LOCATION (c, expr_loc);
+      if (c == boolean_true_node)
+	{
+	  warning_at (expr_loc, 0,
+		      "%<num_workers%> value must be positive");
+	  t = integer_one_node;
+	}
+
+      check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_workers");
+
+      c = build_omp_clause (num_workers_loc, OMP_CLAUSE_NUM_WORKERS);
+      OMP_CLAUSE_NUM_WORKERS_EXPR (c) = t;
+      OMP_CLAUSE_CHAIN (c) = list;
+      list = c;
+    }
+
+  return list;
+}
+
+/* OpenACC:
+   async [( int-expr )] */
+
+static tree
+c_parser_oacc_clause_async (c_parser *parser, tree list)
+{
+  tree c, t;
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  t = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    {
+      c_parser_consume_token (parser);
+
+      t = c_parser_expression (parser).value;
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+	c_parser_error (parser, "expected integer expression");
+      else if (t == error_mark_node
+	  || !c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+	return list;
+    }
+  else
+    t = c_fully_fold (t, false, NULL);
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async");
+
+  c = build_omp_clause (loc, OMP_CLAUSE_ASYNC);
+  OMP_CLAUSE_ASYNC_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  list = c;
+
+  return list;
+}
+
+/* OpenACC:
+   wait ( int-expr-list ) */
+
+static tree
+c_parser_oacc_clause_wait (c_parser *parser, tree list)
+{
+  location_t clause_loc = c_parser_peek_token (parser)->location;
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    list = c_parser_oacc_wait_list (parser, clause_loc, list);
+
+  return list;
+}
+
 /* OpenMP 2.5:
    ordered */
 
@@ -10557,6 +10996,51 @@ c_parser_omp_clause_untied (c_parser *parser ATTRIBUTE_UNUSED, tree list)
   return c;
 }
 
+/* OpenACC:
+   vector_length ( expression ) */
+
+static tree
+c_parser_omp_clause_vector_length (c_parser *parser, tree list)
+{
+  location_t vector_length_loc = c_parser_peek_token (parser)->location;
+  if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    {
+      location_t expr_loc = c_parser_peek_token (parser)->location;
+      tree c, t = c_parser_expression (parser).value;
+      mark_exp_read (t);
+      t = c_fully_fold (t, false, NULL);
+
+      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+	{
+	  c_parser_error (parser, "expected integer expression");
+	  return list;
+	}
+
+      /* Attempt to statically determine when the number isn't positive.  */
+      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
+		       build_int_cst (TREE_TYPE (t), 0));
+      if (CAN_HAVE_LOCATION_P (c))
+	SET_EXPR_LOCATION (c, expr_loc);
+      if (c == boolean_true_node)
+	{
+	  warning_at (expr_loc, 0,
+		      "%<vector_length%> value must be positive");
+	  t = integer_one_node;
+	}
+
+      check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length");
+
+      c = build_omp_clause (vector_length_loc, OMP_CLAUSE_VECTOR_LENGTH);
+      OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t;
+      OMP_CLAUSE_CHAIN (c) = list;
+      list = c;
+    }
+
+  return list;
+}
+
 /* OpenMP 4.0:
    inbranch
    notinbranch */
@@ -10898,7 +11382,7 @@ static tree
 c_parser_omp_clause_map (c_parser *parser, tree list)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
-  enum omp_clause_map_kind kind = OMP_CLAUSE_MAP_TOFROM;
+  enum gomp_map_kind kind = GOMP_MAP_TOFROM;
   tree nl, c;
 
   if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
@@ -10909,13 +11393,13 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
     {
       const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
       if (strcmp ("alloc", p) == 0)
-	kind = OMP_CLAUSE_MAP_ALLOC;
+	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = OMP_CLAUSE_MAP_TO;
+	kind = GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = OMP_CLAUSE_MAP_FROM;
+	kind = GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = OMP_CLAUSE_MAP_TOFROM;
+	kind = GOMP_MAP_TOFROM;
       else
 	{
 	  c_parser_error (parser, "invalid map kind");
@@ -10930,7 +11414,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_MAP_KIND (c) = kind;
+    OMP_CLAUSE_SET_MAP_KIND (c, kind);
 
   c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
   return nl;
@@ -11093,9 +11577,154 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list)
   return list;
 }
 
+/* Parse all OpenACC clauses.  The set clauses allowed by the directive
+   is a bitmask in MASK.  Return the list of clauses found.  */
+
+static tree
+c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
+			   const char *where, bool finish_p = true)
+{
+  tree clauses = NULL;
+  bool first = true;
+
+  while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+    {
+      location_t here;
+      pragma_omp_clause c_kind;
+      const char *c_name;
+      tree prev = clauses;
+
+      if (!first && c_parser_next_token_is (parser, CPP_COMMA))
+	c_parser_consume_token (parser);
+
+      here = c_parser_peek_token (parser)->location;
+      c_kind = c_parser_omp_clause_name (parser);
+
+      switch (c_kind)
+	{
+	case PRAGMA_OACC_CLAUSE_ASYNC:
+	  clauses = c_parser_oacc_clause_async (parser, clauses);
+	  c_name = "async";
+	  break;
+	case PRAGMA_OACC_CLAUSE_COLLAPSE:
+	  clauses = c_parser_omp_clause_collapse (parser, clauses);
+	  c_name = "collapse";
+	  break;
+	case PRAGMA_OACC_CLAUSE_COPY:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "copy";
+	  break;
+	case PRAGMA_OACC_CLAUSE_COPYIN:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "copyin";
+	  break;
+	case PRAGMA_OACC_CLAUSE_COPYOUT:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "copyout";
+	  break;
+	case PRAGMA_OACC_CLAUSE_CREATE:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "create";
+	  break;
+	case PRAGMA_OACC_CLAUSE_DELETE:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "delete";
+	  break;
+	case PRAGMA_OACC_CLAUSE_DEVICE:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "device";
+	  break;
+	case PRAGMA_OACC_CLAUSE_DEVICEPTR:
+	  clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
+	  c_name = "deviceptr";
+	  break;
+	case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
+	  clauses = c_parser_omp_clause_firstprivate (parser, clauses);
+	  c_name = "firstprivate";
+	  break;
+	case PRAGMA_OACC_CLAUSE_HOST:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "host";
+	  break;
+	case PRAGMA_OACC_CLAUSE_IF:
+	  clauses = c_parser_omp_clause_if (parser, clauses);
+	  c_name = "if";
+	  break;
+	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
+	  clauses = c_parser_omp_clause_num_gangs (parser, clauses);
+	  c_name = "num_gangs";
+	  break;
+	case PRAGMA_OACC_CLAUSE_NUM_WORKERS:
+	  clauses = c_parser_omp_clause_num_workers (parser, clauses);
+	  c_name = "num_workers";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRESENT:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "present";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "present_or_copy";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "present_or_copyin";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "present_or_copyout";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "present_or_create";
+	  break;
+	case PRAGMA_OACC_CLAUSE_PRIVATE:
+	  clauses = c_parser_omp_clause_private (parser, clauses);
+	  c_name = "private";
+	  break;
+	case PRAGMA_OACC_CLAUSE_REDUCTION:
+	  clauses = c_parser_omp_clause_reduction (parser, clauses);
+	  c_name = "reduction";
+	  break;
+	case PRAGMA_OACC_CLAUSE_SELF:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "self";
+	  break;
+	case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
+	  clauses = c_parser_omp_clause_vector_length (parser, clauses);
+	  c_name = "vector_length";
+	  break;
+	case PRAGMA_OACC_CLAUSE_WAIT:
+	  clauses = c_parser_oacc_clause_wait (parser, clauses);
+	  c_name = "wait";
+	  break;
+	default:
+	  c_parser_error (parser, "expected %<#pragma acc%> clause");
+	  goto saw_error;
+	}
+
+      first = false;
+
+      if (((mask >> c_kind) & 1) == 0 && !parser->error)
+	{
+	  /* Remove the invalid clause(s) from the list to avoid
+	     confusing the rest of the compiler.  */
+	  clauses = prev;
+	  error_at (here, "%qs is not valid for %qs", c_name, where);
+	}
+    }
+
+ saw_error:
+  c_parser_skip_to_pragma_eol (parser);
+
+  if (finish_p)
+    return c_finish_omp_clauses (clauses);
+
+  return clauses;
+}
+
 /* Parse all OpenMP clauses.  The set clauses allowed by the directive
-   is a bitmask in MASK.  Return the list of clauses found; the result
-   of clause default goes in *pdefault.  */
+   is a bitmask in MASK.  Return the list of clauses found.  */
 
 static tree
 c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
@@ -11323,7 +11952,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
   return clauses;
 }
 
-/* OpenMP 2.5:
+/* OpenACC 2.0, OpenMP 2.5:
    structured-block:
      statement
 
@@ -11339,6 +11968,351 @@ c_parser_omp_structured_block (c_parser *parser)
   return pop_stmt_list (stmt);
 }
 
+/* OpenACC 2.0:
+   # pragma acc cache (variable-list) new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+static tree
+c_parser_oacc_cache (location_t loc, c_parser *parser)
+{
+  tree stmt, clauses;
+
+  clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
+  clauses = c_finish_omp_clauses (clauses);
+
+  c_parser_skip_to_pragma_eol (parser);
+
+  stmt = make_node (OACC_CACHE);
+  TREE_TYPE (stmt) = void_type_node;
+  OACC_CACHE_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, loc);
+  add_stmt (stmt);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
+   # pragma acc data oacc-data-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_DATA_CLAUSE_MASK						\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
+
+static tree
+c_parser_oacc_data (location_t loc, c_parser *parser)
+{
+  tree stmt, clauses, block;
+
+  clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+				       "#pragma acc data");
+
+  block = c_begin_omp_parallel ();
+  add_stmt (c_parser_omp_structured_block (parser));
+
+  stmt = c_finish_oacc_data (loc, clauses, block);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
+   # pragma acc kernels oacc-kernels-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_KERNELS_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+static tree
+c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
+{
+  tree stmt, clauses = NULL_TREE, block;
+
+  strcat (p_name, " kernels");
+
+  if (c_parser_next_token_is (parser, CPP_NAME))
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      if (strcmp (p, "loop") == 0)
+	{
+	  c_parser_consume_token (parser);
+	  block = c_begin_omp_parallel ();
+	  c_parser_oacc_loop (loc, parser, p_name);
+	  stmt = c_finish_oacc_kernels (loc, clauses, block);
+	  OACC_KERNELS_COMBINED (stmt) = 1;
+	  return stmt;
+	}
+    }
+
+  clauses =  c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
+					p_name);
+
+  block = c_begin_omp_parallel ();
+  add_stmt (c_parser_omp_structured_block (parser));
+
+  stmt = c_finish_oacc_kernels (loc, clauses, block);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
+   # pragma acc enter data oacc-enter-data-clause[optseq] new-line
+
+   or
+
+   # pragma acc exit data oacc-exit-data-clause[optseq] new-line
+
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+#define OACC_EXIT_DATA_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) 		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+static void
+c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+  tree clauses, stmt;
+
+  c_parser_consume_pragma (parser);
+
+  if (!c_parser_next_token_is (parser, CPP_NAME))
+    {
+      c_parser_error (parser, enter
+		      ? "expected %<data%> in %<#pragma acc enter data%>"
+		      : "expected %<data%> in %<#pragma acc exit data%>");
+      c_parser_skip_to_pragma_eol (parser);
+      return;
+    }
+
+  const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+  if (strcmp (p, "data") != 0)
+    {
+      c_parser_error (parser, "invalid pragma");
+      c_parser_skip_to_pragma_eol (parser);
+      return;
+    }
+
+  c_parser_consume_token (parser);
+
+  if (enter)
+    clauses = c_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK,
+					 "#pragma acc enter data");
+  else
+    clauses = c_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK,
+					 "#pragma acc exit data");
+
+  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+    {
+      error_at (loc, enter
+		? "%<#pragma acc enter data%> has no data movement clause"
+		: "%<#pragma acc exit data%> has no data movement clause");
+      return;
+    }
+
+  stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA);;
+  TREE_TYPE (stmt) = void_type_node;
+  if (enter)
+    OACC_ENTER_DATA_CLAUSES (stmt) = clauses;
+  else
+    OACC_EXIT_DATA_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, loc);
+  add_stmt (stmt);
+}
+
+
+/* OpenACC 2.0:
+
+   # pragma acc loop oacc-loop-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_LOOP_CLAUSE_MASK						\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
+
+static tree
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+{
+  tree stmt, clauses, block;
+
+  strcat (p_name, " loop");
+
+  clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+
+  block = c_begin_compound_stmt (true);
+  stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+  block = c_end_compound_stmt (loc, block, true);
+  add_stmt (block);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
+   # pragma acc parallel oacc-parallel-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_PARALLEL_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+static tree
+c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
+{
+  tree stmt, clauses = NULL_TREE, block;
+
+  strcat (p_name, " parallel");
+
+  if (c_parser_next_token_is (parser, CPP_NAME))
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      if (strcmp (p, "loop") == 0)
+	{
+	  c_parser_consume_token (parser);
+	  block = c_begin_omp_parallel ();
+	  c_parser_oacc_loop (loc, parser, p_name);
+	  stmt = c_finish_oacc_parallel (loc, clauses, block);
+	  OACC_PARALLEL_COMBINED (stmt) = 1;
+	  return stmt;
+	}
+    }
+
+  clauses =  c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+					p_name);
+
+  block = c_begin_omp_parallel ();
+  add_stmt (c_parser_omp_structured_block (parser));
+
+  stmt = c_finish_oacc_parallel (loc, clauses, block);
+
+  return stmt;
+}
+
+/* OpenACC 2.0:
+   # pragma acc update oacc-update-clause[optseq] new-line
+*/
+
+#define OACC_UPDATE_CLAUSE_MASK						\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
+static void
+c_parser_oacc_update (c_parser *parser)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  c_parser_consume_pragma (parser);
+
+  tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK,
+					    "#pragma acc update");
+  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+    {
+      error_at (loc,
+		"%<#pragma acc update%> must contain at least one "
+		"%<device%> or %<host/self%> clause");
+      return;
+    }
+
+  if (parser->error)
+    return;
+
+  tree stmt = make_node (OACC_UPDATE);
+  TREE_TYPE (stmt) = void_type_node;
+  OACC_UPDATE_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, loc);
+  add_stmt (stmt);
+}
+
+/* OpenACC 2.0:
+   # pragma acc wait [(intseq)] oacc-wait-clause[optseq] new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_WAIT_CLAUSE_MASK						\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) )
+
+static tree
+c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name)
+{
+  tree clauses, list = NULL_TREE, stmt = NULL_TREE;
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    list = c_parser_oacc_wait_list (parser, loc, list);
+
+  strcpy (p_name, " wait");
+  clauses = c_parser_oacc_all_clauses (parser, OACC_WAIT_CLAUSE_MASK, p_name);
+  stmt = c_finish_oacc_wait (loc, list, clauses);
+
+  return stmt;
+}
+
 /* OpenMP 2.5:
    # pragma omp atomic new-line
      expression-stmt
@@ -11815,10 +12789,11 @@ c_parser_omp_flush (c_parser *parser)
   c_finish_omp_flush (loc);
 }
 
-/* Parse the restricted form of the for statement allowed by OpenMP.
+/* Parse the restricted form of loop statements allowed by OpenACC and OpenMP.
    The real trick here is to determine the loop control variable early
    so that we can push a new decl if necessary to make it private.
-   LOC is the location of the OMP in "#pragma omp".  */
+   LOC is the location of the "acc" or "omp" in "#pragma acc" or "#pragma omp",
+   respectively.  */
 
 static tree
 c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
@@ -13660,6 +14635,29 @@ c_parser_omp_construct (c_parser *parser)
 
   switch (p_kind)
     {
+    case PRAGMA_OACC_CACHE:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_cache (loc, parser);
+      break;
+    case PRAGMA_OACC_DATA:
+      stmt = c_parser_oacc_data (loc, parser);
+      break;
+    case PRAGMA_OACC_KERNELS:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_kernels (loc, parser, p_name);
+      break;
+    case PRAGMA_OACC_LOOP:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_loop (loc, parser, p_name);
+      break;
+    case PRAGMA_OACC_PARALLEL:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_parallel (loc, parser, p_name);
+      break;
+    case PRAGMA_OACC_WAIT:
+      strcpy (p_name, "#pragma wait");
+      stmt = c_parser_oacc_wait (loc, parser, p_name);
+      break;
     case PRAGMA_OMP_ATOMIC:
       c_parser_omp_atomic (loc, parser);
       return;