diff --git contrib/ChangeLog contrib/ChangeLog
index 1f02d95..f062ea9 100644
--- contrib/ChangeLog
+++ contrib/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gcc_update (files_and_dependencies): Update rules for new
+	libgomp/plugin/Makefrag.am and libgomp/plugin/configfrag.ac files.
+
 2015-01-12  Yury Gribov  <y.gribov@samsung.com>
 
 	* check_GNU_style.sh: Support patches coming from stdin.
diff --git gcc/ChangeLog gcc/ChangeLog
index d40a3f0..3a27df9 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,183 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    Bernd Schmidt  <bernds@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Tom de Vries  <tom@codesourcery.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+	    Dmitry Bocharnikov  <dmitry.b@samsung.com>
+	    Evgeny Gavrin  <e.gavrin@samsung.com>
+	    Jakub Jelinek  <jakub@redhat.com>
+
+	* 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.
+
 2015-01-15  Prathamesh Kulkarni  <prathamesh.kulkarni@linaro.org>
 
 	* explow.h: Remove duplicate contents.
diff --git gcc/ada/ChangeLog gcc/ada/ChangeLog
index 93efb49..c130f7d 100644
--- gcc/ada/ChangeLog
+++ gcc/ada/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_8)
+	(DEF_FUNCTION_TYPE_VAR_12): New macros.
+
 2015-01-09  Michael Collison  <michael.collison@linaro.org>
 
 	* gcc-interface/cuintp.c: Include hash-set.h, machmode.h,
diff --git gcc/builtin-types.def gcc/builtin-types.def
index 8ab4300..3412677 100644
--- gcc/builtin-types.def
+++ gcc/builtin-types.def
@@ -593,8 +593,9 @@ DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 			 BT_PTR, BT_INT, BT_INT)
 
 DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
-	 BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
-	 BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+			  BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+			  BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT,
+			  BT_INT, BT_INT)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_VAR, BT_FN_VOID_VAR)
 DEF_FUNCTION_TYPE_3 (BT_FN_PTR_PTR_FN_VOID_VAR_PTR_SIZE,
diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog
index 9764045..e0ad215 100644
--- gcc/c-family/ChangeLog
+++ gcc/c-family/ChangeLog
@@ -1,3 +1,41 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    Bernd Schmidt  <bernds@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+	    Jakub Jelinek  <jakub@redhat.com>
+
+	* 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.
+
 2015-01-14  Marcos Diaz <marcos.diaz@tallertechnologies.com>
 
 	* c-cppbuiltin.c (c_cpp_builtins): New cpp define __SSP_EXPLICIT__
diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index eea1e3a..8715045 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -714,7 +714,6 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
   enum c_omp_clause_split s;
   int i;
 
-  gcc_assert (code != OACC_PARALLEL);
   for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
     cclauses[i] = NULL;
   /* Add implicit nowait clause on
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index f42f53b..4652409 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,54 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    Bernd Schmidt  <bernds@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Jakub Jelinek  <jakub@redhat.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+
+	* 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.
+
 2015-01-09  Michael Collison  <michael.collison@linaro.org>
 
 	* c-array-notation.c: Include hash-set.h, machmode.h,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index fc6661b..665ee42 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1250,11 +1250,11 @@ 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_oacc_enter_exit_data (c_parser *, bool);
-static void c_parser_oacc_update (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,
@@ -11699,7 +11699,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  c_name = "wait";
 	  break;
 	default:
-	  c_parser_error (parser, "expected clause");
+	  c_parser_error (parser, "expected %<#pragma acc%> clause");
 	  goto saw_error;
 	}
 
@@ -11928,7 +11928,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  c_name = "simdlen";
 	  break;
 	default:
-	  c_parser_error (parser, "expected clause");
+	  c_parser_error (parser, "expected %<#pragma omp%> clause");
 	  goto saw_error;
 	}
 
@@ -13046,7 +13046,6 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
 	  if (cclauses != NULL
 	      && cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL)
 	    {
-	      gcc_assert (code != OACC_LOOP);
 	      tree *c;
 	      for (c = &cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL]; *c ; )
 		if (OMP_CLAUSE_CODE (*c) != OMP_CLAUSE_FIRSTPRIVATE
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index 02b4fac..543f4d9 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,3 +1,48 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+	    Jakub Jelinek  <jakub@redhat.com>
+
+	* 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.
+
 2015-01-14  Paolo Carlini  <paolo.carlini@oracle.com>
 
 	PR c++/58671
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 1ce7de7..bfa3d81 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -27653,8 +27653,6 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	  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))
 	    result = PRAGMA_OMP_CLAUSE_PROC_BIND;
 	  break;
@@ -29236,8 +29234,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "self";
 	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
-	  clauses =
-		cp_parser_oacc_clause_vector_length (parser, clauses);
+	  clauses = cp_parser_oacc_clause_vector_length (parser, clauses);
 	  c_name = "vector_length";
 	  break;
 	case PRAGMA_OACC_CLAUSE_WAIT:
@@ -29245,7 +29242,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "wait";
 	  break;
 	default:
-	  cp_parser_error (parser, "expected clause");
+	  cp_parser_error (parser, "expected %<#pragma acc%> clause");
 	  goto saw_error;
 	}
 
@@ -29496,7 +29493,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "simdlen";
 	  break;
 	default:
-	  cp_parser_error (parser, "expected clause");
+	  cp_parser_error (parser, "expected %<#pragma omp%> clause");
 	  goto saw_error;
 	}
 
diff --git gcc/doc/invoke.texi gcc/doc/invoke.texi
index 112eb81..510201a 100644
--- gcc/doc/invoke.texi
+++ gcc/doc/invoke.texi
@@ -1895,6 +1895,10 @@ Programming Interface v2.0 @w{@uref{http://www.openacc.org/}}.  This option
 implies @option{-pthread}, and thus is only supported on targets that
 have support for @option{-pthread}.
 
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@w{@uref{https://gcc.gnu.org/wiki/OpenACC}} for more information.
+
 @item -fopenmp
 @opindex fopenmp
 @cindex OpenMP parallel
diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog
index df4a2f3..d8b72a2 100644
--- gcc/fortran/ChangeLog
+++ gcc/fortran/ChangeLog
@@ -1,3 +1,199 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+	    Tobias Burnus  <burnus@net-b.de>
+
+	* 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.
+
 2015-01-15  Janus Weil  <janus@gcc.gnu.org>
 
 	PR fortran/58023
diff --git gcc/fortran/gfortran.texi gcc/fortran/gfortran.texi
index b9ea217..300b8b8 100644
--- gcc/fortran/gfortran.texi
+++ gcc/fortran/gfortran.texi
@@ -474,10 +474,13 @@ The GNU Fortran compiler is able to compile nearly all
 standard-compliant Fortran 95, Fortran 90, and Fortran 77 programs,
 including a number of standard and non-standard extensions, and can be
 used on real-world programs.  In particular, the supported extensions
-include OpenACC, OpenMP, Cray-style pointers, and several Fortran 2003
-and Fortran
+include OpenMP, Cray-style pointers, and several Fortran 2003 and Fortran
 2008 features, including TR 15581.  However, it is still under
 development and has a few remaining rough edges.
+There also is initial support for OpenACC.
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
 
 At present, the GNU Fortran compiler passes the
 @uref{http://www.fortran-2000.com/ArnaudRecipes/fcvs21_f95.html, 
@@ -532,9 +535,13 @@ The current status of the support is can be found in the
 @ref{Fortran 2003 status}, @ref{Fortran 2008 status} and
 @ref{TS 29113 status} sections of the documentation.
 
-Additionally, the GNU Fortran compilers supports the OpenACC specification
-(version 2.0, @url{http://www.openacc.org/}), and OpenMP specification
+Additionally, the GNU Fortran compilers supports the OpenMP specification
 (version 4.0, @url{http://openmp.org/@/wp/@/openmp-specifications/}).
+There also is initial support for the OpenACC specification (targeting
+version 2.0, @uref{http://www.openacc.org/}).
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
 
 @node Varying Length Character Strings
 @subsection Varying Length Character Strings
@@ -1377,8 +1384,8 @@ without warning.
 * Hollerith constants support::
 * Cray pointers::
 * CONVERT specifier::
-* OpenACC::
 * OpenMP::
+* OpenACC::
 * Argument list functions::
 @end menu
 
@@ -1893,33 +1900,6 @@ carries a significant speed overhead.  If speed in this area matters
 to you, it is best if you use this only for data that needs to be
 portable.
 
-@node OpenACC
-@subsection OpenACC
-@cindex OpenACC
-
-OpenACC is an application programming interface (API) that supports
-offloading of code to accelerator devices.  It consists of a set of
-compiler directives, library routines, and environment variables that
-influence run-time behavior.
-
-GNU Fortran strives to be compatible to the
-@uref{http://www.openacc.org/, OpenACC Application Programming
-Interface v2.0}.
-
-To enable the processing of the OpenACC directive @code{!$acc} in
-free-form source code; the @code{c$acc}, @code{*$acc} and @code{!$acc}
-directives in fixed form; the @code{!$} conditional compilation
-sentinels in free form; and the @code{c$}, @code{*$} and @code{!$}
-sentinels in fixed form, @command{gfortran} needs to be invoked with
-the @option{-fopenacc}.  This also arranges for automatic linking of
-the GNU Offloading and Multi Processing Runtime Library
-@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
-Library}.
-
-The OpenACC Fortran runtime library routines are provided both in a
-form of a Fortran 90 module named @code{openacc} and in a form of a
-Fortran @code{include} file named @file{openacc_lib.h}.
-
 @node OpenMP
 @subsection OpenMP
 @cindex OpenMP
@@ -1980,6 +1960,37 @@ to the command line.  However, this is not supported by @command{gcc} and
 thus not recommended.
 @end itemize
 
+@node OpenACC
+@subsection OpenACC
+@cindex OpenACC
+
+OpenACC is an application programming interface (API) that supports
+offloading of code to accelerator devices.  It consists of a set of
+compiler directives, library routines, and environment variables that
+influence run-time behavior.
+
+GNU Fortran strives to be compatible to the
+@uref{http://www.openacc.org/, OpenACC Application Programming
+Interface v2.0}.
+
+To enable the processing of the OpenACC directive @code{!$acc} in
+free-form source code; the @code{c$acc}, @code{*$acc} and @code{!$acc}
+directives in fixed form; the @code{!$} conditional compilation
+sentinels in free form; and the @code{c$}, @code{*$} and @code{!$}
+sentinels in fixed form, @command{gfortran} needs to be invoked with
+the @option{-fopenacc}.  This also arranges for automatic linking of
+the GNU Offloading and Multi Processing Runtime Library
+@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
+Library}.
+
+The OpenACC Fortran runtime library routines are provided both in a
+form of a Fortran 90 module named @code{openacc} and in a form of a
+Fortran @code{include} file named @file{openacc_lib.h}.
+
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
+
 @node Argument list functions
 @subsection Argument list functions @code{%VAL}, @code{%REF} and @code{%LOC}
 @cindex argument list functions
diff --git gcc/fortran/intrinsic.texi gcc/fortran/intrinsic.texi
index 34c525b..06bce15 100644
--- gcc/fortran/intrinsic.texi
+++ gcc/fortran/intrinsic.texi
@@ -13773,8 +13773,8 @@ Fortran 95 elemental function: @ref{IEOR}
 * ISO_FORTRAN_ENV::
 * ISO_C_BINDING::
 * IEEE modules::
-* OpenACC Module OPENACC::
 * OpenMP Modules OMP_LIB and OMP_LIB_KINDS::
+* OpenACC Module OPENACC::
 @end menu
 
 @node ISO_FORTRAN_ENV
@@ -14020,33 +14020,6 @@ with the following options: @code{-fno-unsafe-math-optimizations
 
 
 
-@node OpenACC Module OPENACC
-@section OpenACC Module @code{OPENACC}
-@table @asis
-@item @emph{Standard}:
-OpenACC Application Programming Interface v2.0
-@end table
-
-
-The OpenACC Fortran runtime library routines are provided both in a
-form of a Fortran 90 module, named @code{OPENACC}, and in form of a
-Fortran @code{include} file named @file{openacc_lib.h}.  The
-procedures provided by @code{OPENACC} can be found in the
-@ref{Top,,Introduction,libgomp,GNU Offloading and Multi Processing
-Runtime Library} manual, the named constants defined in the modules
-are listed below.
-
-For details refer to the actual
-@uref{http://www.openacc.org/,
-OpenACC Application Programming Interface v2.0}.
-
-@code{OPENACC} provides the scalar default-integer
-named constant @code{openacc_version} with a value of the form
-@var{yyyymm}, where @code{yyyy} is the year and @var{mm} the month
-of the OpenACC version; for OpenACC v2.0 the value is @code{201306}.
-
-
-
 @node OpenMP Modules OMP_LIB and OMP_LIB_KINDS
 @section OpenMP Modules @code{OMP_LIB} and @code{OMP_LIB_KINDS}
 @table @asis
@@ -14103,3 +14076,30 @@ kind @code{omp_proc_bind_kind}:
 @item @code{omp_proc_bind_close}
 @item @code{omp_proc_bind_spread}
 @end table
+
+
+
+@node OpenACC Module OPENACC
+@section OpenACC Module @code{OPENACC}
+@table @asis
+@item @emph{Standard}:
+OpenACC Application Programming Interface v2.0
+@end table
+
+
+The OpenACC Fortran runtime library routines are provided both in a
+form of a Fortran 90 module, named @code{OPENACC}, and in form of a
+Fortran @code{include} file named @file{openacc_lib.h}.  The
+procedures provided by @code{OPENACC} can be found in the
+@ref{Top,,Introduction,libgomp,GNU Offloading and Multi Processing
+Runtime Library} manual, the named constants defined in the modules
+are listed below.
+
+For details refer to the actual
+@uref{http://www.openacc.org/,
+OpenACC Application Programming Interface v2.0}.
+
+@code{OPENACC} provides the scalar default-integer
+named constant @code{openacc_version} with a value of the form
+@var{yyyymm}, where @code{yyyy} is the year and @var{mm} the month
+of the OpenACC version; for OpenACC v2.0 the value is @code{201306}.
diff --git gcc/fortran/invoke.texi gcc/fortran/invoke.texi
index 60ad5df..9228c78 100644
--- gcc/fortran/invoke.texi
+++ gcc/fortran/invoke.texi
@@ -312,6 +312,10 @@ compilation sentinels in free form and @code{c$}, @code{*$} and
 @code{!$} sentinels in fixed form, and when linking arranges for the
 OpenACC runtime library to be linked in.
 
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@w{@uref{https://gcc.gnu.org/wiki/OpenACC}} for more information.
+
 @item -fopenmp
 @opindex @code{fopenmp}
 @cindex OpenMP
diff --git gcc/fortran/types.def gcc/fortran/types.def
index 40cbc94..fdae28d 100644
--- gcc/fortran/types.def
+++ gcc/fortran/types.def
@@ -218,5 +218,6 @@ DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 			 BT_PTR, BT_INT, BT_INT)
 
 DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
-	 BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
-	 BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+			  BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+			  BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT,
+			  BT_INT, BT_INT)
diff --git gcc/lto/ChangeLog gcc/lto/ChangeLog
index 4eafacc..9b98edd 100644
--- gcc/lto/ChangeLog
+++ gcc/lto/ChangeLog
@@ -1,3 +1,10 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* lto-lang.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
+	New macros.
+	* lto.c: Include "gomp-constants.h".
+
 2015-01-14  Ilya Verbin  <ilya.verbin@intel.com>
 
 	* lto-partition.c (lto_promote_cross_file_statics): Remove argument
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index d9525d5..885a710 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,133 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    Ilmir Usmanov  <i.usmanov@samsung.com>
+
+	* 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.
+
 2015-01-15  David Malcolm  <dmalcolm@redhat.com>
 
 	* jit.dg/test-error-mismatching-types-in-assignment-op.c: New
diff --git gcc/testsuite/c-c++-common/goacc/asyncwait-1.c gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
index 9d69dcf..ccc0106 100644
--- gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
+++ gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
@@ -1,224 +1,159 @@
-/* { dg-skip-if "not yet" { c++ } } */
-
-#include <stdlib.h>
-
-int
-main (int argc, char **argv)
+void
+f (int N, float *a, float *b)
 {
-    int N = 64;
-    float *a, *b;
-    int i;
-
-    a = (float *) malloc (N * sizeof (float));
-    b = (float *) malloc (N * sizeof (float));
-
-    for (i = 0; i < N; i++)
-    {
-        a[i] = 3.0;
-        b[i] = 0.0;
-    }
+    int ii;
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (,1) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 /* { dg-error "expected '\\)' before end of line" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (*) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (a)
-	/* { dg-error "expected integer expression before" "" { target c } 85 } */
-	/* { dg-error "'async' expression must be integral" "" { target c++ } 85 } */
+	/* { dg-error "expected integer expression before" "" { target c } 54 } */
+	/* { dg-error "'async' expression must be integral" "" { target c++ } 54 } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1.0)
-	/* { dg-error "expected integer expression before" "" { target c } 95 } */
-	/* { dg-error "'async' expression must be integral" "" { target c++ } 95 } */
+	/* { dg-error "expected integer expression before" "" { target c } 62 } */
+	/* { dg-error "'async' expression must be integral" "" { target c++ } 62 } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async () /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (,1) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 /* { dg-error "expected '\\\)' before end of line" } */
+    /* { dg-error "expected integer expression before '\\\)'" "" { target c++ } 118 } */
     {
-	/* { dg-error "expected integer expression list before" "" { target c++ } 169 } */
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,*) /* { dg-error "expected (primary-|)expression before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,a) /*{ dg-error "must be integral" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (a) /* { dg-error "must be integral" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1.0) /* { dg-error "must be integral" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait () /* { dg-error "expected (integer |)expression (list |)before" } */
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
 
 #pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait
     {
-        int ii;
-
         for (ii = 0; ii < N; ii++)
             b[ii] = a[ii];
     }
@@ -236,7 +171,7 @@ main (int argc, char **argv)
 #pragma acc wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
 
 #pragma acc wait (1 /* { dg-error "expected '\\\)' before end of line" } */
-    /* { dg-error "expected integer expression list before" "" { target c++ } 238 } */
+    /* { dg-error "expected integer expression before '\\\)'" "" { target c++ } 173 } */
 
 #pragma acc wait (1,*) /* { dg-error "expected (primary-|)expression before" } */
 
@@ -246,9 +181,9 @@ main (int argc, char **argv)
 
 #pragma acc wait (1.0) /* { dg-error "expression must be integral" } */
 
-#pragma acc wait 1 /* { dg-error "expected clause before numeric constant" } */
+#pragma acc wait 1 /* { dg-error "expected '#pragma acc' clause before numeric constant" } */
 
-#pragma acc wait N /* { dg-error "expected clause before 'N'" } */
+#pragma acc wait N /* { dg-error "expected '#pragma acc' clause before 'N'" } */
 
 #pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
 
@@ -269,12 +204,10 @@ main (int argc, char **argv)
 #pragma acc wait async (*) /* { dg-error "expected (primary-|)expression before " } */
 
 #pragma acc wait async (a)
-    /* { dg-error "expected integer expression before" "" { target c } 271 } */
-    /* { dg-error "expression must be integral" "" { target c++ } 271 } */
+    /* { dg-error "expected integer expression before" "" { target c } 206 } */
+    /* { dg-error "expression must be integral" "" { target c++ } 206 } */
 
 #pragma acc wait async (1.0)
-   /* { dg-error "expected integer expression before" "" { target c } 275 } */
-   /* { dg-error "expression must be integral" "" { target c++ } 275 } */
-
-    return 0;
+   /* { dg-error "expected integer expression before" "" { target c } 210 } */
+   /* { dg-error "expression must be integral" "" { target c++ } 210 } */
 }
diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index e8e1278..8990180 100644
--- gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -3,16 +3,16 @@ f (void)
 {
   int i;
 
-#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
+#pragma acc parallel one /* { dg-error "expected '#pragma acc' clause before 'one'" } */
   ;
 
-#pragma acc kernels eins /* { dg-error "expected clause before 'eins'" } */
+#pragma acc kernels eins /* { dg-error "expected '#pragma acc' clause before 'eins'" } */
   ;
 
-#pragma acc data two /* { dg-error "expected clause before 'two'" } */
+#pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */
   ;
 
-#pragma acc loop deux /* { dg-error "expected clause before 'deux'" } */
+#pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */
   for (i = 0; i < 2; ++i)
     ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/data-2.c gcc/testsuite/c-c++-common/goacc/data-2.c
index 9c0a185..a67d8a4 100644
--- gcc/testsuite/c-c++-common/goacc/data-2.c
+++ gcc/testsuite/c-c++-common/goacc/data-2.c
@@ -5,7 +5,7 @@ foo (void)
   int n;
 #pragma acc enter data copyin (a, b) async wait
 #pragma acc enter data create (b[20:30]) async wait
-#pragma acc enter data (a) /* { dg-error "expected clause before '\\\(' token" } */
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
 #pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
 #pragma acc exit data delete (a) if (0)
 #pragma acc exit data copyout (b) if (a)
diff --git include/ChangeLog include/ChangeLog
index c1011b9..0917d94 100644
--- include/ChangeLog
+++ include/ChangeLog
@@ -1,3 +1,9 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    Julian Brown  <julian@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* gomp-constants.h: New file.
+
 2015-12-14  Jan-Benedict Glaw  <jbglaw@lug-owl.de>
 
 	* libiberty.h: Merge Copyright year update from Binutils.
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 6e1e141..9b003cb 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,354 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Tom de Vries  <tom@codesourcery.com>
+	    Julian Brown  <julian@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+	    Nathan Sidwell  <nathan@codesourcery.com>
+	    Tobias Burnus  <burnus@net-b.de>
+
+	* 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.
+
 2015-01-10  Thomas Schwinge  <thomas@codesourcery.com>
 	    Julian Brown  <julian@codesourcery.com>
 	    David Malcolm  <dmalcolm@redhat.com>
diff --git libgomp/Makefile.am libgomp/Makefile.am
index 4471fab..5411278 100644
--- libgomp/Makefile.am
+++ libgomp/Makefile.am
@@ -61,9 +61,9 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS)
 libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
 	iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
 	task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
-	time.c fortran.c affinity.c target.c oacc-parallel.c splay-tree.c \
-	oacc-host.c oacc-init.c oacc-mem.c oacc-async.c oacc-plugin.c \
-	oacc-cuda.c libgomp-plugin.c
+	time.c fortran.c affinity.c target.c 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
 
 include $(top_srcdir)/plugin/Makefrag.am
 
diff --git libgomp/Makefile.in libgomp/Makefile.in
index 8758989..b61b108 100644
--- libgomp/Makefile.in
+++ libgomp/Makefile.in
@@ -151,9 +151,9 @@ am_libgomp_la_OBJECTS = alloc.lo barrier.lo critical.lo env.lo \
 	error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \
 	parallel.lo sections.lo single.lo task.lo team.lo work.lo \
 	lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \
-	fortran.lo affinity.lo target.lo oacc-parallel.lo \
-	splay-tree.lo oacc-host.lo oacc-init.lo oacc-mem.lo \
-	oacc-async.lo oacc-plugin.lo oacc-cuda.lo libgomp-plugin.lo \
+	fortran.lo affinity.lo target.lo splay-tree.lo \
+	libgomp-plugin.lo oacc-parallel.lo oacc-host.lo oacc-init.lo \
+	oacc-mem.lo oacc-async.lo oacc-plugin.lo oacc-cuda.lo \
 	$(am__objects_1)
 libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
 DEFAULT_INCLUDES = -I.@am__isrc@
@@ -396,9 +396,9 @@ libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
 	iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c \
 	single.c task.c team.c work.c lock.c mutex.c proc.c sem.c \
 	bar.c ptrlock.c time.c fortran.c affinity.c target.c \
-	oacc-parallel.c splay-tree.c oacc-host.c oacc-init.c \
-	oacc-mem.c oacc-async.c oacc-plugin.c oacc-cuda.c \
-	libgomp-plugin.c $(am__append_2)
+	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 \
+	$(am__append_2)
 
 # Nvidia PTX OpenACC plugin.
 @PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_version_info = -version-info $(libtool_VERSION)
diff --git libgomp/env.c libgomp/env.c
index 130c52c..6b5e963 100644
--- libgomp/env.c
+++ libgomp/env.c
@@ -79,8 +79,8 @@ unsigned long gomp_bind_var_list_len;
 void **gomp_places_list;
 unsigned long gomp_places_list_len;
 int gomp_debug_var;
+char *goacc_device_type;
 int goacc_device_num;
-char* goacc_device_type;
 
 /* Parse the OMP_SCHEDULE environment variable.  */
 
@@ -1017,10 +1017,10 @@ parse_affinity (bool ignore)
 }
 
 static void
-goacc_parse_device_type (void)
+parse_acc_device_type (void)
 {
   const char *env = getenv ("ACC_DEVICE_TYPE");
-  
+
   if (env && *env != '\0')
     goacc_device_type = strdup (env);
   else
@@ -1287,14 +1287,14 @@ initialize_env (void)
     }
 
   handle_omp_display_env (stacksize, wait_policy);
-  
-  /* Look for OpenACC-specific environment variables.  */
+
+  /* OpenACC.  */
+
   if (!parse_int ("ACC_DEVICE_NUM", &goacc_device_num, true))
     goacc_device_num = 0;
 
-  goacc_parse_device_type ();
+  parse_acc_device_type ();
 
-  /* Initialize OpenACC-specific internal state.  */
   goacc_runtime_initialize ();
 }
 
diff --git libgomp/error.c libgomp/error.c
index fb96fb3..094c24a 100644
--- libgomp/error.c
+++ libgomp/error.c
@@ -38,7 +38,7 @@
 
 #undef gomp_vdebug
 void
-gomp_vdebug (int kind __attribute__((unused)), const char *msg, va_list list)
+gomp_vdebug (int kind __attribute__ ((unused)), const char *msg, va_list list)
 {
   if (gomp_debug_var)
     vfprintf (stderr, msg, list);
diff --git libgomp/libgomp-plugin.c libgomp/libgomp-plugin.c
index 1dd33f5..ffb22e9 100644
--- libgomp/libgomp-plugin.c
+++ libgomp/libgomp-plugin.c
@@ -53,7 +53,7 @@ void
 GOMP_PLUGIN_debug (int kind, const char *msg, ...)
 {
   va_list ap;
-  
+
   va_start (ap, msg);
   gomp_debug (kind, msg, ap);
   va_end (ap);
@@ -63,7 +63,7 @@ void
 GOMP_PLUGIN_error (const char *msg, ...)
 {
   va_list ap;
-  
+
   va_start (ap, msg);
   gomp_verror (msg, ap);
   va_end (ap);
@@ -73,11 +73,8 @@ void
 GOMP_PLUGIN_fatal (const char *msg, ...)
 {
   va_list ap;
-  
+
   va_start (ap, msg);
   gomp_vfatal (msg, ap);
   va_end (ap);
-  
-  /* Unreachable.  */
-  abort ();
 }
diff --git libgomp/libgomp-plugin.h libgomp/libgomp-plugin.h
index c8383e1..d9cbff5 100644
--- libgomp/libgomp-plugin.h
+++ libgomp/libgomp-plugin.h
@@ -62,16 +62,16 @@ struct mapping_table
 };
 
 /* Miscellaneous functions.  */
-extern void *GOMP_PLUGIN_malloc (size_t) __attribute__((malloc));
-extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__((malloc));
+extern void *GOMP_PLUGIN_malloc (size_t) __attribute__ ((malloc));
+extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__ ((malloc));
 extern void *GOMP_PLUGIN_realloc (void *, size_t);
 
 extern void GOMP_PLUGIN_debug (int, const char *, ...)
-	__attribute__((format (printf, 2, 3)));
+	__attribute__ ((format (printf, 2, 3)));
 extern void GOMP_PLUGIN_error (const char *, ...)
-	__attribute__((format (printf, 1, 2)));
+	__attribute__ ((format (printf, 1, 2)));
 extern void GOMP_PLUGIN_fatal (const char *, ...)
-	__attribute__((noreturn, format (printf, 1, 2)));
+	__attribute__ ((noreturn, format (printf, 1, 2)));
 
 #ifdef __cplusplus
 }
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 97732a5..ab8e229 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -260,7 +260,7 @@ extern void **gomp_places_list;
 extern unsigned long gomp_places_list_len;
 extern int gomp_debug_var;
 extern int goacc_device_num;
-extern char* goacc_device_type;
+extern char *goacc_device_type;
 
 enum gomp_task_kind
 {
@@ -542,7 +542,7 @@ extern void *gomp_realloc (void *, size_t);
 
 extern void gomp_vdebug (int, const char *, va_list);
 extern void gomp_debug (int, const char *, ...)
-	__attribute__((format (printf, 2, 3)));
+	__attribute__ ((format (printf, 2, 3)));
 #define gomp_vdebug(KIND, FMT, VALIST) \
   do { \
     if (__builtin_expect (gomp_debug_var, 0)) \
@@ -555,11 +555,11 @@ extern void gomp_debug (int, const char *, ...)
   } while (0)
 extern void gomp_verror (const char *, va_list);
 extern void gomp_error (const char *, ...)
-	__attribute__((format (printf, 1, 2)));
+	__attribute__ ((format (printf, 1, 2)));
 extern void gomp_vfatal (const char *, va_list)
-	__attribute__((noreturn));
+	__attribute__ ((noreturn));
 extern void gomp_fatal (const char *, ...)
-	__attribute__((noreturn, format (printf, 1, 2)));
+	__attribute__ ((noreturn, format (printf, 1, 2)));
 
 /* iter.c */
 
@@ -633,7 +633,9 @@ extern void gomp_free_thread (void *);
 extern void gomp_init_targets_once (void);
 extern int gomp_get_num_devices (void);
 
-#include "splay-tree.h"
+typedef struct splay_tree_node_s *splay_tree_node;
+typedef struct splay_tree_s *splay_tree;
+typedef struct splay_tree_key_s *splay_tree_key;
 
 struct target_mem_desc {
   /* Reference count.  */
@@ -662,18 +664,37 @@ struct target_mem_desc {
   splay_tree_key list[];
 };
 
+struct splay_tree_key_s {
+  /* Address of the host object.  */
+  uintptr_t host_start;
+  /* Address immediately after the host object.  */
+  uintptr_t host_end;
+  /* Descriptor of the target memory.  */
+  struct target_mem_desc *tgt;
+  /* Offset from tgt->tgt_start to the start of the target object.  */
+  uintptr_t tgt_offset;
+  /* Reference count.  */
+  uintptr_t refcount;
+  /* Asynchronous reference count.  */
+  uintptr_t async_refcount;
+  /* True if data should be copied from device to host at the end.  */
+  bool copy_from;
+};
+
+#include "splay-tree.h"
+
 /* Information about mapped memory regions (per device/context).  */
 
 struct gomp_memory_mapping
 {
-  /* Splay tree containing information about mapped memory regions.  */
-  struct splay_tree_s splay_tree;
-
   /* Mutex for operating with the splay tree and other shared structures.  */
   gomp_mutex_t lock;
 
   /* True when tables have been added to this memory map.  */
   bool is_initialized;
+
+  /* Splay tree containing information about mapped memory regions.  */
+  struct splay_tree_s splay_tree;
 };
 
 typedef struct acc_dispatch_t
@@ -747,12 +768,6 @@ struct gomp_device_descr
   /* This is the TYPE of device.  */
   enum offload_target_type type;
 
-  /* Set to true when device is initialized.  */
-  bool is_initialized;
-
-  /* True when offload regions have been registered with this device.  */
-  bool offload_regions_registered;
-
   /* Function handlers.  */
   const char *(*get_name_func) (void);
   unsigned int (*get_caps_func) (void);
@@ -768,17 +783,23 @@ struct gomp_device_descr
   void *(*host2dev_func) (int, void *, const void *, size_t);
   void (*run_func) (int, void *, void *);
 
-  /* OpenACC-specific functions.  */
-  /* This is mutable because of its mutable data_environ and target_data
-     members.  */
-  acc_dispatch_t openacc;
-
   /* Memory-mapping info for this device instance.  */
   /* Uses a separate lock.  */
   struct gomp_memory_mapping mem_map;
 
   /* Mutex for the mutable data.  */
   gomp_mutex_t lock;
+
+  /* Set to true when device is initialized.  */
+  bool is_initialized;
+
+  /* True when offload regions have been registered with this device.  */
+  bool offload_regions_registered;
+
+  /* OpenACC-specific data and functions.  */
+  /* This is mutable because of its mutable data_environ and target_data
+     members.  */
+  acc_dispatch_t openacc;
 };
 
 extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
@@ -787,18 +808,12 @@ extern void gomp_acc_remove_pointer (void *, bool, int, int);
 extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
 					      size_t, void **, void **,
 					      size_t *, void *, bool, bool);
-
 extern void gomp_copy_from_async (struct target_mem_desc *);
-
 extern void gomp_unmap_vars (struct target_mem_desc *, bool);
-
 extern void gomp_init_device (struct gomp_device_descr *);
-
 extern void gomp_init_tables (struct gomp_device_descr *,
 			      struct gomp_memory_mapping *);
-
 extern void gomp_free_memmap (struct gomp_memory_mapping *);
-
 extern void gomp_fini_device (struct gomp_device_descr *);
 
 /* work.c */
diff --git libgomp/libgomp.map libgomp/libgomp.map
index bfdb78c..59ed4ee 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -318,7 +318,6 @@ GOACC_2.0 {
 	GOACC_data_end;
 	GOACC_data_start;
 	GOACC_enter_exit_data;
-	GOACC_kernels;
 	GOACC_parallel;
 	GOACC_update;
 	GOACC_wait;
@@ -326,13 +325,18 @@ GOACC_2.0 {
 	GOACC_get_num_threads;
 };
 
+GOACC_2.0.GOMP_4_BRANCH {
+  global:
+	GOACC_kernels;
+} GOACC_2.0;
+
 GOMP_PLUGIN_1.0 {
   global:
 	GOMP_PLUGIN_malloc;
 	GOMP_PLUGIN_malloc_cleared;
 	GOMP_PLUGIN_realloc;
-	GOMP_PLUGIN_error;
 	GOMP_PLUGIN_debug;
+	GOMP_PLUGIN_error;
 	GOMP_PLUGIN_fatal;
 	GOMP_PLUGIN_async_unmap_vars;
 	GOMP_PLUGIN_acc_thread;
diff --git libgomp/libgomp.texi libgomp/libgomp.texi
index c4713ad..ff093d7 100644
--- libgomp/libgomp.texi
+++ libgomp/libgomp.texi
@@ -143,6 +143,11 @@ A complete description of all OpenACC directives accepted may be found in
 the @uref{http://www.openacc.org/, OpenMP Application Programming
 Interface} manual, version 2.0.
 
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC.  See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
+
+
 
 @c ---------------------------------------------------------------------
 @c OpenACC Runtime Library Routines
diff --git libgomp/oacc-cuda.c libgomp/oacc-cuda.c
index 6f1a06f..c8ef376 100644
--- libgomp/oacc-cuda.c
+++ libgomp/oacc-cuda.c
@@ -74,7 +74,7 @@ acc_set_cuda_stream (int async, void *stream)
 
   if (async < 0 || stream == NULL)
     return 0;
-  
+
   goacc_lazy_initialize ();
 
   if (base_dev && base_dev->openacc.cuda.set_stream_func)
diff --git libgomp/oacc-host.c libgomp/oacc-host.c
index 21c1ead..6aeb1e7 100644
--- libgomp/oacc-host.c
+++ libgomp/oacc-host.c
@@ -33,32 +33,31 @@
 static struct gomp_device_descr host_dispatch =
   {
     .name = "host",
-
-    .type = OFFLOAD_TARGET_TYPE_HOST,
     .capabilities = (GOMP_OFFLOAD_CAP_OPENACC_200
 		     | GOMP_OFFLOAD_CAP_NATIVE_EXEC
 		     | GOMP_OFFLOAD_CAP_SHARED_MEM),
-
-    .is_initialized = false,
-    .offload_regions_registered = false,
+    .target_id = 0,
+    .type = OFFLOAD_TARGET_TYPE_HOST,
 
     .get_name_func = GOMP_OFFLOAD_get_name,
-    .get_type_func = GOMP_OFFLOAD_get_type,
     .get_caps_func = GOMP_OFFLOAD_get_caps,
-
-    .init_device_func = GOMP_OFFLOAD_init_device,
-    .fini_device_func = GOMP_OFFLOAD_fini_device,
+    .get_type_func = GOMP_OFFLOAD_get_type,
     .get_num_devices_func = GOMP_OFFLOAD_get_num_devices,
     .register_image_func = GOMP_OFFLOAD_register_image,
+    .init_device_func = GOMP_OFFLOAD_init_device,
+    .fini_device_func = GOMP_OFFLOAD_fini_device,
     .get_table_func = GOMP_OFFLOAD_get_table,
-
     .alloc_func = GOMP_OFFLOAD_alloc,
     .free_func = GOMP_OFFLOAD_free,
-    .host2dev_func = GOMP_OFFLOAD_host2dev,
     .dev2host_func = GOMP_OFFLOAD_dev2host,
-
+    .host2dev_func = GOMP_OFFLOAD_host2dev,
     .run_func = GOMP_OFFLOAD_run,
 
+    .mem_map.is_initialized = false,
+    .mem_map.splay_tree.root = NULL,
+    .is_initialized = false,
+    .offload_regions_registered = false,
+
     .openacc = {
       .open_device_func = GOMP_OFFLOAD_openacc_open_device,
       .close_device_func = GOMP_OFFLOAD_openacc_close_device,
diff --git libgomp/oacc-init.c libgomp/oacc-init.c
index 6f4a32c..166eb55 100644
--- libgomp/oacc-init.c
+++ libgomp/oacc-init.c
@@ -33,13 +33,12 @@
 #include <stdlib.h>
 #include <strings.h>
 #include <stdbool.h>
-#include <stdio.h>
 #include <string.h>
 
 static gomp_mutex_t acc_device_lock;
 
 /* The dispatch table for the current accelerator device.  This is global, so
-   you can only have one type of device open at any given time in a program. 
+   you can only have one type of device open at any given time in a program.
    This is the "base" device in that several devices that use the same
    dispatch table may be active concurrently: this one (the "zeroth") is used
    for overall initialisation/shutdown, and other instances -- not necessarily
@@ -129,7 +128,7 @@ resolve_device (acc_device_t d)
 	if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
 	  goto found;
       if (d_arg == acc_device_default)
-	{	  
+	{
 	  d = acc_device_host;
 	  goto found;
 	}
@@ -171,7 +170,7 @@ acc_init_1 (acc_device_t d)
     gomp_fatal ("device already active");
 
   /* We need to remember what we were intialized as, to check shutdown etc.  */
-  init_key = d;  
+  init_key = d;
 
   gomp_init_device (acc_dev);
 
@@ -203,9 +202,9 @@ static void
 goacc_destroy_thread (void *data)
 {
   struct goacc_thread *thr = data, *walk, *prev;
-  
+
   gomp_mutex_lock (&goacc_thread_lock);
-  
+
   if (thr)
     {
       if (base_dev && thr->target_tls)
@@ -500,7 +499,7 @@ acc_get_device_num (acc_device_t d)
   num = dev->openacc.get_device_num_func ();
   if (num < 0)
     num = goacc_device_num;
-  
+
   return num;
 }
 
@@ -514,11 +513,11 @@ acc_set_device_num (int n, acc_device_t d)
 
   if (!base_dev)
     gomp_init_targets_once ();
-  
+
   if ((int) d == 0)
     {
       int i;
-      
+
       /* A device setting of zero sets all device types on the system to use
          the Nth instance of that device type.  Only attempt it for initialized
 	 devices though.  */
@@ -571,6 +570,7 @@ acc_on_device (acc_device_t dev)
   /* Just rely on the compiler builtin.  */
   return __builtin_acc_on_device (dev);
 }
+
 ialias (acc_on_device)
 
 attribute_hidden void
diff --git libgomp/oacc-int.h libgomp/oacc-int.h
index 700577d..85619c8 100644
--- libgomp/oacc-int.h
+++ libgomp/oacc-int.h
@@ -35,8 +35,8 @@
    that are part of the external ABI, and the lower case prefix "goacc"
    is used group items that are completely private to the library.  */
 
-#ifndef _OACC_INT_H
-#define _OACC_INT_H 1
+#ifndef OACC_INT_H
+#define OACC_INT_H 1
 
 #include "openacc.h"
 #include "config.h"
@@ -58,16 +58,16 @@ struct goacc_thread
 {
   /* The device for the current thread.  */
   struct gomp_device_descr *dev;
-  
+
   struct gomp_device_descr *saved_bound_dev;
 
   /* This is a linked list of data mapped by the "acc data" pragma, following
      strictly push/pop semantics according to lexical scope.  */
   struct target_mem_desc *mapped_data;
-    
+
   /* These structures form a list: this is the next thread in that list.  */
   struct goacc_thread *next;
-  
+
   /* Target-specific data (used by plugin).  */
   void *target_tls;
 };
@@ -88,8 +88,6 @@ goacc_thread (void)
 }
 #endif
 
-struct gomp_device_descr;
-
 void goacc_register (struct gomp_device_descr *) __GOACC_NOTHROW;
 
 /* Current dispatcher.  */
@@ -104,4 +102,4 @@ void goacc_lazy_initialize (void);
 # pragma GCC visibility pop
 #endif
 
-#endif /* _OACC_INT_H */
+#endif
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 674fb76..0096d51 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -31,15 +31,13 @@
 #include "libgomp.h"
 #include "gomp-constants.h"
 #include "oacc-int.h"
-#include <stdio.h>
+#include "splay-tree.h"
 #include <stdint.h>
 #include <assert.h>
 
-#include "splay-tree.h"
-
 /* Return block containing [H->S), or NULL if not contained.  */
 
-attribute_hidden splay_tree_key
+static splay_tree_key
 lookup_host (struct gomp_memory_mapping *mem_map, void *h, size_t s)
 {
   struct splay_tree_key_s node;
@@ -342,11 +340,11 @@ acc_unmap_data (void *h)
   gomp_unmap_vars (t, true);
 }
 
-#define PCC_Present (1 << 0)
-#define PCC_Create (1 << 1)
-#define PCC_Copy (1 << 2)
+#define FLAG_PRESENT (1 << 0)
+#define FLAG_CREATE (1 << 1)
+#define FLAG_COPY (1 << 2)
 
-attribute_hidden void *
+static void *
 present_create_copy (unsigned f, void *h, size_t s)
 {
   void *d;
@@ -366,13 +364,13 @@ present_create_copy (unsigned f, void *h, size_t s)
       /* Present. */
       d = (void *) (n->tgt->tgt_start + n->tgt_offset);
 
-      if (!(f & PCC_Present))
+      if (!(f & FLAG_PRESENT))
         gomp_fatal ("[%p,+%d] already mapped to [%p,+%d]",
             (void *)h, (int)s, (void *)d, (int)s);
       if ((h + s) > (void *)n->host_end)
         gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
     }
-  else if (!(f & PCC_Create))
+  else if (!(f & FLAG_CREATE))
     {
       gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
     }
@@ -383,7 +381,7 @@ present_create_copy (unsigned f, void *h, size_t s)
       unsigned short kinds;
       void *hostaddrs = h;
 
-      if (f & PCC_Copy)
+      if (f & FLAG_COPY)
 	kinds = GOMP_MAP_TO;
       else
 	kinds = GOMP_MAP_ALLOC;
@@ -406,28 +404,28 @@ present_create_copy (unsigned f, void *h, size_t s)
 void *
 acc_create (void *h, size_t s)
 {
-  return present_create_copy (PCC_Create, h, s);
+  return present_create_copy (FLAG_CREATE, h, s);
 }
 
 void *
 acc_copyin (void *h, size_t s)
 {
-  return present_create_copy (PCC_Create | PCC_Copy, h, s);
+  return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s);
 }
 
 void *
 acc_present_or_create (void *h, size_t s)
 {
-  return present_create_copy (PCC_Present | PCC_Create, h, s);
+  return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
 }
 
 void *
 acc_present_or_copyin (void *h, size_t s)
 {
-  return present_create_copy (PCC_Present | PCC_Create | PCC_Copy, h, s);
+  return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
 }
 
-#define DC_Copyout (1 << 0)
+#define FLAG_COPYOUT (1 << 0)
 
 static void
 delete_copyout (unsigned f, void *h, size_t s)
@@ -454,7 +452,7 @@ delete_copyout (unsigned f, void *h, size_t s)
     gomp_fatal ("[%p,%d] surrounds2 [%p,+%d]",
 		(void *) n->host_start, (int) host_size, (void *) h, (int) s);
 
-  if (f & DC_Copyout)
+  if (f & FLAG_COPYOUT)
     acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
 
   acc_unmap_data (h);
@@ -470,7 +468,7 @@ acc_delete (void *h , size_t s)
 
 void acc_copyout (void *h, size_t s)
 {
-  delete_copyout (DC_Copyout, h, s);
+  delete_copyout (FLAG_COPYOUT, h, s);
 }
 
 static void
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index b6ee7c1..3bbf8de 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -31,7 +31,6 @@
 #include "libgomp_g.h"
 #include "gomp-constants.h"
 #include "oacc-int.h"
-#include <stdio.h>
 #include <string.h>
 #include <stdarg.h>
 #include <assert.h>
diff --git libgomp/oacc-plugin.c libgomp/oacc-plugin.c
index 1fd6b2d..44cd6d6 100644
--- libgomp/oacc-plugin.c
+++ libgomp/oacc-plugin.c
@@ -34,7 +34,7 @@ void
 GOMP_PLUGIN_async_unmap_vars (void *ptr)
 {
   struct target_mem_desc *tgt = ptr;
-  
+
   gomp_unmap_vars (tgt, false);
 }
 
diff --git libgomp/oacc-plugin.h libgomp/oacc-plugin.h
index 33cec79..c60eb9c 100644
--- libgomp/oacc-plugin.h
+++ libgomp/oacc-plugin.h
@@ -24,10 +24,10 @@
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-#ifndef _OACC_PLUGIN_H
-#define _OACC_PLUGIN_H 1
+#ifndef OACC_PLUGIN_H
+#define OACC_PLUGIN_H 1
 
-extern void GOMP_PLUGIN_async_unmap_vars (void *ptr);
+extern void GOMP_PLUGIN_async_unmap_vars (void *);
 extern void *GOMP_PLUGIN_acc_thread (void);
 
 #endif
diff --git libgomp/openacc.h libgomp/openacc.h
index daf7c95..3343241 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -29,7 +29,7 @@
 #ifndef _OPENACC_H
 #define _OPENACC_H 1
 
-/* The OpenACC std is silent on whether or not including openacc.h
+/* The OpenACC standard is silent on whether or not including <openacc.h>
    might or must not include other header files.  We chose to include
    some.  */
 #include <stddef.h>
@@ -46,80 +46,71 @@ extern "C" {
 # define __GOACC_NOTHROW __attribute__ ((__nothrow__))
 #endif
 
-  /* Types */
-  typedef enum acc_device_t
-    {
-      /* Keep in sync with include/gomp-constants.h.  */
-      acc_device_none = 0,
-      acc_device_default = 1,
-      acc_device_host = 2,
-      acc_device_host_nonshm = 3,
-      acc_device_not_host = 4,
-      acc_device_nvidia = 5,
-      _ACC_device_hwm
-    } acc_device_t;
+/* Types */
+typedef enum acc_device_t
+  {
+    /* Keep in sync with include/gomp-constants.h.  */
+    acc_device_none = 0,
+    acc_device_default = 1,
+    acc_device_host = 2,
+    acc_device_host_nonshm = 3,
+    acc_device_not_host = 4,
+    acc_device_nvidia = 5,
+    _ACC_device_hwm
+  } acc_device_t;
 
-  typedef enum acc_async_t
-    {
-      /* Keep in sync with include/gomp-constants.h.  */
-      acc_async_noval = -1,
-      acc_async_sync  = -2
-    } acc_async_t;
+typedef enum acc_async_t
+  {
+    /* Keep in sync with include/gomp-constants.h.  */
+    acc_async_noval = -1,
+    acc_async_sync  = -2
+  } acc_async_t;
 
-  int acc_get_num_devices (acc_device_t __dev) __GOACC_NOTHROW;
-  void acc_set_device_type (acc_device_t __dev) __GOACC_NOTHROW;
-  acc_device_t acc_get_device_type (void) __GOACC_NOTHROW;
-  void acc_set_device_num (int __num, acc_device_t __dev) __GOACC_NOTHROW;
-  int acc_get_device_num (acc_device_t __dev) __GOACC_NOTHROW;
-  int acc_async_test (int __async) __GOACC_NOTHROW;
-  int acc_async_test_all (void) __GOACC_NOTHROW;
-  void acc_wait (int __async) __GOACC_NOTHROW;
-  void acc_wait_async (int __async1, int __async2) __GOACC_NOTHROW;
-  void acc_wait_all (void) __GOACC_NOTHROW;
-  void acc_wait_all_async (int __async) __GOACC_NOTHROW;
-  void acc_init (acc_device_t __dev) __GOACC_NOTHROW;
-  void acc_shutdown (acc_device_t __dev) __GOACC_NOTHROW;
-  int acc_on_device (acc_device_t __dev) __GOACC_NOTHROW;
-  void *acc_malloc (size_t) __GOACC_NOTHROW;
-  void acc_free (void *) __GOACC_NOTHROW;
-  /* Some of these would be more correct with const qualifiers, but
-     the standard specifies otherwise.  */
-  void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
-  void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
-  void *acc_create (void *, size_t) __GOACC_NOTHROW;
-  void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
-  void acc_copyout (void *, size_t) __GOACC_NOTHROW;
-  void acc_delete (void *, size_t) __GOACC_NOTHROW;
-  void acc_update_device (void *, size_t) __GOACC_NOTHROW;
-  void acc_update_self (void *, size_t) __GOACC_NOTHROW;
-  void acc_map_data (void *, void *, size_t) __GOACC_NOTHROW;
-  void acc_unmap_data (void *) __GOACC_NOTHROW;
-  void *acc_deviceptr (void *) __GOACC_NOTHROW;
-  void *acc_hostptr (void *) __GOACC_NOTHROW;
-  int acc_is_present (void *, size_t) __GOACC_NOTHROW;
-  void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
-  void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+int acc_get_num_devices (acc_device_t) __GOACC_NOTHROW;
+void acc_set_device_type (acc_device_t) __GOACC_NOTHROW;
+acc_device_t acc_get_device_type (void) __GOACC_NOTHROW;
+void acc_set_device_num (int, acc_device_t) __GOACC_NOTHROW;
+int acc_get_device_num (acc_device_t) __GOACC_NOTHROW;
+int acc_async_test (int) __GOACC_NOTHROW;
+int acc_async_test_all (void) __GOACC_NOTHROW;
+void acc_wait (int) __GOACC_NOTHROW;
+void acc_wait_async (int, int) __GOACC_NOTHROW;
+void acc_wait_all (void) __GOACC_NOTHROW;
+void acc_wait_all_async (int) __GOACC_NOTHROW;
+void acc_init (acc_device_t) __GOACC_NOTHROW;
+void acc_shutdown (acc_device_t) __GOACC_NOTHROW;
+int acc_on_device (acc_device_t) __GOACC_NOTHROW;
+void *acc_malloc (size_t) __GOACC_NOTHROW;
+void acc_free (void *) __GOACC_NOTHROW;
+/* Some of these would be more correct with const qualifiers, but
+   the standard specifies otherwise.  */
+void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void acc_copyout (void *, size_t) __GOACC_NOTHROW;
+void acc_delete (void *, size_t) __GOACC_NOTHROW;
+void acc_update_device (void *, size_t) __GOACC_NOTHROW;
+void acc_update_self (void *, size_t) __GOACC_NOTHROW;
+void acc_map_data (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_unmap_data (void *) __GOACC_NOTHROW;
+void *acc_deviceptr (void *) __GOACC_NOTHROW;
+void *acc_hostptr (void *) __GOACC_NOTHROW;
+int acc_is_present (void *, size_t) __GOACC_NOTHROW;
+void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
 
-  void ACC_target (int, void (*) (void *), const void *,
-	     size_t, void **, size_t *, unsigned char *, int *) __GOACC_NOTHROW;
-  void ACC_parallel (int, void (*) (void *), const void *,
-	     size_t, void **, size_t *, unsigned char *) __GOACC_NOTHROW;
-  void ACC_add_device_code (void const *, char const *) __GOACC_NOTHROW;
+/* Old names.  OpenACC does not specify whether these can or must
+   not be macros, inlines or aliases for the new names.  */
+#define acc_pcreate acc_present_or_create
+#define acc_pcopyin acc_present_or_copyin
 
-  void ACC_async_copy (int) __GOACC_NOTHROW;
-  void ACC_async_kern (int) __GOACC_NOTHROW;
+/* CUDA-specific routines.  */
+void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
+void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
+void *acc_get_cuda_stream (int) __GOACC_NOTHROW;
+int acc_set_cuda_stream (int, void *) __GOACC_NOTHROW;
 
-  /* Old names.  OpenACC does not specify whether these can or must
-     not be macros, inlines or aliases for the new names.  */
-  #define acc_pcreate acc_present_or_create
-  #define acc_pcopyin acc_present_or_copyin
-
-  /* CUDA-specific routines.  */
-  void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
-  void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
-  void *acc_get_cuda_stream (int __async) __GOACC_NOTHROW;
-  int acc_set_cuda_stream (int __async, void *__stream) __GOACC_NOTHROW;
-  
 #ifdef __cplusplus
 }
 #endif
diff --git libgomp/plugin/plugin-host.c libgomp/plugin/plugin-host.c
index 7437407..ebf7f11 100644
--- libgomp/plugin/plugin-host.c
+++ libgomp/plugin/plugin-host.c
@@ -55,10 +55,6 @@
 #define SELF "host: "
 #endif
 
-#ifndef HOST_NONSHM_PLUGIN
-static struct gomp_device_descr host_dispatch;
-#endif
-
 STATIC const char *
 GOMP_OFFLOAD_get_name (void)
 {
@@ -99,24 +95,24 @@ GOMP_OFFLOAD_get_num_devices (void)
 }
 
 STATIC void
-GOMP_OFFLOAD_register_image (void *host_table __attribute__((unused)),
-			     void *target_data __attribute__((unused)))
+GOMP_OFFLOAD_register_image (void *host_table __attribute__ ((unused)),
+			     void *target_data __attribute__ ((unused)))
 {
 }
 
 STATIC void
-GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_init_device (int n __attribute__ ((unused)))
 {
 }
 
 STATIC void
-GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_fini_device (int n __attribute__ ((unused)))
 {
 }
 
 STATIC int
-GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
-			struct mapping_table **table __attribute__((unused)))
+GOMP_OFFLOAD_get_table (int n __attribute__ ((unused)),
+			struct mapping_table **table __attribute__ ((unused)))
 {
   return 0;
 }
@@ -143,23 +139,23 @@ STATIC void
 GOMP_OFFLOAD_openacc_set_device_num (int n)
 {
   if (n > 0)
-    GOMP(fatal) ("device number %u out of range for host execution", n);
+    GOMP (fatal) ("device number %u out of range for host execution", n);
 }
 
 STATIC void *
-GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t s)
+GOMP_OFFLOAD_alloc (int n __attribute__ ((unused)), size_t s)
 {
-  return GOMP(malloc) (s);
+  return GOMP (malloc) (s);
 }
 
 STATIC void
-GOMP_OFFLOAD_free (int n __attribute__((unused)), void *p)
+GOMP_OFFLOAD_free (int n __attribute__ ((unused)), void *p)
 {
   free (p);
 }
 
 STATIC void *
-GOMP_OFFLOAD_host2dev (int n __attribute__((unused)), void *d, const void *h,
+GOMP_OFFLOAD_host2dev (int n __attribute__ ((unused)), void *d, const void *h,
 		       size_t s)
 {
 #ifdef HOST_NONSHM_PLUGIN
@@ -170,7 +166,7 @@ GOMP_OFFLOAD_host2dev (int n __attribute__((unused)), void *d, const void *h,
 }
 
 STATIC void *
-GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
+GOMP_OFFLOAD_dev2host (int n __attribute__ ((unused)), void *h, const void *d,
 		       size_t s)
 {
 #ifdef HOST_NONSHM_PLUGIN
@@ -181,7 +177,7 @@ GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
 }
 
 STATIC void
-GOMP_OFFLOAD_run (int n __attribute__((unused)), void *fn_ptr, void *vars)
+GOMP_OFFLOAD_run (int n __attribute__ ((unused)), void *fn_ptr, void *vars)
 {
   void (*fn)(void *) = (void (*)(void *)) fn_ptr;
 
@@ -190,16 +186,16 @@ GOMP_OFFLOAD_run (int n __attribute__((unused)), void *fn_ptr, void *vars)
 
 STATIC void
 GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *),
-			       size_t mapnum __attribute__((unused)),
-			       void **hostaddrs __attribute__((unused)),
-			       void **devaddrs __attribute__((unused)),
-			       size_t *sizes __attribute__((unused)),
-			       unsigned short *kinds __attribute__((unused)),
-			       int num_gangs __attribute__((unused)),
-			       int num_workers __attribute__((unused)),
-			       int vector_length __attribute__((unused)),
-			       int async __attribute__((unused)),
-			       void *targ_mem_desc __attribute__((unused)))
+			       size_t mapnum __attribute__ ((unused)),
+			       void **hostaddrs __attribute__ ((unused)),
+			       void **devaddrs __attribute__ ((unused)),
+			       size_t *sizes __attribute__ ((unused)),
+			       unsigned short *kinds __attribute__ ((unused)),
+			       int num_gangs __attribute__ ((unused)),
+			       int num_workers __attribute__ ((unused)),
+			       int vector_length __attribute__ ((unused)),
+			       int async __attribute__ ((unused)),
+			       void *targ_mem_desc __attribute__ ((unused)))
 {
 #ifdef HOST_NONSHM_PLUGIN
   fn (devaddrs);
@@ -219,12 +215,12 @@ GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
 }
 
 STATIC void
-GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__ ((unused)))
 {
 }
 
 STATIC int
-GOMP_OFFLOAD_openacc_async_test (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_test (int async __attribute__ ((unused)))
 {
   return 1;
 }
@@ -236,7 +232,7 @@ GOMP_OFFLOAD_openacc_async_test_all (void)
 }
 
 STATIC void
-GOMP_OFFLOAD_openacc_async_wait (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait (int async __attribute__ ((unused)))
 {
 }
 
@@ -246,25 +242,25 @@ GOMP_OFFLOAD_openacc_async_wait_all (void)
 }
 
 STATIC void
-GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__((unused)),
-				       int async2 __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__ ((unused)),
+				       int async2 __attribute__ ((unused)))
 {
 }
 
 STATIC void
-GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__ ((unused)))
 {
 }
 
 STATIC void *
 GOMP_OFFLOAD_openacc_create_thread_data (void *targ_data
-					 __attribute__((unused)))
+					 __attribute__ ((unused)))
 {
   return NULL;
 }
 
 STATIC void
 GOMP_OFFLOAD_openacc_destroy_thread_data (void *tls_data
-					  __attribute__((unused)))
+					  __attribute__ ((unused)))
 {
 }
diff --git libgomp/plugin/plugin-nvptx.c libgomp/plugin/plugin-nvptx.c
index ee0c818..483cb75 100644
--- libgomp/plugin/plugin-nvptx.c
+++ libgomp/plugin/plugin-nvptx.c
@@ -738,7 +738,7 @@ link_ptx (CUmodule *module, char *ptx_code)
   CUlinkState linkstate;
   CUresult r;
   void *linkout;
-  size_t linkoutsize __attribute__((unused));
+  size_t linkoutsize __attribute__ ((unused));
 
   GOMP_PLUGIN_debug (0, "attempting to load:\n---\n%s\n---\n", ptx_code);
 
@@ -1518,19 +1518,19 @@ GOMP_OFFLOAD_register_image (void *host_table, void *target_data)
 }
 
 void
-GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_init_device (int n __attribute__ ((unused)))
 {
   (void) nvptx_init ();
 }
 
 void
-GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_fini_device (int n __attribute__ ((unused)))
 {
   nvptx_fini ();
 }
 
 int
-GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
+GOMP_OFFLOAD_get_table (int n __attribute__ ((unused)),
 			struct mapping_table **tablep)
 {
   CUmodule module;
@@ -1589,26 +1589,26 @@ GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
 }
 
 void *
-GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t size)
+GOMP_OFFLOAD_alloc (int n __attribute__ ((unused)), size_t size)
 {
   return nvptx_alloc (size);
 }
 
 void
-GOMP_OFFLOAD_free (int n __attribute__((unused)), void *ptr)
+GOMP_OFFLOAD_free (int n __attribute__ ((unused)), void *ptr)
 {
   nvptx_free (ptr);
 }
 
 void *
-GOMP_OFFLOAD_dev2host (int ord __attribute__((unused)), void *dst,
+GOMP_OFFLOAD_dev2host (int ord __attribute__ ((unused)), void *dst,
 		       const void *src, size_t n)
 {
   return nvptx_dev2host (dst, src, n);
 }
 
 void *
-GOMP_OFFLOAD_host2dev (int ord __attribute__((unused)), void *dst,
+GOMP_OFFLOAD_host2dev (int ord __attribute__ ((unused)), void *dst,
 		       const void *src, size_t n)
 {
   return nvptx_host2dev (dst, src, n);
diff --git libgomp/splay-tree.c libgomp/splay-tree.c
index 3e4a904..030ca8f 100644
--- libgomp/splay-tree.c
+++ libgomp/splay-tree.c
@@ -26,14 +26,7 @@
 
 /* The splay tree code copied from include/splay-tree.h and adjusted,
    so that all the data lives directly in splay_tree_node_s structure
-   and no extra allocations are needed.
-
-   Files including this header should before including it add:
-typedef struct splay_tree_node_s *splay_tree_node;
-typedef struct splay_tree_s *splay_tree;
-typedef struct splay_tree_key_s *splay_tree_key;
-   define splay_tree_key_s structure, and define
-   splay_compare inline function.  */
+   and no extra allocations are needed.  */
 
 /* For an easily readable description of splay-trees, see:
 
diff --git libgomp/splay-tree.h libgomp/splay-tree.h
index d04e13c..085021c 100644
--- libgomp/splay-tree.h
+++ libgomp/splay-tree.h
@@ -46,27 +46,6 @@ typedef struct splay_tree_key_s *splay_tree_key;
 #ifndef _SPLAY_TREE_H
 #define _SPLAY_TREE_H 1
 
-typedef struct splay_tree_node_s *splay_tree_node;
-typedef struct splay_tree_s *splay_tree;
-typedef struct splay_tree_key_s *splay_tree_key;
-
-struct splay_tree_key_s {
-  /* Address of the host object.  */
-  uintptr_t host_start;
-  /* Address immediately after the host object.  */
-  uintptr_t host_end;
-  /* Descriptor of the target memory.  */
-  struct target_mem_desc *tgt;
-  /* Offset from tgt->tgt_start to the start of the target object.  */
-  uintptr_t tgt_offset;
-  /* Reference count.  */
-  uintptr_t refcount;
-  /* Asynchronous reference count.  */
-  uintptr_t async_refcount;
-  /* True if data should be copied from device to host at the end.  */
-  bool copy_from;
-};
-
 /* The nodes in the splay tree.  */
 struct splay_tree_node_s {
   struct splay_tree_key_s key;
diff --git libgomp/target.c libgomp/target.c
index 6871e7b..83ad511 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -34,7 +34,6 @@
 #include <stdbool.h>
 #include <stdlib.h>
 #include <string.h>
-#include <stdio.h>
 #include <assert.h>
 
 #ifdef PLUGIN_SUPPORT
@@ -168,6 +167,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
       tgt_align = align;
       tgt_size = mapnum * sizeof (void *);
     }
+
   gomp_mutex_lock (&mm->lock);
 
   for (i = 0; i < mapnum; i++)
@@ -288,9 +288,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		tgt->refcount++;
 		array->left = NULL;
 		array->right = NULL;
-
 		splay_tree_insert (&mm->splay_tree, array);
-
 		switch (kind & typemask)
 		  {
 		  case GOMP_MAP_ALLOC:
@@ -302,7 +300,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		  case GOMP_MAP_TOFROM:
 		  case GOMP_MAP_FORCE_TO:
 		  case GOMP_MAP_FORCE_TOFROM:
-		    /* Copy from host to device memory.  */
 		    /* FIXME: Perhaps add some smarts, like if copying
 		       several adjacent fields from host to target, use some
 		       host buffer to avoid sending each var individually.  */
@@ -318,7 +315,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		    if (cur_node.host_start == (uintptr_t) NULL)
 		      {
 			cur_node.tgt_offset = (uintptr_t) NULL;
-			/* Copy from host to device memory.  */
 			/* FIXME: see above FIXME comment.  */
 			devicep->host2dev_func (devicep->target_id,
 						(void *) (tgt->tgt_start
@@ -346,7 +342,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		    if (n == NULL)
 		      gomp_fatal ("Pointer target of array section "
 				  "wasn't mapped");
-
 		    cur_node.host_start -= n->host_start;
 		    cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
 					  + cur_node.host_start;
@@ -354,7 +349,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		       array section.  Now subtract bias to get what we want
 		       to initialize the pointer with.  */
 		    cur_node.tgt_offset -= sizes[i];
-		    /* Copy from host to device memory.  */
 		    /* FIXME: see above FIXME comment.  */
 		    devicep->host2dev_func (devicep->target_id,
 					    (void *) (tgt->tgt_start
@@ -363,7 +357,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 					    sizeof (void *));
 		    break;
 		  case GOMP_MAP_TO_PSET:
-		    /* Copy from host to device memory.  */
 		    /* FIXME: see above FIXME comment.  */
 		    devicep->host2dev_func (devicep->target_id,
 					    (void *) (tgt->tgt_start
@@ -388,7 +381,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 			  if (cur_node.host_start == (uintptr_t) NULL)
 			    {
 			      cur_node.tgt_offset = (uintptr_t) NULL;
-			      /* Copy from host to device memory.  */
 			      /* FIXME: see above FIXME comment.  */
 			      devicep->host2dev_func (devicep->target_id,
 				 (void *) (tgt->tgt_start + k->tgt_offset
@@ -428,7 +420,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 			     array section.  Now subtract bias to get what we
 			     want to initialize the pointer with.  */
 			  cur_node.tgt_offset -= sizes[j];
-			  /* Copy from host to device memory.  */
 			  /* FIXME: see above FIXME comment.  */
 			  devicep->host2dev_func (devicep->target_id,
 			     (void *) (tgt->tgt_start + k->tgt_offset
@@ -476,7 +467,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	  else
 	    cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start
 				  + tgt->list[i]->tgt_offset;
-	  /* Copy from host to device memory.  */
 	  /* FIXME: see above FIXME comment.  */
 	  devicep->host2dev_func (devicep->target_id,
 				  (void *) (tgt->tgt_start
@@ -528,7 +518,6 @@ gomp_copy_from_async (struct target_mem_desc *tgt)
       {
 	splay_tree_key k = tgt->list[i];
 	if (k->copy_from)
-	  /* Copy from device to host memory.  */
 	  devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
 				  (void *) (k->tgt->tgt_start + k->tgt_offset),
 				  k->host_end - k->host_start);
@@ -567,7 +556,6 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
       {
 	splay_tree_key k = tgt->list[i];
 	if (k->copy_from && do_copyfrom)
-	  /* Copy from device to host memory.  */
 	  devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
 				  (void *) (k->tgt->tgt_start + k->tgt_offset),
 				  k->host_end - k->host_start);
@@ -621,7 +609,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
 			  (void *) n->host_start,
 			  (void *) n->host_end);
 	    if (GOMP_MAP_COPY_TO_P (kind & typemask))
-	      /* Copy from host to device memory.  */
 	      devicep->host2dev_func (devicep->target_id,
 				      (void *) (n->tgt->tgt_start
 						+ n->tgt_offset
@@ -630,7 +617,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
 				      (void *) cur_node.host_start,
 				      cur_node.host_end - cur_node.host_start);
 	    if (GOMP_MAP_COPY_FROM_P (kind & typemask))
-	      /* Copy from device to host memory.  */
 	      devicep->dev2host_func (devicep->target_id,
 				      (void *) cur_node.host_start,
 				      (void *) (n->tgt->tgt_start
@@ -647,9 +633,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
   gomp_mutex_unlock (&mm->lock);
 }
 
-static void gomp_register_image_for_device (struct gomp_device_descr *device,
-					    struct offload_image_descr *image);
-
 /* This function should be called from every offload image.
    It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
    the target, and TARGET_DATA needed by target plugin.  */
@@ -662,9 +645,6 @@ GOMP_offload_register (void *host_table, enum offload_target_type target_type,
 				 (num_offload_images + 1)
 				 * sizeof (struct offload_image_descr));
 
-  if (offload_images == NULL)
-    return;
-
   offload_images[num_offload_images].type = target_type;
   offload_images[num_offload_images].host_table = host_table;
   offload_images[num_offload_images].target_data = target_data;
@@ -691,9 +671,10 @@ gomp_init_tables (struct gomp_device_descr *devicep,
 {
   /* Get address mapping table for device.  */
   struct mapping_table *table = NULL;
-  int i, num_entries = devicep->get_table_func (devicep->target_id, &table);
+  int num_entries = devicep->get_table_func (devicep->target_id, &table);
 
   /* Insert host-target address mapping into dev_splay_tree.  */
+  int i;
   for (i = 0; i < num_entries; i++)
     {
       struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
@@ -1116,12 +1097,13 @@ gomp_target_init (void)
 		    break;
 		  }
 
-		current_device.type = current_device.get_type_func ();
 		current_device.name = current_device.get_name_func ();
+		/* current_device.capabilities has already been set.  */
+		current_device.type = current_device.get_type_func ();
+		current_device.mem_map.is_initialized = false;
+		current_device.mem_map.splay_tree.root = NULL;
 		current_device.is_initialized = false;
 		current_device.offload_regions_registered = false;
-		current_device.mem_map.splay_tree.root = NULL;
-		current_device.mem_map.is_initialized = false;
 		current_device.openacc.data_environ = NULL;
 		current_device.openacc.target_data = NULL;
 		for (i = 0; i < new_num_devices; i++)
diff --git libgomp/testsuite/lib/libgomp.exp libgomp/testsuite/lib/libgomp.exp
index ca62897..dded36a 100644
--- libgomp/testsuite/lib/libgomp.exp
+++ libgomp/testsuite/lib/libgomp.exp
@@ -36,12 +36,17 @@ load_gcc_lib fortran-modules.exp
 load_file libgomp-test-support.exp
 
 # Populate offload_targets_s (offloading targets separated by a space), and
-# offload_targets_s_openacc (the same, but with OpenACC names).
+# offload_targets_s_openacc (the same, but with OpenACC names; OpenACC spells
+# some of them a little differently).
 set offload_targets_s [split $offload_targets ","]
-# OpenACC spells some of them a little differently.
 set offload_targets_s_openacc {}
 foreach offload_target_openacc $offload_targets_s {
     switch $offload_target_openacc {
+	intelmic {
+	    # TODO.  Skip; will all FAIL because of missing
+	    # GOMP_OFFLOAD_CAP_OPENACC_200.
+	    continue
+	}
 	nvptx {
 	    set offload_target_openacc "nvidia"
 	}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/abort.c libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/abort.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
index 80c51d5..747109f 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -140,7 +140,7 @@ main (int argc, char **argv)
     e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 
 #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
-#pragma acc delete (N)
+#pragma acc exit data delete (N)
 #pragma acc wait (1)
 
 
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 2a77936..184b355 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -538,7 +538,76 @@ main(int argc, char **argv)
     {
         if (b[i] != 18.0)
             abort ();
-	}
+    }
+
+#pragma acc enter data copyin (b[0:N]) if (0)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (0)
+
+#pragma acc enter data copyin (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (zero)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (zero)
+
+#pragma acc enter data copyin (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (one == 0)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 0)
+
+#pragma acc enter data copyin (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
 
     return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index 0579185..c164598 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
@@ -6,30 +6,136 @@ int
 main (int argc, char *argv[])
 {
 #define N 10
-  char a[N];
-
-  {
+    char a[N];
     int i;
+
     for (i = 0; i < N; ++i)
-      a[i] = 0;
-  }
+       a[i] = 0;
 
-#pragma acc data copyout (a)
-  {
-#pragma acc parallel /* will result in a "dummy frame" */ present (a)
+#pragma acc data copy (a)
     {
-      int i;
-      for (i = 0; i < N; ++i)
-	a[i] = i;
+#pragma acc parallel present (a)
+        {
+            int j;
+
+            for (j = 0; j < N; ++j)
+                a[j] = j;
+        }
     }
-  }
 
-  {
-    int i;
     for (i = 0; i < N; ++i)
-      if (a[i] != i)
-	abort ();
-  }
+    {
+        if (a[i] != i)
+            abort ();
+    }
+
+    for (i = 0; i < N; ++i)
+       a[i] = 0;
+
+#pragma acc data copy (a)
+    {
+#pragma acc kernels present (a)
+        {
+            int j;
+
+            for (j = 0; j < N; ++j)
+                a[j] = j;
+        }
+    }
+
+    for (i = 0; i < N; ++i)
+    {
+        if (a[i] != i)
+            abort ();
+    }
+
+    for (i = 0; i < N; ++i)
+        a[i] = 0;
+
+#pragma acc data copy (a)
+    {
+#pragma acc data present (a)
+        {
+#pragma acc parallel present (a)
+            {
+                int j;
+
+                for (j = 0; j < N; ++j)
+                    a[j] = j;
+            }
+        }
+    }
+
+    for (i = 0; i < N; ++i)
+    {
+        if (a[i] != i)
+            abort ();
+    }
+
+#pragma acc data copy (a)
+    {
+#pragma acc data present (a)
+        {
+#pragma acc kernels present (a)
+            {
+                int j;
+
+                for (j = 0; j < N; ++j)
+                    a[j] = j;
+            }
+        }
+    }
+
+    for (i = 0; i < N; ++i)
+    {
+        if (a[i] != i)
+            abort ();
+    }
+
+    for (i = 0; i < N; ++i)
+        a[i] = 0;
+
+#pragma acc enter data copyin (a)
+
+#pragma acc data present (a)
+    {
+#pragma acc parallel present (a)
+        {
+            int j;
+
+            for (j = 0; j < N; ++j)
+                a[j] = j;
+        }
+    }
+
+#pragma acc exit data copyout (a)
+
+    for (i = 0; i < N; ++i)
+    {
+        if (a[i] != i)
+            abort ();
+    }
+
+#pragma acc enter data copyin (a)
+
+#pragma acc data present (a)
+    {
+#pragma acc kernels present (a)
+        {
+            int j;
+
+            for (j = 0; j < N; ++j)
+                a[j] = j;
+        }
+    }
+
+#pragma acc exit data copyout (a)
+
+    for (i = 0; i < N; ++i)
+    {
+        if (a[i] != i)
+            abort ();
+    }
 
-  return 0;
+    return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90 libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90
deleted file mode 100644
index a54d6a7..0000000
--- libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90
+++ /dev/null
@@ -1,82 +0,0 @@
-! { dg-do run }
-
-program main
-  implicit none
-  include "openacc_lib.h"
-
-  integer, target :: a_3d_i(10, 10, 10)
-  complex a_3d_c(10, 10, 10)
-  real a_3d_r(10, 10, 10)
-
-  integer i, j, k
-  complex c
-  real r
-  integer, parameter :: i_size = sizeof (i)
-  integer, parameter :: c_size = sizeof (c)
-  integer, parameter :: r_size = sizeof (r)
-
-  if (acc_get_num_devices (acc_device_nvidia) .eq. 0) call exit
-
-  call acc_init (acc_device_nvidia)
-
-  call set3d (.FALSE., a_3d_i, a_3d_c, a_3d_r)
-
-  call acc_copyin (a_3d_i)
-  call acc_copyin (a_3d_c)
-  call acc_copyin (a_3d_r)
-
-  if (acc_is_present (a_3d_i) .neqv. .TRUE.) call abort
-  if (acc_is_present (a_3d_c) .neqv. .TRUE.) call abort
-  if (acc_is_present (a_3d_r) .neqv. .TRUE.) call abort
-
-  do i = 1, 10
-    do j = 1, 10
-      do k = 1, 10
-        if (acc_is_present (a_3d_i(i, j, k), i_size) .neqv. .TRUE.) call abort
-        if (acc_is_present (a_3d_c(i, j, k), i_size) .neqv. .TRUE.) call abort
-        if (acc_is_present (a_3d_r(i, j, k), i_size) .neqv. .TRUE.) call abort
-      end do
-    end do
-  end do
-
-  call acc_shutdown (acc_device_nvidia)
-
-contains
-
-  subroutine set3d (clear, a_i, a_c, a_r)
-  logical clear
-  integer, dimension (:,:,:), intent (inout) :: a_i
-  complex, dimension (:,:,:), intent (inout) :: a_c
-  real, dimension (:,:,:), intent (inout) :: a_r
-
-  integer i, j, k
-  integer lb1, ub1, lb2, ub2, lb3, ub3
-
-  lb1 = lbound (a_i, 1)
-  ub1 = ubound (a_i, 1)
-
-  lb2 = lbound (a_i, 2)
-  ub2 = ubound (a_i, 2)
-
-  lb3 = lbound (a_i, 3)
-  ub3 = ubound (a_i, 3)
-
-  do i = lb1, ub1
-    do j = lb2, ub2
-      do k = lb3, ub3
-        if (clear) then
-          a_i(i, j, k) = 0
-          a_c(i, j, k) = cmplx (0.0, 0.0)
-          a_r(i, j, k) = 0.0
-        else
-          a_i(i, j, k) = i
-          a_c(i, j, k) = cmplx (i, j)
-          a_r(i, j, k) = i
-        end if
-      end do
-    end do
-  end do
-
-  end subroutine
-
-end program
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90 libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90
deleted file mode 100644
index ad758b2..0000000
--- libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90
+++ /dev/null
@@ -1,83 +0,0 @@
-! { dg-do run }
-
-program main
-  use openacc
-  use iso_c_binding
-  implicit none
-
-  integer, target :: a_3d_i(10, 10, 10)
-  complex a_3d_c(10, 10, 10)
-  real a_3d_r(10, 10, 10)
-
-  integer i, j, k
-  complex c
-  real r
-  integer, parameter :: i_size = sizeof (i)
-  integer, parameter :: c_size = sizeof (c)
-  integer, parameter :: r_size = sizeof (r)
-
-  if (acc_get_num_devices (acc_device_nvidia) .eq. 0) call exit
-
-  call acc_init (acc_device_nvidia)
-
-  call set3d (.FALSE., a_3d_i, a_3d_c, a_3d_r)
-
-  call acc_copyin (a_3d_i)
-  call acc_copyin (a_3d_c)
-  call acc_copyin (a_3d_r)
-
-  if (acc_is_present (a_3d_i) .neqv. .TRUE.) call abort
-  if (acc_is_present (a_3d_c) .neqv. .TRUE.) call abort
-  if (acc_is_present (a_3d_r) .neqv. .TRUE.) call abort
-
-  do i = 1, 10
-    do j = 1, 10
-      do k = 1, 10
-        if (acc_is_present (a_3d_i(i, j, k), i_size) .neqv. .TRUE.) call abort
-        if (acc_is_present (a_3d_c(i, j, k), i_size) .neqv. .TRUE.) call abort
-        if (acc_is_present (a_3d_r(i, j, k), i_size) .neqv. .TRUE.) call abort
-      end do
-    end do
-  end do
-
-  call acc_shutdown (acc_device_nvidia)
-
-contains
-
-  subroutine set3d (clear, a_i, a_c, a_r)
-  logical clear
-  integer, dimension (:,:,:), intent (inout) :: a_i
-  complex, dimension (:,:,:), intent (inout) :: a_c
-  real, dimension (:,:,:), intent (inout) :: a_r
-
-  integer i, j, k
-  integer lb1, ub1, lb2, ub2, lb3, ub3
-
-  lb1 = lbound (a_i, 1)
-  ub1 = ubound (a_i, 1)
-
-  lb2 = lbound (a_i, 2)
-  ub2 = ubound (a_i, 2)
-
-  lb3 = lbound (a_i, 3)
-  ub3 = ubound (a_i, 3)
-
-  do i = lb1, ub1
-    do j = lb2, ub2
-      do k = lb3, ub3
-        if (clear) then
-          a_i(i, j, k) = 0
-          a_c(i, j, k) = cmplx (0.0, 0.0)
-          a_r(i, j, k) = 0.0
-        else
-          a_i(i, j, k) = i
-          a_c(i, j, k) = cmplx (i, j)
-          a_r(i, j, k) = i
-        end if
-      end do
-    end do
-  end do
-
-  end subroutine
-
-end program
diff --git liboffloadmic/ChangeLog liboffloadmic/ChangeLog
index 9faa452..074926e 100644
--- liboffloadmic/ChangeLog
+++ liboffloadmic/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_name)
+	(GOMP_OFFLOAD_get_caps, GOMP_OFFLOAD_fini_device): New functions.
+
 2014-11-13  Ilya Verbin  <ilya.verbin@intel.com>
 	    Andrey Turetskiy  <andrey.turetskiy@intel.com>
 
