gcc/

2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* builtin-types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR,
	BT_FN_BOOL_UINT_ULLPTR_ULLPTR_ULLPTR,
	BT_FN_BOOL_UINT_LONGPTR_LONG_LONGPTR_LONGPTR,
	BT_FN_BOOL_UINT_ULLPTR_ULL_ULLPTR_ULLPTR,
	BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
	BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_ULL_ULL_ULL,
	BT_FN_VOID_LONG_VAR, BT_FN_VOID_ULL_VAR): New.
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR): Remove.
	* cgraph.h (enum cgraph_simd_clone_arg_type): Add
	SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP,
	SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP and
	SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP.
	(struct cgraph_simd_clone_arg): Adjust comment.
	* coretypes.h (struct gomp_ordered): New forward decl.
	* gimple.c (gimple_build_omp_critical): Add CLAUSES argument,
	set critical clauses to it.
	(gimple_build_omp_ordered): Return gomp_ordered * instead of
	gimple *.  Add CLAUSES argument, set ordered clauses to it.
	(gimple_copy): Unshare clauses on GIMPLE_OMP_CRITICAL and
	GIMPLE_OMP_ORDERED.
	* gimple.def (GIMPLE_OMP_ORDERED): Change from GSS_OMP to
	GSS_OMP_SINGLE_LAYOUT, move it after GIMPLE_OMP_TEAMS.
	* gimple.h (enum gf_mask): Add GF_OMP_TASK_TASKLOOP.  Add another bit
	to GF_OMP_FOR_KIND_MASK mask. Add GF_OMP_FOR_KIND_TASKLOOP, renumber
	GF_OMP_FOR_KIND_CILKFOR and GF_OMP_FOR_KIND_OACC_LOOP.  Adjust
	GF_OMP_FOR_SIMD, GF_OMP_FOR_COMBINED and GF_OMP_FOR_COMBINED_INTO.
	Add another bit to GF_OMP_TARGET_KIND_MASK mask.  Add
	GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA,
	renumber
	GF_OMP_TARGET_KIND_OACC_{PARALLEL,KERNELS,DATA,UPDATE,ENTER_EXIT_DATA}.
	(gomp_critical): Add clauses field.
	(gomp_ordered): New struct.
	(is_a_helper <gomp_ordered *>::test): New inline.
	(gimple_build_omp_critical): Add CLAUSES argument.
	(gimple_build_omp_ordered): Likewise.  Return gomp_ordered *
	instead of gimple *.
	(gimple_omp_critical_clauses, gimple_omp_critical_clauses_ptr,
	gimple_omp_critical_set_clauses, gimple_omp_ordered_clauses,
	gimple_omp_ordered_clauses_ptr, gimple_omp_ordered_set_clauses,
	gimple_omp_task_taskloop_p, gimple_omp_task_set_taskloop_p): New
	inline functions.
	* gimple-pretty-print.c (dump_gimple_omp_for): Handle taskloop.
	(dump_gimple_omp_target): Handle enter data and exit data.
	(dump_gimple_omp_block): Don't handle GIMPLE_OMP_ORDERED here.
	(dump_gimple_omp_critical): Print clauses.
	(dump_gimple_omp_ordered): New function.
	(dump_gimple_omp_task): Handle taskloop.
	(pp_gimple_stmt_1): Use dump_gimple_omp_ordered for
	GIMPLE_OMP_ORDERED.
	* gimple-walk.c (walk_gimple_op): Walk clauses on
	GIMPLE_OMP_CRITICAL and GIMPLE_OMP_ORDERED.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_0LEN_ARRAY.
	(enum omp_region_type): Add ORT_COMBINED_TARGET and ORT_NONE.
	(struct gimplify_omp_ctx): Add loop_iter_var,
	target_map_scalars_firstprivate, target_map_pointers_as_0len_arrays
	and target_firstprivatize_array_bases fields.
	(delete_omp_context): Release loop_iter_var.
	(gimplify_bind_expr): Handle ORT_NONE.
	(maybe_fold_stmt): Adjust check for ORT_TARGET for the addition of
	ORT_COMBINED_TARGET.
	(is_gimple_stmt): Return true for OMP_TASKLOOP, OMP_TEAMS and
	OMP_TARGET{,_DATA,_UPDATE,_ENTER_DATA,_EXIT_DATA}.
	(omp_firstprivatize_variable): Handle ORT_NONE.  Adjust check for
	ORT_TARGET for the addition of ORT_COMBINED_TARGET.  Handle
	ctx->target_map_scalars_firstprivate.
	(omp_add_variable): Handle ORT_NONE.  Allow map clause together with
	data sharing clauses.  For data sharing clause with VLA decl
	on omp target/target data don't add firstprivate for the pointer.
	Call omp_notice_variable on TYPE_SIZE_UNIT only if it is a DECL_P.
	(omp_notice_threadprivate_variable): Adjust check for ORT_TARGET for
	the addition of ORT_COMBINED_TARGET.
	(omp_notice_variable): Handle ORT_NONE.  Adjust check for ORT_TARGET
	for the addition of ORT_COMBINED_TARGET.  Handle implicit mapping of
	pointers as zero length array sections and
	ctx->target_map_scalars_firstprivate mapping of scalars as firstprivate
	data sharing.
	(omp_check_private): Handle omp_member_access_dummy_var vars.
	(find_decl_expr): New function.
	(gimplify_scan_omp_clauses): Add CODE argument.  For OMP_CLAUSE_IF
	complain if OMP_CLAUSE_IF_MODIFIER is present and does not match code.
	Handle OMP_CLAUSE_GANG separately.  Handle
	OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD,SIMDLEN}
	clauses.  Diagnose linear clause on combined
	distribute {, parallel for} simd construct, unless it is the loop
	iterator.  Handle struct element GOMP_MAP_FIRSTPRIVATE_POINTER.
	Handle map clauses with COMPONENT_REF.  Initialize
	ctx->target_map_scalars_firstprivate,
	ctx->target_firstprivatize_array_bases and
	ctx->target_map_pointers_as_0len_arrays.  Add firstprivate for
	linear clause even to target region if combined.  Remove
	map clauses with GOMP_MAP_FIRSTPRIVATE_POINTER kind from
	OMP_TARGET_{,ENTER_,EXIT_}DATA.  For GOMP_MAP_FIRSTPRIVATE_POINTER
	map kind with non-INTEGER_CST OMP_CLAUSE_SIZE firstprivatize the bias.
	Handle OMP_CLAUSE_DEPEND_{SINK,SOURCE}.  Handle
	OMP_CLAUSE_{{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}.
	For linear clause on worksharing loop combined with parallel add
	shared clause on the parallel.  Handle OMP_CLAUSE_REDUCTION
	with MEM_REF OMP_CLAUSE_DECL.  Set DECL_NAME on
	omp_member_access_dummy_var vars.  Add lastprivate clause to outer
	taskloop if needed.
	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_0LEN_ARRAY.
	If gimplify_omp_ctxp->target_firstprivatize_array_bases, use
	GOMP_MAP_FIRSTPRIVATE_POINTER map kind instead of
	GOMP_MAP_POINTER.
	(gimplify_adjust_omp_clauses): Add CODE argument.  Handle removal
	of GOMP_MAP_FIRSTPRIVATE_POINTER struct elements for struct not seen
	in target body.  Handle removal of struct mapping if struct is not
	seen in target body.  Remove GOMP_MAP_STRUCT map clause on
	OMP_TARGET_EXIT_DATA.  Adjust check for ORT_TARGET for the
	addition of ORT_COMBINED_TARGET.  Use GOMP_MAP_FIRSTPRIVATE_POINTER
	instead of GOMP_MAP_POINTER if ctx->target_firstprivatize_array_bases
	for VLAs.  Set OMP_CLAUSE_MAP_PRIVATE if both data sharing and map
	clause appear together.  Handle
	OMP_CLAUSE_{{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}.  Don't remove map
	clause if it has map-type-modifier always.  Handle
	OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD,SIMDLEN}
	clauses.
	(gimplify_oacc_cache, gimplify_omp_parallel, gimplify_omp_task):
	Adjust gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses
	callers.
	(gimplify_omp_for): Likewise.  Handle OMP_TASKLOOP.  Initialize
	loop_iter_var.  Use OMP_FOR_ORIG_DECLS.  Fix handling of lastprivate
	iterators in doacross loops.
	(gimplify_omp_workshare): Adjust gimplify_scan_omp_clauses and
	gimplify_adjust_omp_clauses callers.  Use ORT_COMBINED_TARGET
	for OMP_TARGET_COMBINED.  Adjust check for ORT_TARGET
	for the addition of ORT_COMBINED_TARGET.
	(gimplify_omp_target_update): Adjust gimplify_scan_omp_clauses and
	gimplify_adjust_omp_clauses callers.  Handle OMP_TARGET_ENTER_DATA
	and OMP_TARGET_EXIT_DATA.
	(gimplify_omp_ordered): New function.
	(gimplify_expr): Handle OMP_TASKLOOP, OMP_TARGET_ENTER_DATA and
	OMP_TARGET_EXIT_DATA.  Use gimplify_omp_ordered for OMP_ORDERED.
	Gimplify clauses on OMP_CRITICAL.
	* internal-fn.c (expand_GOMP_SIMD_ORDERED_START,
	expand_GOMP_SIMD_ORDERED_END): New functions.
	* internal-fn.def (GOMP_SIMD_ORDERED_START,
	GOMP_SIMD_ORDERED_END): New internal functions.
	* omp-builtins.def (BUILT_IN_GOMP_LOOP_DOACROSS_STATIC_START,
	BUILT_IN_GOMP_LOOP_DOACROSS_DYNAMIC_START,
	BUILT_IN_GOMP_LOOP_DOACROSS_GUIDED_START,
	BUILT_IN_GOMP_LOOP_DOACROSS_RUNTIME_START,
	BUILT_IN_GOMP_LOOP_ULL_DOACROSS_STATIC_START,
	BUILT_IN_GOMP_LOOP_ULL_DOACROSS_DYNAMIC_START,
	BUILT_IN_GOMP_LOOP_ULL_DOACROSS_GUIDED_START,
	BUILT_IN_GOMP_LOOP_ULL_DOACROSS_RUNTIME_START,
	BUILT_IN_GOMP_DOACROSS_POST, BUILT_IN_GOMP_DOACROSS_WAIT,
	BUILT_IN_GOMP_DOACROSS_ULL_POST, BUILT_IN_GOMP_DOACROSS_ULL_WAIT,
	BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA, BUILT_IN_GOMP_TASKLOOP,
	BUILT_IN_GOMP_TASKLOOP_ULL): New built-ins.
	(BUILT_IN_GOMP_TASK): Add INT argument to the end.
	(BUILT_IN_GOMP_TARGET): Rename from GOMP_target to GOMP_target_41,
	adjust type.
	(BUILT_IN_GOMP_TARGET_DATA): Rename from GOMP_target_data to
	GOMP_target_data_41, adjust type.
	(BUILT_IN_GOMP_TARGET_UPDATE): Rename from GOMP_target_update to
	GOMP_target_update_41, adjust type.
	* omp-low.c (struct omp_region): Adjust comments, add ord_stmt
	field.
	(struct omp_for_data): Add ordered and simd_schedule fields.
	(omp_member_access_dummy_var, unshare_and_remap_1,
	unshare_and_remap, is_taskloop_ctx): New functions.
	(is_taskreg_ctx): Use is_parallel_ctx and is_task_ctx.
	(extract_omp_for_data): Handle taskloops and doacross loops
	and simd schedule modifier.
	(omp_adjust_chunk_size): New function.
	(get_ws_args_for): Use it.
	(lookup_sfield): Change first argument to splay_tree_key,
	add overload with first argument tree.
	(maybe_lookup_field): Likewise.
	(use_pointer_for_field): Handle omp_member_access_dummy_var.
	(omp_copy_decl_2): If var is TREE_ADDRESSABLE listed in
	task_shared_vars, clear TREE_ADDRESSABLE on the copy.
	(build_outer_var_ref): Add LASTPRIVATE argument, handle
	taskloops and omp_member_access_dummy_var vars.
	(build_sender_ref): Change first argument to splay_tree_key,
	add overload with first argument tree.
	(install_var_field): For mask & 8 use &DECL_UID as key instead
	of the tree itself.
	(fixup_child_record_type): Const qualify *.omp_data_i.
	(scan_sharing_clauses): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE,
	C/C++ array reductions, OMP_CLAUSE_{IS,USE}_DEVICE_PTR clauses,
	OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,SIMDLEN,THREADS,SIMD} and
	OMP_CLAUSE_{NOGROUP,DEFAULTMAP} clauses, OMP_CLAUSE__LOOPTEMP_ clause
	on taskloop, GOMP_MAP_FIRSTPRIVATE_POINTER, OMP_CLAUSE_MAP_PRIVATE.
	(create_omp_child_function): Set TREE_READONLY on .omp_data_i.
	(find_combined_for): Allow searching for different GIMPLE_OMP_FOR
	kinds.
	(add_taskreg_looptemp_clauses): New function.
	(scan_omp_parallel): Use it.
	(scan_omp_task): Likewise.
	(finish_taskreg_scan): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE.
	For taskloop, move fields for the first two _LOOPTEMP_ clauses first.
	(check_omp_nesting_restrictions): Handle GF_OMP_TARGET_KIND_ENTER_DATA
	and GF_OMP_TARGET_KIND_EXIT_DATA.  Formatting fixes.  Allow the
	sandwiched taskloop constructs.  Type check
	OMP_CLAUSE_DEPEND_{KIND,SOURCE}.  Allow ordered simd inside of simd
	region.  Diagnose depend(source) or depend(sink:...) on
	target constructs or task/taskloop.
	(handle_simd_reference): Use get_name.
	(lower_rec_input_clauses): Likewise.  Ignore all
	OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE clauses on taskloop construct.
	Allow _LOOPTEMP_ clause on GOMP_TASK.  Unshare new_var
	before passing it to omp_clause_{default,copy}_ctor.  Handle
	OMP_CLAUSE_REDUCTION with MEM_REF OMP_CLAUSE_DECL.  Set
	lastprivate_firstprivate flag for linear that needs copyin and
	copyout.  Use BUILT_IN_ALLOCA_WITH_ALIGN instead of BUILT_IN_ALLOCA.
	(lower_lastprivate_clauses): For OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE
	on taskloop lookup decl in outer context.  Pass true to
	build_outer_var_ref lastprivate argument.  Handle
	OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV lastprivate if the decl is global
	outside of outer taskloop for.
	(lower_reduction_clauses): Handle OMP_CLAUSE_REDUCTION with MEM_REF
	OMP_CLAUSE_DECL.
	(lower_send_clauses): Ignore first two _LOOPTEMP_ clauses in taskloop
	GOMP_TASK.  Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE.  Handle
	omp_member_access_dummy_var vars.  Handle OMP_CLAUSE_REDUCTION
	with MEM_REF OMP_CLAUSE_DECL.  Use new lookup_sfield overload.
	(lower_send_shared_vars): Ignore fields with NULL or FIELD_DECL
	abstract origin.  Handle omp_member_access_dummy_var vars.
	(expand_parallel_call): Use expand_omp_build_assign.
	(expand_task_call): Handle taskloop construct expansion.  Add
	REGION argument.  Use GOMP_TASK_* defines instead of hardcoded
	integers.  Add priority argument to GOMP_task* calls.  Or in
	GOMP_TASK_FLAG_PRIORITY into flags if priority is present for
	GOMP_task call.
	(expand_omp_build_assign): Add prototype.  Add AFTER
	argument, if true emit statements after *GSI_P and continue linking.
	(expand_omp_taskreg): Adjust expand_task_call caller.
	(expand_omp_for_init_counts): Rename zero_iter_bb argument to
	zero_iter1_bb and first_zero_iter to first_zero_iter1, add
	zero_iter2_bb and first_zero_iter2 arguments, handle computation
	of counts even for ordered loops.
	(expand_omp_for_init_vars): Handle GOMP_TASK inner_stmt.
	(expand_omp_ordered_source, expand_omp_ordered_sink,
	expand_omp_ordered_source_sink, expand_omp_for_ordered_loops): New
	functions.
	(expand_omp_for_generic): Use omp_adjust_chunk_size.  Handle linear
	clauses on worksharing loop.  Handle DOACROSS loop expansion.
	(expand_omp_for_static_nochunk): Handle linear clauses on
	worksharing loop.  Adjust expand_omp_for_init_counts
	callers.
	(expand_omp_for_static_chunk): Likewise.  Use omp_adjust_chunk_size.
	(expand_omp_simd): Handle addressable fd->loop.v.  Adjust
	expand_omp_for_init_counts callers.
	(expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): New
	functions.
	(expand_omp_for): Call expand_omp_taskloop_for_* for taskloop.
	Handle doacross loops.
	(expand_omp_target): Handle GF_OMP_TARGET_KIND_ENTER_DATA and
	GF_OMP_TARGET_KIND_EXIT_DATA.  Pass flags and depend arguments to
	GOMP_target_{41,update_41,enter_exit_data} libcalls.
	(expand_omp): Don't expand ordered depend constructs here, record
	ord_stmt instead for later expand_omp_for_generic.
	(build_omp_regions_1): Handle GF_OMP_TARGET_KIND_ENTER_DATA and
	GF_OMP_TARGET_KIND_EXIT_DATA.  Treat GIMPLE_OMP_ORDERED with depend
	clause as stand-alone directive.
	(lower_omp_ordered_clauses): New function.
	(lower_omp_ordered): Handle OMP_CLAUSE_SIMD, for OMP_CLAUSE_DEPEND
	don't lower anything.
	(lower_omp_for_lastprivate): Use last _looptemp_ clause
	on taskloop for comparison.
	(lower_omp_for): Handle taskloop constructs.  Adjust OMP_CLAUSE_DECL
	and OMP_CLAUSE_LINEAR_STEP so that expand_omp_for_* can use it during
	expansion for linear adjustments.
	(create_task_copyfn): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE.
	(lower_depend_clauses): Assert not seeing sink/source depend kinds.
	Set TREE_ADDRESSABLE on array.  Change first argument from gimple *
	to tree * pointing to the stmt's clauses.
	(lower_omp_taskreg): Adjust lower_depend_clauses caller.
	(lower_omp_target): Handle GF_OMP_TARGET_KIND_ENTER_DATA
	and GF_OMP_TARGET_KIND_EXIT_DATA, depend clauses,
	GOMP_MAP_{RELEASE,ALWAYS_{TO,FROM,TOFROM},FIRSTPRIVATE_POINTER,STRUCT}
	map kinds, OMP_CLAUSE_{FIRSTPRIVATE,PRIVATE,{IS,USE}_DEVICE_PTR
	clauses.  Always use short kind and 8-bit align shift.
	(lower_omp_regimplify_p): Use IS_TYPE_OR_DECL_P macro.
	(struct lower_omp_regimplify_operands_data): New type.
	(lower_omp_regimplify_operands_p, lower_omp_regimplify_operands):
	New functions.
	(lower_omp_1): Use lower_omp_regimplify_operands instead of
	gimple_regimplify_operands.
	(make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_ENTER_DATA and
	GF_OMP_TARGET_KIND_EXIT_DATA.  Treat GIMPLE_OMP_ORDERED with depend
	clause as stand-alone directive.
	(simd_clone_clauses_extract): Honor OMP_CLAUSE_LINEAR_KIND.
	(simd_clone_mangle): Mangle the various linear kinds
	per the new ABI.
	(simd_clone_adjust_argument_types): Handle
	SIMD_CLONE_ARG_TYPE_LINEAR_*_CONSTANT_STEP.
	(simd_clone_init_simd_arrays): Don't do anything for uval.
	(simd_clone_adjust): Handle
	SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP like
	SIMD_CLONE_ARG_TYPE_LINEAR_CONSTANT_STEP.
	Handle SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP.
	* omp-low.h (omp_member_access_dummy_var): New prototype.
	* passes.def (pass_simduid_cleanup): Schedule another copy of the
	pass after all optimizations.
	* tree.c (omp_clause_code_name): Add entries for
	OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}
	and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD}.
	(omp_clause_num_ops): Likewise.  Bump number of OMP_CLAUSE_REDUCTION
	arguments to 5 and for OMP_CLAUSE_ORDERED to 1.
	(walk_tree_1): Adjust for OMP_CLAUSE_ORDERED having 1 argument and
	OMP_CLAUSE_REDUCTION 5 arguments.  Handle
	OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}
	and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD}
	clauses.
	* tree-core.h (enum omp_clause_linear_kind): New.
	(struct tree_omp_clause): Change type of map_kind
	from unsigned char to unsigned int.  Add subcode.if_modifier
	and subcode.linear_kind fields.
	(enum omp_clause_code): Add
	OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}
	and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD}.
	(OMP_CLAUSE_REDUCTION): Document
	OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER.
	(enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_{SOURCE,SINK}.
	* tree.def (OMP_FOR): Add OMP_FOR_ORIG_DECLS operand.
	(OMP_CRITICAL): Move before OMP_SINGLE.  Add OMP_CRITICAL_CLAUSES
	operand.
	(OMP_ORDERED): Move before OMP_SINGLE.  Add OMP_ORDERED_CLAUSES
	operand.
	(OMP_TASKLOOP, OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA): New tree
	codes.
	* tree.h (OMP_BODY): Replace OMP_CRITICAL with OMP_TASKGROUP.
	(OMP_CLAUSE_SET_MAP_KIND): Cast to unsigned int rather than unsigned
	char.
	(OMP_CRITICAL_NAME): Adjust to be 3rd operand instead of 2nd.
	(OMP_CLAUSE_NUM_TASKS_EXPR): Formatting fix.
	(OMP_STANDALONE_CLAUSES): Adjust to cover OMP_TARGET_{ENTER,EXIT}_DATA.
	(OMP_CLAUSE_DEPEND_SINK_NEGATIVE, OMP_TARGET_COMBINED,
	OMP_CLAUSE_MAP_PRIVATE, OMP_FOR_ORIG_DECLS, OMP_CLAUSE_IF_MODIFIER,
	OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION, OMP_CRITICAL_CLAUSES,
	OMP_CLAUSE_PRIVATE_TASKLOOP_IV, OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV,
	OMP_CLAUSE_HINT_EXPR, OMP_CLAUSE_SCHEDULE_SIMD,
	OMP_CLAUSE_LINEAR_KIND, OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER,
	OMP_CLAUSE_SHARED_FIRSTPRIVATE, OMP_ORDERED_CLAUSES,
	OMP_TARGET_ENTER_DATA_CLAUSES, OMP_TARGET_EXIT_DATA_CLAUSES,
	OMP_CLAUSE_NUM_TASKS_EXPR, OMP_CLAUSE_GRAINSIZE_EXPR,
	OMP_CLAUSE_PRIORITY_EXPR, OMP_CLAUSE_ORDERED_EXPR): Define.
	* tree-inline.c (remap_gimple_stmt): Handle clauses on
	GIMPLE_OMP_ORDERED and GIMPLE_OMP_CRITICAL.  For
	IFN_GOMP_SIMD_ORDERED_{START,END} set has_simduid_loops.
	* tree-nested.c (convert_nonlocal_omp_clauses): Handle
	OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,SIMDLEN,PRIORITY,SIMD}
	and OMP_CLAUSE_{GRAINSIZE,NUM_TASKS,HINT,NOGROUP,THREADS,DEFAULTMAP}
	clauses.  Handle OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER.
	(convert_local_omp_clauses): Likewise.
	* tree-pretty-print.c (dump_omp_clause): Handle
	OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,SIMDLEN,PRIORITY,SIMD}
	and OMP_CLAUSE_{GRAINSIZE,NUM_TASKS,HINT,NOGROUP,THREADS,DEFAULTMAP}
	clauses.  Handle OMP_CLAUSE_IF_MODIFIER, OMP_CLAUSE_ORDERED_EXPR,
	OMP_CLAUSE_SCHEDULE_SIMD, OMP_CLAUSE_LINEAR_KIND,
	OMP_CLAUSE_DEPEND_{SOURCE,SINK}.  Use "delete" for
	GOMP_MAP_FORCE_DEALLOC.  Handle
	GOMP_MAP_{ALWAYS_{TO,FROM,TOFROM},RELEASE,FIRSTPRIVATE_POINTER,STRUCT}.
	(dump_generic_node): Handle OMP_TASKLOOP, OMP_TARGET_{ENTER,EXIT}_DATA
	and clauses on OMP_ORDERED and OMP_CRITICAL.
	* tree-vectorizer.c (adjust_simduid_builtins): Adjust comment.
	Remove IFN_GOMP_SIMD_ORDERED_{START,END}.
	(vectorize_loops): Adjust comments.
	(pass_simduid_cleanup::execute): Likewise.
	* tree-vect-stmts.c (vectorizable_simd_clone_call): Handle
	SIMD_CLONE_ARG_TYPE_LINEAR_{REF,VAL,UVAL}_CONSTANT_STEP.
	* wide-int.h (wi::gcd): New.
gcc/c-family/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>

	* c-common.c (enum c_builtin_type): Define DEF_FUNCTION_TYPE_9,
	DEF_FUNCTION_TYPE_10 and DEF_FUNCTION_TYPE_11.
	(c_define_builtins): Likewise.
	* c-common.h (enum c_omp_clause_split): Add
	C_OMP_CLAUSE_SPLIT_TASKLOOP.
	(c_finish_omp_critical, c_finish_omp_ordered): Add CLAUSES argument.
	(c_finish_omp_for): Add ORIG_DECLV argument.
	* c-cppbuiltin.c (c_cpp_builtins): Predefine _OPENMP as
	201511 instead of 201307.
	* c-omp.c (c_finish_omp_critical): Add CLAUSES argument, set
	OMP_CRITICAL_CLAUSES to it.
	(c_finish_omp_ordered): Add CLAUSES argument, set
	OMP_ORDERED_CLAUSES to it.
	(c_finish_omp_for): Add ORIG_DECLV argument, set OMP_FOR_ORIG_DECLS
	to it if OMP_FOR.  Clear DECL_INITIAL on the IVs.
	(c_omp_split_clauses): Handle OpenMP 4.5 combined/composite
	constructs and new OpenMP 4.5 clauses.  Clear
	OMP_CLAUSE_SCHEDULE_SIMD if not combined with OMP_SIMD.  Add
	verification code.
	* c-pragma.c (omp_pragmas_simd): Add taskloop.
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_TASKLOOP.
	(enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_{DEFAULTMAP,GRAINSIZE,HINT,{IS,USE}_DEVICE_PTR}
	and PRAGMA_OMP_CLAUSE_{LINK,NOGROUP,NUM_TASKS,PRIORITY,SIMD,THREADS}.
gcc/c/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>

	* c-parser.c (c_parser_pragma): Handle PRAGMA_OMP_ORDERED here.
	(c_parser_omp_clause_name): Handle OpenMP 4.5 clauses.
	(c_parser_omp_variable_list): Handle structure elements for
	map, to and from clauses.  Handle array sections in reduction
	clause.  Formatting fixes.
	(c_parser_omp_clause_if): Add IS_OMP argument, handle parsing of
	if clause modifiers.
	(c_parser_omp_clause_num_tasks, c_parser_omp_clause_grainsize,
	c_parser_omp_clause_priority, c_parser_omp_clause_hint,
	c_parser_omp_clause_defaultmap, c_parser_omp_clause_use_device_ptr,
	c_parser_omp_clause_is_device_ptr): New functions.
	(c_parser_omp_clause_ordered): Parse optional parameter.
	(c_parser_omp_clause_reduction): Handle array reductions.
	(c_parser_omp_clause_schedule): Parse optional simd modifier.
	(c_parser_omp_clause_nogroup, c_parser_omp_clause_orderedkind): New
	functions.
	(c_parser_omp_clause_linear): Parse linear clause modifiers.
	(c_parser_omp_clause_depend_sink): New function.
	(c_parser_omp_clause_depend): Parse source/sink depend kinds.
	(c_parser_omp_clause_map): Parse release/delete map kinds and
	optional always modifier.
	(c_parser_oacc_all_clauses): Adjust c_parser_omp_clause_if
	and c_finish_omp_clauses callers.
	(c_parser_omp_all_clauses): Likewise.  Parse OpenMP 4.5 clauses.
	Parse "to" as OMP_CLAUSE_TO_DECLARE if on declare target directive.
	(c_parser_oacc_cache): Adjust c_finish_omp_clauses caller.
	(OMP_CRITICAL_CLAUSE_MASK): Define.
	(c_parser_omp_critical): Parse critical clauses.
	(c_parser_omp_for_loop): Handle doacross loops, adjust
	c_finish_omp_for and c_finish_omp_clauses callers.
	(OMP_SIMD_CLAUSE_MASK): Add simdlen clause.
	(c_parser_omp_simd): Allow ordered clause if it has no parameter.
	(OMP_FOR_CLAUSE_MASK): Add linear clause.
	(c_parser_omp_for): Disallow ordered clause when combined with
	distribute.  Disallow linear clause when combined with distribute
	and not combined with simd.
	(OMP_ORDERED_CLAUSE_MASK, OMP_ORDERED_DEPEND_CLAUSE_MASK): Define.
	(c_parser_omp_ordered): Add CONTEXT argument, remove LOC argument,
	parse clauses and if depend clause is found, don't parse a body.
	(c_parser_omp_parallel): Disallow copyin clause on target parallel.
	Allow target parallel without for after it.
	(OMP_TASK_CLAUSE_MASK): Add priority clause.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add use_device_ptr clause.
	(c_parser_omp_target_data): Diagnose no map clauses or clauses with
	invalid kinds.
	(OMP_TARGET_UPDATE_CLAUSE_MASK): Add depend and nowait clauses.
	(OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
	OMP_TARGET_EXIT_DATA_CLAUSE_MASK): Define.
	(c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): New
	functions.
	(OMP_TARGET_CLAUSE_MASK): Add depend, nowait, private, firstprivate,
	defaultmap and is_device_ptr clauses.
	(c_parser_omp_target): Parse target parallel and target simd.  Set
	OMP_TARGET_COMBINED on combined constructs.  Parse target enter data
	and target exit data.  Diagnose invalid map kinds.
	(OMP_DECLARE_TARGET_CLAUSE_MASK): Define.
	(c_parser_omp_declare_target): Parse OpenMP 4.5 forms of this
	construct.
	(c_parser_omp_declare_reduction): Use STRIP_NOPS when checking for
	&omp_priv.
	(OMP_TASKLOOP_CLAUSE_MASK): Define.
	(c_parser_omp_taskloop): New function.
	(c_parser_omp_construct): Don't handle PRAGMA_OMP_ORDERED here,
	handle PRAGMA_OMP_TASKLOOP.
	(c_parser_cilk_for): Adjust c_finish_omp_clauses callers.
	* c-tree.h (c_finish_omp_clauses): Add two new arguments.
	* c-typeck.c (handle_omp_array_sections_1): Fix comment typo.
	Add IS_OMP argument, handle structure element bases, diagnose
	bitfields, pass IS_OMP recursively, diagnose known zero length
	array sections in depend clauses, handle array sections in reduction
	clause, diagnose negative length even for pointers.
	(handle_omp_array_sections): Add IS_OMP argument, use auto_vec for
	types, pass IS_OMP down to handle_omp_array_sections_1, handle
	array sections in reduction clause, set
	OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if map could be zero
	length array section, use GOMP_MAP_FIRSTPRIVATE_POINTER for IS_OMP.
	(c_finish_omp_clauses): Add IS_OMP and DECLARE_SIMD arguments.
	Handle new OpenMP 4.5 clauses and new restrictions for the old ones.
gcc/cp/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>

	* class.c (finish_struct_1): Call finish_omp_declare_simd_methods.
	* cp-gimplify.c (cp_gimplify_expr): Handle OMP_TASKLOOP.
	(cp_genericize_r): Likewise.
	(cxx_omp_finish_clause): Don't diagnose references.
	(cxx_omp_disregard_value_expr): New function.
	* cp-objcp-common.h (LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR): Redefine.
	* cp-tree.h (OMP_FOR_GIMPLIFYING_P): Document for OMP_TASKLOOP.
	(DECL_OMP_PRIVATIZED_MEMBER): Define.
	(finish_omp_declare_simd_methods, push_omp_privatization_clauses,
	pop_omp_privatization_clauses, save_omp_privatization_clauses,
	restore_omp_privatization_clauses, omp_privatize_field,
	cxx_omp_disregard_value_expr): New prototypes.
	(finish_omp_clauses): Add two new arguments.
	(finish_omp_for): Add ORIG_DECLV argument.
	* parser.c (cp_parser_lambda_body): Call
	save_omp_privatization_clauses and restore_omp_privatization_clauses.
	(cp_parser_omp_clause_name): Handle OpenMP 4.5 clauses.
	(cp_parser_omp_var_list_no_open): Handle structure elements for
	map, to and from clauses.  Handle array sections in reduction
	clause.  Parse this keyword.  Formatting fixes.
	(cp_parser_omp_clause_if): Add IS_OMP argument, handle parsing of
	if clause modifiers.
	(cp_parser_omp_clause_num_tasks, cp_parser_omp_clause_grainsize,
	cp_parser_omp_clause_priority, cp_parser_omp_clause_hint,
	cp_parser_omp_clause_defaultmap): New functions.
	(cp_parser_omp_clause_ordered): Parse optional parameter.
	(cp_parser_omp_clause_reduction): Handle array reductions.
	(cp_parser_omp_clause_schedule): Parse optional simd modifier.
	(cp_parser_omp_clause_nogroup, cp_parser_omp_clause_orderedkind):
	New functions.
	(cp_parser_omp_clause_linear): Parse linear clause modifiers.
	(cp_parser_omp_clause_depend_sink): New function.
	(cp_parser_omp_clause_depend): Parse source/sink depend kinds.
	(cp_parser_omp_clause_map): Parse release/delete map kinds and
	optional always modifier.
	(cp_parser_oacc_all_clauses): Adjust cp_parser_omp_clause_if
	and finish_omp_clauses callers.
	(cp_parser_omp_all_clauses): Likewise.  Parse OpenMP 4.5 clauses.
	Parse "to" as OMP_CLAUSE_TO_DECLARE if on declare target directive.
	(OMP_CRITICAL_CLAUSE_MASK): Define.
	(cp_parser_omp_critical): Parse critical clauses.
	(cp_parser_omp_for_incr): Use cp_tree_equal if
	processing_template_decl.
	(cp_parser_omp_for_loop_init): Return tree instead of bool.  Handle
	non-static data member iterators.
	(cp_parser_omp_for_loop): Handle doacross loops, adjust
	finish_omp_for and finish_omp_clauses callers.
	(cp_omp_split_clauses): Adjust finish_omp_clauses caller.
	(OMP_SIMD_CLAUSE_MASK): Add simdlen clause.
	(cp_parser_omp_simd): Allow ordered clause if it has no parameter.
	(OMP_FOR_CLAUSE_MASK): Add linear clause.
	(cp_parser_omp_for): Disallow ordered clause when combined with
	distribute.  Disallow linear clause when combined with distribute
	and not combined with simd.
	(OMP_ORDERED_CLAUSE_MASK, OMP_ORDERED_DEPEND_CLAUSE_MASK): Define.
	(cp_parser_omp_ordered): Add CONTEXT argument, return bool instead
	of tree, parse clauses and if depend clause is found, don't parse
	a body.
	(cp_parser_omp_parallel): Disallow copyin clause on target parallel.
	Allow target parallel without for after it.
	(OMP_TASK_CLAUSE_MASK): Add priority clause.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add use_device_ptr clause.
	(cp_parser_omp_target_data): Diagnose no map clauses or clauses with
	invalid kinds.
	(OMP_TARGET_UPDATE_CLAUSE_MASK): Add depend and nowait clauses.
	(OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
	OMP_TARGET_EXIT_DATA_CLAUSE_MASK): Define.
	(cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data): New
	functions.
	(OMP_TARGET_CLAUSE_MASK): Add depend, nowait, private, firstprivate,
	defaultmap and is_device_ptr clauses.
	(cp_parser_omp_target): Parse target parallel and target simd.  Set
	OMP_TARGET_COMBINED on combined constructs.  Parse target enter data
	and target exit data.  Diagnose invalid map kinds.
	(cp_parser_oacc_cache): Adjust finish_omp_clauses caller.
	(OMP_DECLARE_TARGET_CLAUSE_MASK): Define.
	(cp_parser_omp_declare_target): Parse OpenMP 4.5 forms of this
	construct.
	(OMP_TASKLOOP_CLAUSE_MASK): Define.
	(cp_parser_omp_taskloop): New function.
	(cp_parser_omp_construct): Don't handle PRAGMA_OMP_ORDERED here,
	handle PRAGMA_OMP_TASKLOOP.
	(cp_parser_pragma): Handle PRAGMA_OMP_ORDERED here directly,
	handle PRAGMA_OMP_TASKLOOP, call push_omp_privatization_clauses
	and pop_omp_privatization_clauses around parsing calls.
	(cp_parser_cilk_for): Adjust finish_omp_clauses caller.
	* pt.c (apply_late_template_attributes): Adjust tsubst_omp_clauses
	and finish_omp_clauses callers.
	(tsubst_omp_clause_decl): Return NULL if decl is NULL.
	For TREE_LIST, copy over OMP_CLAUSE_DEPEND_SINK_NEGATIVE bit.
	Use tsubst_expr instead of tsubst_copy, undo convert_from_reference
	effects.
	(tsubst_omp_clauses): Add ALLOW_FIELDS argument.  Handle new
	OpenMP 4.5 clauses.  Use tsubst_omp_clause_decl for more clauses.
	If ALLOW_FIELDS, handle non-static data members in the clauses.
	Clear OMP_CLAUSE_LINEAR_STEP if it has been cleared before.
	(omp_parallel_combined_clauses): New variable.
	(tsubst_omp_for_iterator): Add ORIG_DECLV argument, recur on
	OMP_FOR_ORIG_DECLS, handle non-static data member iterators.
	Improve handling of clauses on combined constructs.
	(tsubst_expr): Call push_omp_privatization_clauses and
	pop_omp_privatization_clauses around instantiation of certain
	OpenMP constructs, improve handling of clauses on combined
	constructs, handle OMP_TASKLOOP, adjust tsubst_omp_for_iterator,
	tsubst_omp_clauses and finish_omp_for callers, handle clauses on
	critical and ordered, handle OMP_TARGET_{ENTER,EXIT}_DATA.
	(instantiate_decl): Call save_omp_privatization_clauses and
	restore_omp_privatization_clauses around instantiation.
	(dependent_omp_for_p): Fix up comment typo.  Handle SCOPE_REF.
	* semantics.c (omp_private_member_map, omp_private_member_vec,
	omp_private_member_ignore_next): New variables.
	(finish_non_static_data_member): Return dummy decl for privatized
	non-static data members.
	(omp_clause_decl_field, omp_clause_printable_decl,
	omp_note_field_privatization, omp_privatize_field): New functions.
	(handle_omp_array_sections_1): Fix comment typo.
	Add IS_OMP argument, handle structure element bases, diagnose
	bitfields, pass IS_OMP recursively, diagnose known zero length
	array sections in depend clauses, handle array sections in reduction
	clause, diagnose negative length even for pointers.
	(handle_omp_array_sections): Add IS_OMP argument, use auto_vec for
	types, pass IS_OMP down to handle_omp_array_sections_1, handle
	array sections in reduction clause, set
	OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if map could be zero
	length array section, use GOMP_MAP_FIRSTPRIVATE_POINTER for IS_OMP.
	(finish_omp_reduction_clause): Handle array sections and arrays.
	Use omp_clause_printable_decl.
	(finish_omp_declare_simd_methods, cp_finish_omp_clause_depend_sink):
	New functions.
	(finish_omp_clauses): Add ALLOW_FIELDS and DECLARE_SIMD arguments.
	Handle new OpenMP 4.5 clauses and new restrictions for the old
	ones, handle non-static data members, reject this keyword when not
	allowed.
	(push_omp_privatization_clauses, pop_omp_privatization_clauses,
	save_omp_privatization_clauses, restore_omp_privatization_clauses):
	New functions.
	(handle_omp_for_class_iterator): Handle OMP_TASKLOOP class iterators.
	Add collapse and ordered arguments.  Fix handling of lastprivate
	iterators in doacross loops.
	(finish_omp_for): Add ORIG_DECLV argument, handle doacross loops,
	adjust c_finish_omp_for, handle_omp_for_class_iterator and
	finish_omp_clauses callers.  Fill in OMP_CLAUSE_LINEAR_STEP on simd
	loops with non-static data member iterators.
gcc/fortran/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* f95-lang.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10,
	DEF_FUNCTION_TYPE_11, DEF_FUNCTION_TYPE_VAR_1): Define.
	* trans-openmp.c (gfc_trans_omp_clauses): Set
	OMP_CLAUSE_IF_MODIFIER to ERROR_MARK, OMP_CLAUSE_ORDERED_EXPR
	to NULL.
	(gfc_trans_omp_critical): Adjust for addition of clauses.
	(gfc_trans_omp_ordered): Likewise.
	* types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR,
	BT_FN_BOOL_UINT_ULLPTR_ULLPTR_ULLPTR,
	BT_FN_BOOL_UINT_LONGPTR_LONG_LONGPTR_LONGPTR,
	BT_FN_BOOL_UINT_ULLPTR_ULL_ULLPTR_ULLPTR,
	BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
	BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_ULL_ULL_ULL,
	BT_FN_VOID_LONG_VAR, BT_FN_VOID_ULL_VAR): New.
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR): Remove.
gcc/lto/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>

	* lto-lang.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10,
	DEF_FUNCTION_TYPE_11): Define.
gcc/jit/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>

	* jit-builtins.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10,
	DEF_FUNCTION_TYPE_11): Define.
	* jit-builtins.h (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10,
	DEF_FUNCTION_TYPE_11): Define.
gcc/ada/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>

	* gcc-interface/utils.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10,
	DEF_FUNCTION_TYPE_11): Define.
gcc/testsuite/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>

	* c-c++-common/gomp/cancel-1.c (f2): Add map clause to target data.
	* c-c++-common/gomp/clauses-1.c: New test.
	* c-c++-common/gomp/clauses-2.c: New test.
	* c-c++-common/gomp/clauses-3.c: New test.
	* c-c++-common/gomp/clauses-4.c: New test.
	* c-c++-common/gomp/declare-target-1.c: New test.
	* c-c++-common/gomp/declare-target-2.c: New test.
	* c-c++-common/gomp/depend-3.c: New test.
	* c-c++-common/gomp/depend-4.c: New test.
	* c-c++-common/gomp/doacross-1.c: New test.
	* c-c++-common/gomp/if-1.c: New test.
	* c-c++-common/gomp/if-2.c: New test.
	* c-c++-common/gomp/linear-1.c: New test.
	* c-c++-common/gomp/map-2.c: New test.
	* c-c++-common/gomp/map-3.c: New test.
	* c-c++-common/gomp/nesting-1.c (f_omp_parallel,
	f_omp_target_data): Add map clause to target data.
	* c-c++-common/gomp/nesting-warn-1.c (f_omp_target): Likewise.
	* c-c++-common/gomp/ordered-1.c: New test.
	* c-c++-common/gomp/ordered-2.c: New test.
	* c-c++-common/gomp/ordered-3.c: New test.
	* c-c++-common/gomp/pr61486-1.c (foo): Remove linear clause
	on non-iterator.
	* c-c++-common/gomp/pr61486-2.c (test, test2): Remove ordered
	clause and ordered construct where no longer allowed.
	* c-c++-common/gomp/priority-1.c: New test.
	* c-c++-common/gomp/reduction-1.c: New test.
	* c-c++-common/gomp/schedule-simd-1.c: New test.
	* c-c++-common/gomp/sink-1.c: New test.
	* c-c++-common/gomp/sink-2.c: New test.
	* c-c++-common/gomp/sink-3.c: New test.
	* c-c++-common/gomp/sink-4.c: New test.
	* c-c++-common/gomp/udr-1.c: New test.
	* c-c++-common/taskloop-1.c: New test.
	* c-c++-common/cpp/openmp-define-3.c: Adjust for the new
	value of _OPENMP macro.
	* c-c++-common/cilk-plus/PS/body.c (foo): Adjust expected diagnostics.
	* c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_parallel,
	f_acc_kernels, f_acc_data, f_acc_loop): Add map clause to target data.
	* gcc.dg/gomp/clause-1.c:
	* gcc.dg/gomp/reduction-1.c: New test.
	* gcc.dg/gomp/sink-fold-1.c: New test.
	* gcc.dg/gomp/sink-fold-2.c: New test.
	* gcc.dg/gomp/sink-fold-3.c: New test.
	* gcc.dg/vect/vect-simd-clone-15.c: New test.
	* g++.dg/gomp/clause-1.C (T::test): Remove dg-error on privatization
	of non-static data members.
	* g++.dg/gomp/clause-3.C (foo): Remove one dg-error directive.
	Add some linear clause tests.
	* g++.dg/gomp/declare-simd-3.C: New test.
	* g++.dg/gomp/linear-1.C: New test.
	* g++.dg/gomp/member-1.C: New test.
	* g++.dg/gomp/member-2.C: New test.
	* g++.dg/gomp/pr66571-2.C: New test.
	* g++.dg/gomp/pr67504.C (foo): Add test for ordered clause with
	dependent argument.
	* g++.dg/gomp/pr67522.C (foo): Add test for invalid array section
	in reduction clause.
	* g++.dg/gomp/reference-1.C: New test.
	* g++.dg/gomp/sink-1.C: New test.
	* g++.dg/gomp/sink-2.C: New test.
	* g++.dg/gomp/sink-3.C: New test.
	* g++.dg/gomp/task-1.C: Remove both dg-error directives.
	* g++.dg/gomp/this-1.C: New test.
	* g++.dg/gomp/this-2.C: New test.
	* g++.dg/vect/simd-clone-2.cc: New test.
	* g++.dg/vect/simd-clone-2.h: New test.
	* g++.dg/vect/simd-clone-3.cc: New test.
	* g++.dg/vect/simd-clone-4.cc: New test.
	* g++.dg/vect/simd-clone-4.h: New test.
	* g++.dg/vect/simd-clone-5.cc: New test.
include/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* gomp-constants.h (GOMP_MAP_FLAG_ALWAYS): Define.
	(enum gomp_map_kind): Add GOMP_MAP_FIRSTPRIVATE,
	GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_USE_DEVICE_PTR,
	GOMP_MAP_ZERO_LEN_ARRAY_SECTION, GOMP_MAP_ALWAYS_TO,
	GOMP_MAP_ALWAYS_FROM, GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_STRUCT,
	GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION, GOMP_MAP_DELETE,
	GOMP_MAP_RELEASE, GOMP_MAP_FIRSTPRIVATE_POINTER.
	(GOMP_MAP_ALWAYS_TO_P, GOMP_MAP_ALWAYS_FROM_P): Define.
	(GOMP_TASK_FLAG_UNTIED, GOMP_TASK_FLAG_FINAL, GOMP_TASK_FLAG_MERGEABLE,
	GOMP_TASK_FLAG_DEPEND, GOMP_TASK_FLAG_PRIORITY, GOMP_TASK_FLAG_UP,
	GOMP_TASK_FLAG_GRAINSIZE, GOMP_TASK_FLAG_IF, GOMP_TASK_FLAG_NOGROUP,
	GOMP_TARGET_FLAG_NOWAIT, GOMP_TARGET_FLAG_EXIT_DATA,
	GOMP_TARGET_FLAG_UPDATE): Define.
libgomp/
2015-10-13  Jakub Jelinek  <jakub@redhat.com>
	    Aldy Hernandez  <aldyh@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* config/linux/affinity.c (omp_get_place_num_procs,
	omp_get_place_proc_ids, gomp_get_place_proc_ids_8): New functions.
	* config/linux/doacross.h: New file.
	* config/posix/affinity.c (omp_get_place_num_procs,
	omp_get_place_proc_ids, gomp_get_place_proc_ids_8): New functions.
	* config/posix/doacross.h: New file.
	* env.c: Include gomp-constants.h.
	(struct gomp_task_icv): Rename run_sched_modifier to
	run_sched_chunk_size.
	(gomp_max_task_priority_var): New variable.
	(parse_schedule): Rename run_sched_modifier to run_sched_chunk_size.
	(handle_omp_display_env): Change _OPENMP value from 201307 to
	201511.  Print OMP_MAX_TASK_PRIORITY.
	(initialize_env): Parse OMP_MAX_TASK_PRIORITY.
	(omp_set_schedule, omp_get_schedule): Rename modifier argument to
	chunk_size and run_sched_modifier to run_sched_chunk_size.
	(omp_get_max_task_priority, omp_get_initial_device,
	omp_get_num_places, omp_get_place_num, omp_get_partition_num_places,
	omp_get_partition_place_nums): New functions.
	* fortran.c (omp_set_schedule_, omp_set_schedule_8_,
	omp_get_schedule_, omp_get_schedule_8_): Rename modifier argument
	to chunk_size.
	(omp_get_num_places_, omp_get_place_num_procs_,
	omp_get_place_num_procs_8_, omp_get_place_proc_ids_,
	omp_get_place_proc_ids_8_, omp_get_place_num_,
	omp_get_partition_num_places_, omp_get_partition_place_nums_,
	omp_get_partition_place_nums_8_, omp_get_initial_device_,
	omp_get_max_task_priority_): New functions.
	* libgomp_g.h (GOMP_loop_doacross_static_start,
	GOMP_loop_doacross_dynamic_start, GOMP_loop_doacross_guided_start,
	GOMP_loop_doacross_runtime_start, GOMP_loop_ull_doacross_static_start,
	GOMP_loop_ull_doacross_dynamic_start,
	GOMP_loop_ull_doacross_guided_start,
	GOMP_loop_ull_doacross_runtime_start, GOMP_doacross_post,
	GOMP_doacross_wait, GOMP_doacross_ull_post, GOMP_doacross_wait,
	GOMP_taskloop, GOMP_taskloop_ull, GOMP_target_41,
	GOMP_target_data_41, GOMP_target_update_41,
	GOMP_target_enter_exit_data): New prototypes.
	(GOMP_task): Add prototype argument.
	* libgomp.h (_LIBGOMP_CHECKING_): Define to 0 if not yet defined.
	(struct gomp_doacross_work_share): New type.
	(struct gomp_work_share): Add doacross field.
	(struct gomp_task_icv): Rename run_sched_modifier to
	run_sched_chunk_size.
	(enum gomp_task_kind): Rename GOMP_TASK_IFFALSE to
	GOMP_TASK_UNDEFERRED.  Add comments.
	(struct gomp_task_depend_entry): Add comments.
	(struct gomp_task): Likewise.
	(struct gomp_taskgroup): Likewise.
	(struct gomp_target_task): New type.
	(struct gomp_team): Add comment.
	(gomp_get_place_proc_ids_8, gomp_doacross_init,
	gomp_doacross_ull_init, gomp_task_maybe_wait_for_dependencies,
	gomp_create_target_task, gomp_target_task_fn): New prototypes.
	(struct target_var_desc): New type.
	(struct target_mem_desc): Adjust comment.  Use struct
	target_var_desc instead of splay_tree_key for list.
	(REFCOUNT_INFINITY): Define.
	(struct splay_tree_key_s): Remove copy_from field.
	(struct gomp_device_descr): Add dev2dev_func field.
	(enum gomp_map_vars_kind): New enum.
	(gomp_map_vars): Add one argument.
	* libgomp.map (OMP_4.5): Export omp_get_max_task_priority,
	omp_get_max_task_priority_, omp_get_num_places, omp_get_num_places_,
	omp_get_place_num_procs, omp_get_place_num_procs_,
	omp_get_place_num_procs_8_, omp_get_place_proc_ids,
	omp_get_place_proc_ids_, omp_get_place_proc_ids_8_, omp_get_place_num,
	omp_get_place_num_, omp_get_partition_num_places,
	omp_get_partition_num_places_, omp_get_partition_place_nums,
	omp_get_partition_place_nums_, omp_get_partition_place_nums_8_,
	omp_get_initial_device, omp_get_initial_device_, omp_target_alloc,
	omp_target_free, omp_target_is_present, omp_target_memcpy,
	omp_target_memcpy_rect, omp_target_associate_ptr and
	omp_target_disassociate_ptr.
	(GOMP_4.0.2): Renamed to ...
	(GOMP_4.5): ... this.  Export GOMP_target_41, GOMP_target_data_41,
	GOMP_target_update_41, GOMP_target_enter_exit_data, GOMP_taskloop,
	GOMP_taskloop_ull, GOMP_loop_doacross_dynamic_start,
	GOMP_loop_doacross_guided_start, GOMP_loop_doacross_runtime_start,
	GOMP_loop_doacross_static_start, GOMP_doacross_post,
	GOMP_doacross_wait, GOMP_loop_ull_doacross_dynamic_start,
	GOMP_loop_ull_doacross_guided_start,
	GOMP_loop_ull_doacross_runtime_start,
	GOMP_loop_ull_doacross_static_start, GOMP_doacross_ull_post and
	GOMP_doacross_ull_wait.
	* libgomp.texi: Document omp_get_max_task_priority.
	Rename modifier argument to chunk_size for omp_set_schedule and
	omp_get_schedule.  Document OMP_MAX_TASK_PRIORITY env var.
	* loop.c (GOMP_loop_runtime_start): Adjust for run_sched_modifier
	to run_sched_chunk_size renaming.
	(GOMP_loop_ordered_runtime_start): Likewise.
	(gomp_loop_doacross_static_start, gomp_loop_doacross_dynamic_start,
	gomp_loop_doacross_guided_start, GOMP_loop_doacross_runtime_start,
	GOMP_parallel_loop_runtime_start): New functions.
	(GOMP_parallel_loop_runtime): Adjust for run_sched_modifier
	to run_sched_chunk_size renaming.
	(GOMP_loop_doacross_static_start, GOMP_loop_doacross_dynamic_start,
	GOMP_loop_doacross_guided_start): New functions or aliases.
	* loop_ull.c (GOMP_loop_ull_runtime_start): Adjust for
	run_sched_modifier to run_sched_chunk_size renaming.
	(GOMP_loop_ull_ordered_runtime_start): Likewise.
	(gomp_loop_ull_doacross_static_start,
	gomp_loop_ull_doacross_dynamic_start,
	gomp_loop_ull_doacross_guided_start,
	GOMP_loop_ull_doacross_runtime_start): New functions.
	(GOMP_loop_ull_doacross_static_start,
	GOMP_loop_ull_doacross_dynamic_start,
	GOMP_loop_ull_doacross_guided_start): New functions or aliases.
	* oacc-mem.c (acc_map_data, present_create_copy,
	gomp_acc_insert_pointer): Pass GOMP_MAP_VARS_OPENACC instead of false
	to gomp_map_vars.
	(gomp_acc_remove_pointer): Use copy_from from target_var_desc.
	* oacc-parallel.c (GOACC_data_start): Pass GOMP_MAP_VARS_OPENACC
	instead of false to gomp_map_vars.
	(GOACC_parallel_keyed): Likewise.  Use copy_from from target_var_desc.
	* omp.h.in (omp_lock_hint_t): New type.
	(omp_init_lock_with_hint, omp_init_nest_lock_with_hint,
	omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids,
	omp_get_place_num, omp_get_partition_num_places,
	omp_get_partition_place_nums, omp_get_initial_device,
	omp_get_max_task_priority, omp_target_alloc, omp_target_free,
	omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect,
	omp_target_associate_ptr, omp_target_disassociate_ptr): New
	prototypes.
	* omp_lib.f90.in (omp_lock_hint_kind): New parameter.
	(omp_lock_hint_none, omp_lock_hint_uncontended,
	omp_lock_hint_contended, omp_lock_hint_nonspeculative,
	omp_lock_hint_speculative): New parameters.
	(omp_init_lock_with_hint, omp_init_nest_lock_with_hint,
	omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids,
	omp_get_place_num, omp_get_partition_num_places,
	omp_get_partition_place_nums, omp_get_initial_device,
	omp_get_max_task_priority): New interfaces.
	(omp_set_schedule, omp_get_schedule): Rename modifier argument
	to chunk_size.
	* omp_lib.h.in (omp_lock_hint_kind): New parameter.
	(omp_lock_hint_none, omp_lock_hint_uncontended,
	omp_lock_hint_contended, omp_lock_hint_nonspeculative,
	omp_lock_hint_speculative): New parameters.
	(omp_init_lock_with_hint, omp_init_nest_lock_with_hint,
	omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids,
	omp_get_place_num, omp_get_partition_num_places,
	omp_get_partition_place_nums, omp_get_initial_device,
	omp_get_max_task_priority): New functions and subroutines.
	* ordered.c: Include stdarg.h and string.h.
	(MAX_COLLAPSED_BITS): Define.
	(gomp_doacross_init, GOMP_doacross_post, GOMP_doacross_wait,
	gomp_doacross_ull_init, GOMP_doacross_ull_post,
	GOMP_doacross_ull_wait): New functions.
	* target.c: Include errno.h.
	(resolve_device): If device is not initialized, call
	gomp_init_device on it.
	(gomp_map_lookup): New function.
	(gomp_map_vars_existing): Add tgt_var argument, fill it in.
	Don't bump refcount if REFCOUNT_INFINITY.  Handle
	GOMP_MAP_ALWAYS_TO_P.
	(get_kind): Rename is_openacc argument to short_mapkind.
	(gomp_map_pointer): Use gomp_map_lookup.
	(gomp_map_fields_existing): New function.
	(gomp_map_vars): Rename is_openacc argument to short_mapkind
	and is_target to pragma_kind.  Handle GOMP_MAP_VARS_ENTER_DATA,
	handle GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_STRUCT,
	GOMP_MAP_USE_DEVICE_PTR, GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
	Adjust for tgt->list changed type and copy_from living in there.
	(gomp_copy_from_async): Adjust for tgt->list changed type and
	copy_from living in there.
	(gomp_unmap_vars): Likewise.
	(gomp_update): Likewise.  Rename is_openacc argument to
	short_mapkind.  Don't fail if object is not mapped.
	(gomp_load_image_to_device): Initialize refcount to
	REFCOUNT_INFINITY.
	(gomp_target_fallback): New function.
	(gomp_get_target_fn_addr): Likewise.
	(GOMP_target): Adjust gomp_map_vars caller, use
	gomp_get_target_fn_addr and gomp_target_fallback.
	(GOMP_target_41): New function.
	(gomp_target_data_fallback): New function.
	(GOMP_target_data): Use it, adjust gomp_map_vars caller.
	(GOMP_target_data_41): New function.
	(GOMP_target_update): Adjust gomp_update caller.
	(GOMP_target_update_41): New function.
	(gomp_exit_data, GOMP_target_enter_exit_data,
	gomp_target_task_fn, omp_target_alloc, omp_target_free,
	omp_target_is_present, omp_target_memcpy,
	omp_target_memcpy_rect_worker, omp_target_memcpy_rect,
	omp_target_associate_ptr, omp_target_disassociate_ptr,
	gomp_load_plugin_for_device): New functions.
	* task.c: Include gomp-constants.h.  Include taskloop.c
	twice to get GOMP_taskloop and GOMP_taskloop_ull definitions.
	(gomp_task_handle_depend): New function.
	(GOMP_task): Use it.  Add priority argument.  Use
	gomp-constant.h constants instead of hardcoded numbers.
	Rename GOMP_TASK_IFFALSE to GOMP_TASK_UNDEFERRED.
	(gomp_create_target_task): New function.
	(verify_children_queue, verify_taskgroup_queue,
	verify_task_queue): New functions.
	(gomp_task_run_pre): Call verify_*_queue functions.
	If an upcoming tied task is about to leave the sibling or
	taskgroup queues in an invalid state, adjust appropriately.
	Remove taskgroup argument.  Add comments.
	(gomp_task_run_post_handle_dependers): Add comments.
	(gomp_task_run_post_remove_parent): Likewise.
	(gomp_barrier_handle_tasks): Adjust gomp_task_run_pre caller.
	(GOMP_taskwait): Likewise.  Add comments.
	(gomp_task_maybe_wait_for_dependencies): Fix scheduling
	problem such that the first non parent_depends_on task does not
	end up at the end of the children queue.
	(GOMP_taskgroup_start): Rename GOMP_TASK_IFFALSE to
	GOMP_TASK_UNDEFERRED.
	(GOMP_taskgroup_end): Adjust gomp_task_run_pre caller.
	* taskloop.c: New file.
	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_device_nonshared_as): New proc.
	* testsuite/libgomp.c/affinity-2.c: New test.
	* testsuite/libgomp.c/doacross-1.c: New test.
	* testsuite/libgomp.c/doacross-2.c: New test.
	* testsuite/libgomp.c/examples-4/declare_target-1.c (fib_wrapper):
	Add map clause to target.
	* testsuite/libgomp.c/examples-4/declare_target-4.c (accum): Likewise.
	* testsuite/libgomp.c/examples-4/declare_target-5.c (accum): Likewise.
	* testsuite/libgomp.c/examples-4/device-1.c (main): Likewise.
	* testsuite/libgomp.c/examples-4/device-3.c (main): Likewise.
	* testsuite/libgomp.c/examples-4/target_data-3.c (gramSchmidt):
	Likewise.
	* testsuite/libgomp.c/examples-4/teams-2.c (dotprod): Likewise.
	* testsuite/libgomp.c/examples-4/teams-3.c (dotprod): Likewise.
	* testsuite/libgomp.c/examples-4/teams-4.c (dotprod): Likewise.
	* testsuite/libgomp.c/for-2.h (OMPTGT, OMPTO, OMPFROM): Define if
	not defined.  Use those where needed.
	* testsuite/libgomp.c/for-4.c: New test.
	* testsuite/libgomp.c/for-5.c: New test.
	* testsuite/libgomp.c/for-6.c: New test.
	* testsuite/libgomp.c/linear-1.c: New test.
	* testsuite/libgomp.c/ordered-4.c: New test.
	* testsuite/libgomp.c/pr66199-2.c (f2): Adjust for linear clause
	only allowed on the loop iterator.
	* testsuite/libgomp.c/pr66199-3.c: New test.
	* testsuite/libgomp.c/pr66199-4.c: New test.
	* testsuite/libgomp.c/reduction-7.c: New test.
	* testsuite/libgomp.c/reduction-8.c: New test.
	* testsuite/libgomp.c/reduction-9.c: New test.
	* testsuite/libgomp.c/reduction-10.c: New test.
	* testsuite/libgomp.c/target-1.c (fn2, fn3, fn4): Add
	map(tofrom:s).
	* testsuite/libgomp.c/target-2.c (fn2, fn3, fn4): Likewise.
	* testsuite/libgomp.c/target-7.c (foo): Add map(h) where needed.
	* testsuite/libgomp.c/target-11.c: New test.
	* testsuite/libgomp.c/target-12.c: New test.
	* testsuite/libgomp.c/target-13.c: New test.
	* testsuite/libgomp.c/target-14.c: New test.
	* testsuite/libgomp.c/target-15.c: New test.
	* testsuite/libgomp.c/target-16.c: New test.
	* testsuite/libgomp.c/target-17.c: New test.
	* testsuite/libgomp.c/target-18.c: New test.
	* testsuite/libgomp.c/target-19.c: New test.
	* testsuite/libgomp.c/target-20.c: New test.
	* testsuite/libgomp.c/target-21.c: New test.
	* testsuite/libgomp.c/target-22.c: New test.
	* testsuite/libgomp.c/target-23.c: New test.
	* testsuite/libgomp.c/target-24.c: New test.
	* testsuite/libgomp.c/target-25.c: New test.
	* testsuite/libgomp.c/target-26.c: New test.
	* testsuite/libgomp.c/target-27.c: New test.
	* testsuite/libgomp.c/taskloop-1.c: New test.
	* testsuite/libgomp.c/taskloop-2.c: New test.
	* testsuite/libgomp.c/taskloop-3.c: New test.
	* testsuite/libgomp.c/taskloop-4.c: New test.
	* testsuite/libgomp.c++/ctor-13.C: New test.
	* testsuite/libgomp.c++/doacross-1.C: New test.
	* testsuite/libgomp.c++/examples-4/declare_target-2.C:
	Replace offload_device with offload_device_nonshared_as.
	* testsuite/libgomp.c++/for-12.C: New test.
	* testsuite/libgomp.c++/for-13.C: New test.
	* testsuite/libgomp.c++/for-14.C: New test.
	* testsuite/libgomp.c++/linear-1.C: New test.
	* testsuite/libgomp.c++/member-1.C: New test.
	* testsuite/libgomp.c++/member-2.C: New test.
	* testsuite/libgomp.c++/member-3.C: New test.
	* testsuite/libgomp.c++/member-4.C: New test.
	* testsuite/libgomp.c++/member-5.C: New test.
	* testsuite/libgomp.c++/ordered-1.C: New test.
	* testsuite/libgomp.c++/reduction-5.C: New test.
	* testsuite/libgomp.c++/reduction-6.C: New test.
	* testsuite/libgomp.c++/reduction-7.C: New test.
	* testsuite/libgomp.c++/reduction-8.C: New test.
	* testsuite/libgomp.c++/reduction-9.C: New test.
	* testsuite/libgomp.c++/reduction-10.C: New test.
	* testsuite/libgomp.c++/reference-1.C: New test.
	* testsuite/libgomp.c++/simd14.C: New test.
	* testsuite/libgomp.c++/target-2.C (fn2): Add map(tofrom: s) clause.
	* testsuite/libgomp.c++/target-5.C: New test.
	* testsuite/libgomp.c++/target-6.C: New test.
	* testsuite/libgomp.c++/target-7.C: New test.
	* testsuite/libgomp.c++/target-8.C: New test.
	* testsuite/libgomp.c++/target-9.C: New test.
	* testsuite/libgomp.c++/target-10.C: New test.
	* testsuite/libgomp.c++/target-11.C: New test.
	* testsuite/libgomp.c++/target-12.C: New test.
	* testsuite/libgomp.c++/taskloop-1.C: New test.
	* testsuite/libgomp.c++/taskloop-2.C: New test.
	* testsuite/libgomp.c++/taskloop-3.C: New test.
	* testsuite/libgomp.c++/taskloop-4.C: New test.
	* testsuite/libgomp.c++/taskloop-5.C: New test.
	* testsuite/libgomp.c++/taskloop-6.C: New test.
	* testsuite/libgomp.c++/taskloop-7.C: New test.
	* testsuite/libgomp.c++/taskloop-8.C: New test.
	* testsuite/libgomp.c++/taskloop-9.C: New test.
	* testsuite/libgomp.fortran/affinity1.f90: New test.
	* testsuite/libgomp.fortran/affinity2.f90: New test.
liboffloadmic/
2015-10-13  Ilya Verbin  <ilya.verbin@intel.com>

	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_dev2dev): New
	function.
	* plugin/offload_target_main.cpp (__offload_target_tgt2tgt): New
	static function, register it in liboffloadmic.


git-svn-id: https://gcc.gnu.org/svn/gcc/trunk@228777 138bc75d-0d04-0410-961f-82ee72b054a4

49 files changed, 3265 insertions, 33 deletions

diff --git a/libgomp/testsuite/libgomp.c/affinity-2.c b/libgomp/testsuite/libgomp.c/affinity-2.c
new file mode 100644
index 00000000000..f8216574704
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/affinity-2.c
@@ -0,0 +1,89 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_PROC_BIND "spread,close" } */
+/* { dg-set-target-env-var OMP_PLACES "{6,7}:4:-2,!{2,3}" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS "2" } */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+int *
+get_buf (int nump)
+{
+  static int *buf;
+  static size_t buf_size;
+  if ((size_t) nump > buf_size)
+    {
+      buf_size *= 2;
+      if (nump > buf_size)
+	buf_size = nump + 64;
+      int *bufn = realloc (buf, buf_size * sizeof (int));
+      if (bufn == NULL)
+	{
+	  fprintf (stderr, "memory allocation error\n");
+	  exit (1);
+	}
+      buf = bufn;
+    }
+  return buf;
+}
+
+void
+print_place (int count, int *ids)
+{
+  int i, j;
+  printf ("{");
+  for (i = 0; i < count; i++)
+    {
+      for (j = i + 1; j < count; j++)
+	if (ids[j] != ids[i] + (j - i))
+	  break;
+      if (i)
+	printf (",");
+      if (j == i + 1)
+	printf ("%d", ids[i]);
+      else
+	{
+	  printf ("%d:%d", ids[i], j - i);
+	  i = j - 1;
+	}
+    }
+  printf ("}\n");
+}
+
+void
+print_place_var (void)
+{
+  int place = omp_get_place_num ();
+  int num_places = omp_get_partition_num_places ();
+  int *ids = get_buf (num_places);
+  omp_get_partition_place_nums (ids);
+  printf ("place %d\n", place);
+  if (num_places)
+    printf ("partition %d-%d\n", ids[0], ids[num_places - 1]);
+}
+
+int
+main ()
+{
+  int i, num = omp_get_num_places (), nump, *ids;
+  printf ("omp_get_num_places () == %d\n", num);
+  for (i = 0; i < num; i++)
+    {
+      printf ("place %d ", i);
+      nump = omp_get_place_num_procs (i);
+      ids = get_buf (nump);
+      omp_get_place_proc_ids (i, ids);
+      print_place (nump, ids);
+    }
+  print_place_var ();
+  omp_set_nested (1);
+  #pragma omp parallel
+    if (omp_get_thread_num () == omp_get_num_threads () - 1)
+      {
+      #pragma omp parallel
+	if (omp_get_thread_num () == omp_get_num_threads () - 1)
+	  print_place_var ();
+      }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/doacross-1.c b/libgomp/testsuite/libgomp.c/doacross-1.c
new file mode 100644
index 00000000000..0794c80ec2e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-1.c
@@ -0,0 +1,181 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8];
+volatile int d, e;
+
+int
+main ()
+{
+  int i, j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(static, 1) ordered (1) nowait
+    for (i = 0; i < N; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered depend(sink: i - 1)
+	if (i)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered depend(source)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(static, 0) ordered (3) nowait
+    for (i = 2; i < N / 16 - 1; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered depend(sink: i, j - 2, k - 1) \
+				depend(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered depend(sink: i - 3, j + 2, k - 2)
+	    if (j >= 2 && k > 1)
+	      {
+		#pragma omp atomic read
+		l = b[i][j - 2][k - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 4 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i >= 5 && j < N / 16 - 3 && k == 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 3][j + 2][k - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered depend(source)
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15)
+    for (i = 0; i < N / 32; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+#define A(n) ,n
+#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321)
+	      #pragma omp ordered depend (sink: i, j, k + 2 E(m)) \
+				  depend (sink:i - 2, j + 1, k - 4 E(m)) \
+				  depend(sink: i - 1, j - 2, k - 2 E(m))
+	      if (k <= 4)
+		{
+		  l = c[i][j][k + 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      if (i >= 2 && j < 7 && k >= 4)
+		{
+		  l = c[i - 2][j + 1][k - 4];
+		  if (l < 2)
+		    abort ();
+		}
+	      if (i >= 1 && j >= 4 && k >= 2)
+		{
+		  l = c[i - 1][j - 2][k - 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp ordered depend (source)
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 0; i < d + 1; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 1 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 0; i < d + 1; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, m)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 1 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 0; i < d + 1; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3)
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 2 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (j >= 2 && (k & 1) == 0))
+	    abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/doacross-2.c b/libgomp/testsuite/libgomp.c/doacross-2.c
new file mode 100644
index 00000000000..e491bb22965
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-2.c
@@ -0,0 +1,225 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8], g[N / 16][8][6];
+volatile int d, e;
+volatile unsigned long long f;
+
+int
+main ()
+{
+  unsigned long long i;
+  int j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(static, 1) ordered (1) nowait
+    for (i = 1; i < N + f; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered depend(sink: i - 1)
+	if (i > 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered depend(source)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(static, 0) ordered (3) nowait
+    for (i = 3; i < N / 16 - 1 + f; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered depend(sink: i, j - 2, k - 1) \
+				depend(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered depend(sink: i - 3, j + 2, k - 2)
+	    if (j >= 2 && k > 1)
+	      {
+		#pragma omp atomic read
+		l = b[i][j - 2][k - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 5 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i >= 6 && j < N / 16 - 3 && k == 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 3][j + 2][k - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered depend(source)
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15)
+    for (i = 2; i < N / 32 + f; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+#define A(n) ,n
+#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321)
+	      #pragma omp ordered depend (sink: i, j, k + 2 E(m)) \
+				  depend (sink:i - 2, j + 1, k - 4 E(m)) \
+				  depend(sink: i - 1, j - 2, k - 2 E(m))
+	      if (k <= 4)
+		{
+		  l = c[i][j][k + 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      if (i >= 4 && j < 7 && k >= 4)
+		{
+		  l = c[i - 2][j + 1][k - 4];
+		  if (l < 2)
+		    abort ();
+		}
+	      if (i >= 3 && j >= 4 && k >= 2)
+		{
+		  l = c[i - 1][j - 2][k - 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp ordered depend (source)
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+    #pragma omp for schedule(static, 0) ordered (3) nowait
+    for (j = 0; j < N / 16 - 1; j++)
+      for (k = 0; k < 8; k += 2)
+	for (i = 3; i <= 5 + f; i++)
+	  {
+	    #pragma omp atomic write
+	    g[j][k][i] = 1;
+	    #pragma omp ordered depend(sink: j, k - 2, i - 1) \
+				depend(sink: j - 2, k - 2, i + 1)
+	    #pragma omp ordered depend(sink: j - 3, k + 2, i - 2)
+	    if (k >= 2 && i > 3)
+	      {
+		#pragma omp atomic read
+		l = g[j][k - 2][i - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    g[j][k][i] = 2;
+	    if (j >= 2 && k >= 2 && i < 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 2][k - 2][i + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (j >= 3 && k < N / 16 - 3 && i == 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 3][k + 2][i - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered depend(source)
+	    #pragma omp atomic write
+	    g[j][k][i] = 3;
+	  }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 3 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, m)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 3 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 2; i < f + 3; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered depend (source)
+	      #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3 * (i >= 1))
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 3 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (i >= 2 && j >= 2 && (k & 1) == 0))
+	    abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 6; k++)
+	  if (g[i][j][k] != 3 * (i < N / 16 - 1 && (j & 1) == 0 && k >= 3))
+	    abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c
index beca8555780..6d4bc4fac12 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c
@@ -20,7 +20,7 @@ int fib_wrapper (int n)
 {
   int x = 0;
 
-  #pragma omp target if(n > THRESHOLD)
+  #pragma omp target if(n > THRESHOLD) map(from:x)
     x = fib (n);
 
   return x;
diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c
index db70460b309..f2414366951 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c
@@ -41,7 +41,7 @@ float accum (int k)
   int i;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	tmp += Pfun (i, k);
diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c
index b550f1ff540..33d6137afd5 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c
@@ -48,7 +48,7 @@ float accum ()
   int i, k;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	{
diff --git a/libgomp/testsuite/libgomp.c/examples-4/device-1.c b/libgomp/testsuite/libgomp.c/examples-4/device-1.c
index f7c84fb4c14..dad8572f8f0 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/device-1.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/device-1.c
@@ -10,11 +10,11 @@ int main ()
   int b = 0;
   int c, d;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -26,11 +26,11 @@ int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -42,11 +42,11 @@ int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
diff --git a/libgomp/testsuite/libgomp.c/examples-4/device-3.c b/libgomp/testsuite/libgomp.c/examples-4/device-3.c
index 8a0cf7c200d..af086533278 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/device-3.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/device-3.c
@@ -9,7 +9,7 @@ int main ()
   int res;
   int default_device = omp_get_default_device ();
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (res)
@@ -17,7 +17,7 @@ int main ()
 
   omp_set_default_device (omp_get_num_devices ());
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (!res)
diff --git a/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c b/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c
index abb283801f8..46b674013d0 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c
@@ -47,7 +47,7 @@ void gramSchmidt (int Q[][COLS], const int rows, const int cols)
       {
 	int tmp = 0;
 
-	#pragma omp target
+	#pragma omp target map(tofrom:tmp)
 	  #pragma omp parallel for reduction(+:tmp)
 	    for (i = 0; i < rows; i++)
 	      tmp += (Q[i][k] * Q[i][k]);
diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-2.c b/libgomp/testsuite/libgomp.c/examples-4/teams-2.c
index 8bbbc355b17..7d0a60ebb51 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/teams-2.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/teams-2.c
@@ -32,7 +32,7 @@ float dotprod (float B[], float C[], int n, int block_size,
   int i, i0;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp teams num_teams(num_teams) thread_limit(block_threads) \
 		      reduction(+:sum)
       #pragma omp distribute
diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-3.c b/libgomp/testsuite/libgomp.c/examples-4/teams-3.c
index b6708785884..5fe63a68a4b 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/teams-3.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/teams-3.c
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int n)
   int i;
   float sum = 0;
 
-  #pragma omp target teams map(to: B[0:n], C[0:n])
+  #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp distribute parallel for reduction(+:sum)
       for (i = 0; i < n; i++)
 	sum += B[i] * C[i];
diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-4.c b/libgomp/testsuite/libgomp.c/examples-4/teams-4.c
index 9aef78ecfba..6136eabef66 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/teams-4.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/teams-4.c
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int n)
   int i;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum)
     #pragma omp teams num_teams(8) thread_limit(16)
       #pragma omp distribute parallel for reduction(+:sum) \
 					  dist_schedule(static, 1024) \
diff --git a/libgomp/testsuite/libgomp.c/for-2.h b/libgomp/testsuite/libgomp.c/for-2.h
index 920d23b5202..0bd116c5aec 100644
--- a/libgomp/testsuite/libgomp.c/for-2.h
+++ b/libgomp/testsuite/libgomp.c/for-2.h
@@ -11,11 +11,21 @@ noreturn (void)
 #ifndef SC
 #define SC
 #endif
+#ifndef OMPTGT
+#define OMPTGT
+#endif
+#ifndef OMPTO
+#define OMPTO(v) do {} while (0)
+#endif
+#ifndef OMPFROM
+#define OMPFROM(v) do {} while (0)
+#endif
 
 __attribute__((noinline, noclone)) void
 N(f0) (void)
 {
   int i;
+  OMPTGT
 #pragma omp F S
   for (i = 0; i < 1500; i++)
     a[i] += 2;
@@ -24,6 +34,7 @@ N(f0) (void)
 __attribute__((noinline, noclone)) void
 N(f1) (void)
 {
+  OMPTGT
 #pragma omp F S
   for (unsigned int i = __INT_MAX__; i < 3000U + __INT_MAX__; i += 2)
     a[(i - __INT_MAX__) >> 1] -= 2;
@@ -33,6 +44,7 @@ __attribute__((noinline, noclone)) void
 N(f2) (void)
 {
   unsigned long long i;
+  OMPTGT
 #pragma omp F S
   for (i = __LONG_LONG_MAX__ + 4500ULL - 27;
        i > __LONG_LONG_MAX__ - 27ULL; i -= 3)
@@ -42,6 +54,7 @@ N(f2) (void)
 __attribute__((noinline, noclone)) void
 N(f3) (long long n1, long long n2, long long s3)
 {
+  OMPTGT
 #pragma omp F S
   for (long long i = n1 + 23; i > n2 - 25; i -= s3)
     a[i + 48] += 7;
@@ -51,6 +64,7 @@ __attribute__((noinline, noclone)) void
 N(f4) (void)
 {
   unsigned int i;
+  OMPTGT
 #pragma omp F S
   for (i = 30; i < 20; i += 2)
     a[i] += 10;
@@ -61,6 +75,7 @@ N(f5) (int n11, int n12, int n21, int n22, int n31, int n32,
        int s1, int s2, int s3)
 {
   SC int v1, v2, v3;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (v1 = n11; v1 < n12; v1 += s1)
     for (v2 = n21; v2 < n22; v2 += s2)
@@ -74,6 +89,7 @@ N(f6) (int n11, int n12, int n21, int n22, long long n31, long long n32,
 {
   SC int v1, v2;
   SC long long v3;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (v1 = n11; v1 > n12; v1 += s1)
     for (v2 = n21; v2 > n22; v2 += s2)
@@ -86,6 +102,7 @@ N(f7) (void)
 {
   SC unsigned int v1, v3;
   SC unsigned long long v2;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (v1 = 0; v1 < 20; v1 += 2)
     for (v2 = __LONG_LONG_MAX__ + 16ULL;
@@ -98,6 +115,7 @@ __attribute__((noinline, noclone)) void
 N(f8) (void)
 {
   SC long long v1, v2, v3;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (v1 = 0; v1 < 20; v1 += 2)
     for (v2 = 30; v2 < 20; v2++)
@@ -109,6 +127,7 @@ __attribute__((noinline, noclone)) void
 N(f9) (void)
 {
   int i;
+  OMPTGT
 #pragma omp F S
   for (i = 20; i < 10; i++)
     {
@@ -122,6 +141,7 @@ __attribute__((noinline, noclone)) void
 N(f10) (void)
 {
   SC int i;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (i = 0; i < 10; i++)
     for (int j = 10; j < 8; j++)
@@ -137,6 +157,7 @@ __attribute__((noinline, noclone)) void
 N(f11) (int n)
 {
   int i;
+  OMPTGT
 #pragma omp F S
   for (i = 20; i < n; i++)
     {
@@ -150,6 +171,7 @@ __attribute__((noinline, noclone)) void
 N(f12) (int n)
 {
   SC int i;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (i = 0; i < 10; i++)
     for (int j = n; j < 8; j++)
@@ -165,6 +187,7 @@ __attribute__((noinline, noclone)) void
 N(f13) (void)
 {
   int *i;
+  OMPTGT
 #pragma omp F S
   for (i = a; i < &a[1500]; i++)
     i[0] += 2;
@@ -174,6 +197,7 @@ __attribute__((noinline, noclone)) void
 N(f14) (void)
 {
   SC float *i;
+  OMPTGT
 #pragma omp F S collapse(3)
   for (i = &b[0][0][0]; i < &b[0][0][10]; i++)
     for (float *j = &b[0][15][0]; j > &b[0][0][0]; j -= 10)
@@ -188,27 +212,34 @@ N(test) (void)
   int i, j, k;
   for (i = 0; i < 1500; i++)
     a[i] = i - 25;
+  OMPTO (a);
   N(f0) ();
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 23)
       return 1;
   N(f1) ();
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 25)
       return 1;
   N(f2) ();
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 29)
       return 1;
   N(f3) (1500LL - 1 - 23 - 48, -1LL + 25 - 48, 1LL);
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 22)
       return 1;
   N(f3) (1500LL - 1 - 23 - 48, 1500LL - 1, 7LL);
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 22)
       return 1;
   N(f4) ();
+  OMPFROM (a);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 22)
       return 1;
@@ -216,31 +247,37 @@ N(test) (void)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
 	b[i][j][k] = i - 2.5 + 1.5 * j - 1.5 * k;
+  OMPTO (b);
   N(f5) (0, 10, 0, 15, 0, 10, 1, 1, 1);
+  OMPFROM (b);
   for (i = 0; i < 10; i++)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
 	if (b[i][j][k] != i + 1.5 * j - 1.5 * k)
 	  return 1;
   N(f5) (0, 10, 30, 15, 0, 10, 4, 5, 6);
+  OMPFROM (b);
   for (i = 0; i < 10; i++)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
 	if (b[i][j][k] != i + 1.5 * j - 1.5 * k)
 	  return 1;
   N(f6) (9, -1, 29, 0, 9, -1, -1, -2, -1);
+  OMPFROM (b);
   for (i = 0; i < 10; i++)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
 	if (b[i][j][k] != i - 4.5 + 1.5 * j - 1.5 * k)
 	  return 1;
   N(f7) ();
+  OMPFROM (b);
   for (i = 0; i < 10; i++)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
 	if (b[i][j][k] != i + 1.0 + 1.5 * j - 1.5 * k)
 	  return 1;
   N(f8) ();	  
+  OMPFROM (b);
   for (i = 0; i < 10; i++)
     for (j = 0; j < 15; j++)
       for (k = 0; k < 10; k++)
@@ -250,6 +287,8 @@ N(test) (void)
   N(f10) ();
   N(f11) (10);
   N(f12) (12);
+  OMPFROM (a);
+  OMPFROM (b);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 22)
       return 1;
@@ -260,6 +299,8 @@ N(test) (void)
 	  return 1;
   N(f13) ();
   N(f14) ();
+  OMPFROM (a);
+  OMPFROM (b);
   for (i = 0; i < 1500; i++)
     if (a[i] != i - 20)
       return 1;
diff --git a/libgomp/testsuite/libgomp.c/for-4.c b/libgomp/testsuite/libgomp.c/for-4.c
new file mode 100644
index 00000000000..ef5465e1e76
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/for-4.c
@@ -0,0 +1,42 @@
+/* { dg-options "-std=gnu99 -fopenmp" } */
+
+extern void abort (void);
+
+#define M(x, y, z) O(x, y, z)
+#define O(x, y, z) x ## _ ## y ## _ ## z
+
+#define F taskloop
+#define G taskloop
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F taskloop simd
+#define G taskloop_simd
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+int
+main ()
+{
+  int err = 0;
+  #pragma omp parallel reduction(|:err)
+    #pragma omp single
+      {
+	if (test_taskloop_normal ()
+	    || test_taskloop_simd_normal ())
+	  err = 1;
+      }
+  if (err)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/for-5.c b/libgomp/testsuite/libgomp.c/for-5.c
new file mode 100644
index 00000000000..84e636ab0f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/for-5.c
@@ -0,0 +1,154 @@
+/* { dg-additional-options "-std=gnu99" } */
+
+extern void abort ();
+
+#define M(x, y, z) O(x, y, z)
+#define O(x, y, z) x ## _ ## y ## _ ## z
+
+#pragma omp declare target
+
+#define F for
+#define G f
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#pragma omp end declare target
+
+#undef OMPFROM
+#undef OMPTO
+#define DO_PRAGMA(x) _Pragma (#x)
+#define OMPFROM(v) DO_PRAGMA (omp target update from(v))
+#define OMPTO(v) DO_PRAGMA (omp target update to(v))
+
+#define F target parallel for
+#define G tpf
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F target simd
+#define G t_simd
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F target parallel for simd
+#define G tpf_simd
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F target teams distribute
+#define G ttd
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F target teams distribute
+#define G ttd_ds128
+#define S dist_schedule(static, 128)
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F target teams distribute simd
+#define G ttds
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F target teams distribute simd
+#define G ttds_ds128
+#define S dist_schedule(static, 128)
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F target teams distribute parallel for
+#define G ttdpf
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F target teams distribute parallel for dist_schedule(static, 128)
+#define G ttdpf_ds128
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F target teams distribute parallel for simd
+#define G ttdpfs
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F target teams distribute parallel for simd dist_schedule(static, 128)
+#define G ttdpfs_ds128
+#include "for-1.h"
+#undef F
+#undef G
+
+int
+main ()
+{
+  if (test_tpf_static ()
+      || test_tpf_static32 ()
+      || test_tpf_auto ()
+      || test_tpf_guided32 ()
+      || test_tpf_runtime ()
+      || test_t_simd_normal ()
+      || test_tpf_simd_static ()
+      || test_tpf_simd_static32 ()
+      || test_tpf_simd_auto ()
+      || test_tpf_simd_guided32 ()
+      || test_tpf_simd_runtime ()
+      || test_ttd_normal ()
+      || test_ttd_ds128_normal ()
+      || test_ttds_normal ()
+      || test_ttds_ds128_normal ()
+      || test_ttdpf_static ()
+      || test_ttdpf_static32 ()
+      || test_ttdpf_auto ()
+      || test_ttdpf_guided32 ()
+      || test_ttdpf_runtime ()
+      || test_ttdpf_ds128_static ()
+      || test_ttdpf_ds128_static32 ()
+      || test_ttdpf_ds128_auto ()
+      || test_ttdpf_ds128_guided32 ()
+      || test_ttdpf_ds128_runtime ()
+      || test_ttdpfs_static ()
+      || test_ttdpfs_static32 ()
+      || test_ttdpfs_auto ()
+      || test_ttdpfs_guided32 ()
+      || test_ttdpfs_runtime ()
+      || test_ttdpfs_ds128_static ()
+      || test_ttdpfs_ds128_static32 ()
+      || test_ttdpfs_ds128_auto ()
+      || test_ttdpfs_ds128_guided32 ()
+      || test_ttdpfs_ds128_runtime ())
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/for-6.c b/libgomp/testsuite/libgomp.c/for-6.c
new file mode 100644
index 00000000000..7f3c65e82b1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/for-6.c
@@ -0,0 +1,123 @@
+/* { dg-additional-options "-std=gnu99" } */
+
+extern void abort ();
+
+#define M(x, y, z) O(x, y, z)
+#define O(x, y, z) x ## _ ## y ## _ ## z
+
+#pragma omp declare target
+
+#define F for
+#define G f
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#pragma omp end declare target
+
+#undef OMPTGT
+#undef OMPFROM
+#undef OMPTO
+#define DO_PRAGMA(x) _Pragma (#x)
+#define OMPTGT DO_PRAGMA (omp target)
+#define OMPFROM(v) DO_PRAGMA (omp target update from(v))
+#define OMPTO(v) DO_PRAGMA (omp target update to(v))
+
+#define F teams distribute
+#define G td
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F teams distribute
+#define G td_ds128
+#define S dist_schedule(static, 128)
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F teams distribute simd
+#define G tds
+#define S
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F teams distribute simd
+#define G tds_ds128
+#define S dist_schedule(static, 128)
+#define N(x) M(x, G, normal)
+#include "for-2.h"
+#undef S
+#undef N
+#undef F
+#undef G
+
+#define F teams distribute parallel for
+#define G tdpf
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F teams distribute parallel for dist_schedule(static, 128)
+#define G tdpf_ds128
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F teams distribute parallel for simd
+#define G tdpfs
+#include "for-1.h"
+#undef F
+#undef G
+
+#define F teams distribute parallel for simd dist_schedule(static, 128)
+#define G tdpfs_ds128
+#include "for-1.h"
+#undef F
+#undef G
+
+int
+main ()
+{
+  if (test_td_normal ()
+      || test_td_ds128_normal ()
+      || test_tds_normal ()
+      || test_tds_ds128_normal ()
+      || test_tdpf_static ()
+      || test_tdpf_static32 ()
+      || test_tdpf_auto ()
+      || test_tdpf_guided32 ()
+      || test_tdpf_runtime ()
+      || test_tdpf_ds128_static ()
+      || test_tdpf_ds128_static32 ()
+      || test_tdpf_ds128_auto ()
+      || test_tdpf_ds128_guided32 ()
+      || test_tdpf_ds128_runtime ()
+      || test_tdpfs_static ()
+      || test_tdpfs_static32 ()
+      || test_tdpfs_auto ()
+      || test_tdpfs_guided32 ()
+      || test_tdpfs_runtime ()
+      || test_tdpfs_ds128_static ()
+      || test_tdpfs_ds128_static32 ()
+      || test_tdpfs_ds128_auto ()
+      || test_tdpfs_ds128_guided32 ()
+      || test_tdpfs_ds128_runtime ())
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/linear-1.c b/libgomp/testsuite/libgomp.c/linear-1.c
new file mode 100644
index 00000000000..f86fb33c5da
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/linear-1.c
@@ -0,0 +1,250 @@
+int a[256];
+
+__attribute__((noinline, noclone)) int
+f1 (int i)
+{
+  #pragma omp parallel for linear (i: 4)
+  for (int j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f2 (short int i, char k)
+{
+  #pragma omp parallel for linear (i: k + 1)
+  for (long j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f3 (long long int i, long long int k)
+{
+  #pragma omp parallel for linear (i: k)
+  for (short j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) int
+f4 (int i)
+{
+  #pragma omp parallel for linear (i: 4) schedule(static, 3)
+  for (int j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f5 (short int i, char k)
+{
+  #pragma omp parallel for linear (i: k + 1) schedule(static, 5)
+  for (long j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f6 (long long int i, long long int k)
+{
+  #pragma omp parallel for linear (i: k) schedule(static, 7)
+  for (short j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) int
+f7 (int i)
+{
+  #pragma omp parallel for linear (i: 4) schedule(dynamic, 3)
+  for (int j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f8 (short int i, char k)
+{
+  #pragma omp parallel for linear (i: k + 1) schedule(dynamic, 5)
+  for (long j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f9 (long long int i, long long int k)
+{
+  #pragma omp parallel for linear (i: k) schedule(dynamic, 7)
+  for (short j = 16; j < 64; j++)
+    {
+      a[i] = j;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) int
+f10 (int i, long step)
+{
+  #pragma omp parallel for linear (i: 4)
+  for (int j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f11 (short int i, char k, char step)
+{
+  #pragma omp parallel for linear (i: k + 1)
+  for (long j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f12 (long long int i, long long int k, int step)
+{
+  #pragma omp parallel for linear (i: k)
+  for (short j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) int
+f13 (int i, long long int step)
+{
+  #pragma omp parallel for linear (i: 4) schedule(static, 3)
+  for (int j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f14 (short int i, char k, int step)
+{
+  #pragma omp parallel for linear (i: k + 1) schedule(static, 5)
+  for (long j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f15 (long long int i, long long int k, long int step)
+{
+  #pragma omp parallel for linear (i: k) schedule(static, 7)
+  for (short j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) int
+f16 (int i, long long int step)
+{
+  #pragma omp parallel for linear (i: 4) schedule(dynamic, 3)
+  for (int j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) short int
+f17 (short int i, char k, int step)
+{
+  #pragma omp parallel for linear (i: k + 1) schedule(dynamic, 5)
+  for (long j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+__attribute__((noinline, noclone)) long long int
+f18 (long long int i, long long int k, long int step)
+{
+  #pragma omp parallel for linear (i: k) schedule(dynamic, 7)
+  for (short j = 16; j < 112; j += step)
+    {
+      a[i] = j / 2 + 8;
+      i += 4;
+    }
+  return i;
+}
+
+int
+main ()
+{
+#define TEST(x) \
+  if (x != 8 + 48 * 4)				\
+    __builtin_abort ();				\
+  for (int i = 0; i < 256; i++)			\
+    if (a[i] != (((i & 3) == 0 && i >= 8	\
+		  && i < 8 + 48 * 4)		\
+		 ? ((i - 8) / 4) + 16 : 0))	\
+      __builtin_abort ();			\
+  __builtin_memset (a, 0, sizeof (a))
+  TEST (f1 (8));
+  TEST (f2 (8, 3));
+  TEST (f3 (8LL, 4LL));
+  TEST (f4 (8));
+  TEST (f5 (8, 3));
+  TEST (f6 (8LL, 4LL));
+  TEST (f7 (8));
+  TEST (f8 (8, 3));
+  TEST (f9 (8LL, 4LL));
+  TEST (f10 (8, 2));
+  TEST (f11 (8, 3, 2));
+  TEST (f12 (8LL, 4LL, 2));
+  TEST (f13 (8, 2));
+  TEST (f14 (8, 3, 2));
+  TEST (f15 (8LL, 4LL, 2));
+  TEST (f16 (8, 2));
+  TEST (f17 (8, 3, 2));
+  TEST (f18 (8LL, 4LL, 2));
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/ordered-4.c b/libgomp/testsuite/libgomp.c/ordered-4.c
new file mode 100644
index 00000000000..8412d4715c3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/ordered-4.c
@@ -0,0 +1,83 @@
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+foo (int i, char *j)
+{
+  #pragma omp atomic
+  j[i]++;
+  #pragma omp ordered threads
+  {
+    int t;
+    #pragma omp atomic read
+    t = j[i];
+    if (t != 3)
+      abort ();
+    if (i > 1)
+      {
+	#pragma omp atomic read
+	t = j[i - 1];
+	if (t == 2)
+	  abort ();
+      }
+    if (i < 127)
+      {
+	#pragma omp atomic read
+	t = j[i + 1];
+	if (t == 4)
+	  abort ();
+      }
+  }
+  #pragma omp atomic
+  j[i]++;
+}
+
+int
+main ()
+{
+  int i;
+  char j[128];
+  #pragma omp parallel
+  {
+    #pragma omp for
+    for (i = 0; i < 128; i++)
+      j[i] = 0;
+    #pragma omp for ordered schedule(dynamic, 1)
+    for (i = 0; i < 128; i++)
+      {
+	#pragma omp atomic
+	j[i]++;
+	#pragma omp ordered threads
+	{
+	  int t;
+	  #pragma omp atomic read
+	  t = j[i];
+	  if (t != 1)
+	    abort ();
+	  if (i > 1)
+	    {
+	      #pragma omp atomic read
+	      t = j[i - 1];
+	      if (t == 0)
+		abort ();
+	    }
+	  if (i < 127)
+	    {
+	      #pragma omp atomic read
+	      t = j[i + 1];
+	      if (t == 2)
+		abort ();
+	    }
+	}
+	#pragma omp atomic
+	j[i]++;
+      }
+    #pragma omp for ordered schedule(static, 1)
+    for (i = 0; i < 128; i++)
+      foo (i, j);
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-2.c b/libgomp/testsuite/libgomp.c/pr66199-2.c
index ddb79de8943..2fc9eec529a 100644
--- a/libgomp/testsuite/libgomp.c/pr66199-2.c
+++ b/libgomp/testsuite/libgomp.c/pr66199-2.c
@@ -18,12 +18,11 @@ __attribute__((noinline, noclone)) void
 f2 (long a, long b, long c)
 {
   long d, e;
-  #pragma omp target teams distribute parallel for simd default(none) firstprivate (a, b) shared(u, v, w) linear(d) linear(c:5) lastprivate(e)
+  #pragma omp target teams distribute parallel for simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
   for (d = a; d < b; d++)
     {
       u[d] = v[d] + w[d];
-      c += 5;
-      e = c;
+      e = c + d * 5;
     }
 }
 
diff --git a/libgomp/testsuite/libgomp.c/pr66199-3.c b/libgomp/testsuite/libgomp.c/pr66199-3.c
new file mode 100644
index 00000000000..fe0ccb47197
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-3.c
@@ -0,0 +1,50 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp" } */
+
+int u[1024], v[1024], w[1024];
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp parallel for lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w)
+  for (d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+  return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp parallel for lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w) linear(c:5) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      c += 5;
+      e = c;
+    }
+  return d + c + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + 2 * (17 + 5 * 1024)
+      || f3 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-4.c b/libgomp/testsuite/libgomp.c/pr66199-4.c
new file mode 100644
index 00000000000..a9b1bb8a59e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-4.c
@@ -0,0 +1,58 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp" } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) void
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target teams distribute parallel for default(none) firstprivate (a, b) shared(u, v, w)
+  for (d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+}
+
+__attribute__((noinline, noclone)) void
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target teams distribute parallel for default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+}
+
+__attribute__((noinline, noclone)) void
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+}
+
+__attribute__((noinline, noclone)) void
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+}
+
+int
+main ()
+{
+  f1 (0, 1024);
+  f2 (0, 1024, 17);
+  f3 (0, 32, 0, 32);
+  f4 (0, 32, 0, 32);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/reduction-10.c b/libgomp/testsuite/libgomp.c/reduction-10.c
new file mode 100644
index 00000000000..3c95ebd4a4b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-10.c
@@ -0,0 +1,105 @@
+struct A { int t; };
+struct B { char t; };
+struct C { unsigned long long t; };
+struct D { long t; };
+void
+add (struct B *x, struct B *y)
+{
+  x->t += y->t;
+}
+void
+zero (struct B *x)
+{
+  x->t = 0;
+}
+void
+orit (struct C *x, struct C *y)
+{
+  y->t |= x->t;
+}
+#pragma omp declare reduction(+:struct A:omp_out.t += omp_in.t)
+#pragma omp declare reduction(+:struct B:add (&omp_out, &omp_in)) initializer(zero (&omp_priv))
+#pragma omp declare reduction(*:struct A:omp_out.t *= omp_in.t) initializer(omp_priv = { 1 })
+#pragma omp declare reduction(|:struct C:orit (&omp_in, &omp_out))
+#pragma omp declare reduction(&:struct D:omp_out.t = omp_out.t & omp_in.t) initializer(omp_priv = { ~0L })
+#pragma omp declare reduction(maxb:short:omp_out = omp_in > omp_out ? omp_in : omp_out) initializer(omp_priv = -6)
+
+struct B z[10];
+
+__attribute__((noinline, noclone)) void
+foo (struct A (*x)[3][2], struct A *y, struct D w[1][2], int p1, long p2, long p3, int p4,
+     int p5, long p6, short p7)
+{
+  struct C a[p7 + 4];
+  short b[p7];
+  int i;
+  for (i = 0; i < p7 + 4; i++)
+    {
+      if (i < p7)
+	b[i] = -6;
+      a[i].t = 0;
+    }
+  #pragma omp parallel for reduction(+:x[0:p1 + 1][:p2], z[:p3]) \
+			   reduction(*:y[:p4]) reduction(|:a[:p5]) \
+			   reduction(&:w[0:p6 - 1][:p6]) reduction(maxb:b)
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[0].t *= 3;
+      if ((i & 31) == 2)
+	y[1].t *= 7;
+      if ((i & 63) == 3)
+	y[2].t *= 17;
+      z[i / 32].t += (i & 3);
+      if (i < 4)
+	z[i].t += i;
+      a[i / 32].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      if ((i % 79) > b[0])
+	b[0] = i % 79;
+      if ((i % 13) > b[1])
+	b[1] = i % 13;
+      if ((i % 23) > b[2])
+	b[2] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+  for (i = 0; i < 9; i++)
+    if (a[i].t != (i < 4 ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  struct A a[4][3][2] = {};
+  static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }},
+			    {{ 312, 381 }, { 295, 356 }, { 337, 335 }},
+			    {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }},
+			    {{ 0, 0 }, { 0, 0 }, { 0, 0 }}};
+  struct A y[5] = { { 0 }, { 1 }, { 1 }, { 1 }, { 0 } };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 };
+  struct D w[1][2] = { { { ~0L }, { ~0L } } };
+  foo (&a[1], y + 1, w, 1, 3L, 4L, 3, 4, 2L, 5);
+  int i, j, k;
+  for (i = 0; i < 4; i++)
+    for (j = 0; j < 3; j++)
+      for (k = 0; k < 2; k++)
+	if (a[i][j][k].t != a2[i][j][k])
+	  __builtin_abort ();
+  for (i = 0; i < 5; i++)
+    if (y[i].t != y2[i])
+      __builtin_abort ();
+  for (i = 0; i < 10; i++)
+    if (z[i].t != z2[i])
+      __builtin_abort ();
+  if (w[0][0].t != ~0x249249L || w[0][1].t != ~0x249249L)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/reduction-7.c b/libgomp/testsuite/libgomp.c/reduction-7.c
new file mode 100644
index 00000000000..347c26f46d3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-7.c
@@ -0,0 +1,64 @@
+char z[10] = { 0 };
+
+__attribute__((noinline, noclone)) void
+foo (int (*x)[3][2], int *y, long w[1][2])
+{
+  unsigned long long a[9] = {};
+  short b[5] = {};
+  int i;
+  #pragma omp parallel for reduction(+:x[0:2][:][0:2], z[:4]) \
+			   reduction(*:y[:3]) reduction(|:a[:4]) \
+			   reduction(&:w[0:1][:2]) reduction(max:b)
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[0] *= 3;
+      if ((i & 31) == 2)
+	y[1] *= 7;
+      if ((i & 63) == 3)
+	y[2] *= 17;
+      z[i / 32] += (i & 3);
+      if (i < 4)
+	z[i] += i;
+      a[i / 32] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      if ((i % 79) > b[0])
+	b[0] = i % 79;
+      if ((i % 13) > b[1])
+	b[1] = i % 13;
+      if ((i % 23) > b[2])
+	b[2] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+  for (i = 0; i < 9; i++)
+    if (a[i] != (i < 4 ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  int a[4][3][2] = {};
+  static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }},
+			    {{ 312, 381 }, { 295, 356 }, { 337, 335 }},
+			    {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }},
+			    {{ 0, 0 }, { 0, 0 }, { 0, 0 }}};
+  int y[5] = { 0, 1, 1, 1, 0 };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 };
+  long w[1][2] = { ~0L, ~0L };
+  foo (&a[1], y + 1, w);
+  if (__builtin_memcmp (a, a2, sizeof (a))
+      || __builtin_memcmp (y, y2, sizeof (y))
+      || __builtin_memcmp (z, z2, sizeof (z))
+      || w[0][0] != ~0x249249L
+      || w[0][1] != ~0x249249L)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/reduction-8.c b/libgomp/testsuite/libgomp.c/reduction-8.c
new file mode 100644
index 00000000000..f4ec03aabea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-8.c
@@ -0,0 +1,98 @@
+struct A { int t; };
+struct B { char t; };
+struct C { unsigned long long t; };
+struct D { long t; };
+void
+add (struct B *x, struct B *y)
+{
+  x->t += y->t;
+}
+void
+zero (struct B *x)
+{
+  x->t = 0;
+}
+void
+orit (struct C *x, struct C *y)
+{
+  y->t |= x->t;
+}
+#pragma omp declare reduction(+:struct A:omp_out.t += omp_in.t)
+#pragma omp declare reduction(+:struct B:add (&omp_out, &omp_in)) initializer(zero (&omp_priv))
+#pragma omp declare reduction(*:struct A:omp_out.t *= omp_in.t) initializer(omp_priv = { 1 })
+#pragma omp declare reduction(|:struct C:orit (&omp_in, &omp_out))
+#pragma omp declare reduction(&:struct D:omp_out.t = omp_out.t & omp_in.t) initializer(omp_priv = { ~0L })
+#pragma omp declare reduction(maxb:short:omp_out = omp_in > omp_out ? omp_in : omp_out) initializer(omp_priv = -6)
+
+struct B z[10];
+
+__attribute__((noinline, noclone)) void
+foo (struct A (*x)[3][2], struct A *y, struct D w[1][2])
+{
+  struct C a[9] = {};
+  short b[5] = {};
+  int i;
+  #pragma omp parallel for reduction(+:x[0:2][:][0:2], z[:4]) \
+			   reduction(*:y[:3]) reduction(|:a[:4]) \
+			   reduction(&:w[0:1][:2]) reduction(maxb:b)
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[0].t *= 3;
+      if ((i & 31) == 2)
+	y[1].t *= 7;
+      if ((i & 63) == 3)
+	y[2].t *= 17;
+      z[i / 32].t += (i & 3);
+      if (i < 4)
+	z[i].t += i;
+      a[i / 32].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      if ((i % 79) > b[0])
+	b[0] = i % 79;
+      if ((i % 13) > b[1])
+	b[1] = i % 13;
+      if ((i % 23) > b[2])
+	b[2] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+  for (i = 0; i < 9; i++)
+    if (a[i].t != (i < 4 ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  struct A a[4][3][2] = {};
+  static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }},
+			    {{ 312, 381 }, { 295, 356 }, { 337, 335 }},
+			    {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }},
+			    {{ 0, 0 }, { 0, 0 }, { 0, 0 }}};
+  struct A y[5] = { { 0 }, { 1 }, { 1 }, { 1 }, { 0 } };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 };
+  struct D w[1][2] = { { { ~0L }, { ~0L } } };
+  foo (&a[1], y + 1, w);
+  int i, j, k;
+  for (i = 0; i < 4; i++)
+    for (j = 0; j < 3; j++)
+      for (k = 0; k < 2; k++)
+	if (a[i][j][k].t != a2[i][j][k])
+	  __builtin_abort ();
+  for (i = 0; i < 5; i++)
+    if (y[i].t != y2[i])
+      __builtin_abort ();
+  for (i = 0; i < 10; i++)
+    if (z[i].t != z2[i])
+      __builtin_abort ();
+  if (w[0][0].t != ~0x249249L || w[0][1].t != ~0x249249L)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/reduction-9.c b/libgomp/testsuite/libgomp.c/reduction-9.c
new file mode 100644
index 00000000000..13605c1ab88
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-9.c
@@ -0,0 +1,71 @@
+char z[10] = { 0 };
+
+__attribute__((noinline, noclone)) void
+foo (int (*x)[3][2], int *y, long w[1][2], int p1, long p2, long p3, int p4,
+     int p5, long p6, short p7)
+{
+  unsigned long long a[p7 + 4];
+  short b[p7];
+  int i;
+  for (i = 0; i < p7 + 4; i++)
+    {
+      if (i < p7)
+	b[i] = -6;
+      a[i] = 0;
+    }
+  #pragma omp parallel for reduction(+:x[0:p1 + 1][:p2], z[:p3]) \
+			   reduction(*:y[:p4]) reduction(|:a[:p5]) \
+			   reduction(&:w[0:p6 - 1][:p6]) reduction(max:b)
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[0] *= 3;
+      if ((i & 31) == 2)
+	y[1] *= 7;
+      if ((i & 63) == 3)
+	y[2] *= 17;
+      z[i / 32] += (i & 3);
+      if (i < 4)
+	z[i] += i;
+      a[i / 32] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      if ((i % 79) > b[0])
+	b[0] = i % 79;
+      if ((i % 13) > b[1])
+	b[1] = i % 13;
+      if ((i % 23) > b[2])
+	b[2] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+  for (i = 0; i < 9; i++)
+    if (a[i] != (i < 4 ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  int a[4][3][2] = {};
+  static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }},
+			    {{ 312, 381 }, { 295, 356 }, { 337, 335 }},
+			    {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }},
+			    {{ 0, 0 }, { 0, 0 }, { 0, 0 }}};
+  int y[5] = { 0, 1, 1, 1, 0 };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 };
+  long w[1][2] = { ~0L, ~0L };
+  foo (&a[1], y + 1, w, 1, 3L, 4L, 3, 4, 2L, 5);
+  if (__builtin_memcmp (a, a2, sizeof (a))
+      || __builtin_memcmp (y, y2, sizeof (y))
+      || __builtin_memcmp (z, z2, sizeof (z))
+      || w[0][0] != ~0x249249L
+      || w[0][1] != ~0x249249L)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-1.c b/libgomp/testsuite/libgomp.c/target-1.c
index f734d3c279d..c7abb008182 100644
--- a/libgomp/testsuite/libgomp.c/target-1.c
+++ b/libgomp/testsuite/libgomp.c/target-1.c
@@ -34,7 +34,7 @@ fn2 (int x, int y, int z)
   fn1 (b, c, x);
   #pragma omp target data map(to: b)
   {
-    #pragma omp target map(tofrom: c)
+    #pragma omp target map(tofrom: c, s)
       #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
 	#pragma omp distribute dist_schedule(static, 4) collapse(1)
 	  for (j=0; j < x; j += y)
@@ -52,7 +52,7 @@ fn3 (int x)
   double b[1024], c[1024], s = 0;
   int i;
   fn1 (b, c, x);
-  #pragma omp target map(to: b, c)
+  #pragma omp target map(to: b, c) map(tofrom:s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	tgt (), s += b[i] * c[i];
@@ -66,7 +66,8 @@ fn4 (int x, double *p)
   int i;
   fn1 (b, c, x);
   fn1 (d + x, p + x, x);
-  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
+  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+		     map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + p[x + i];
diff --git a/libgomp/testsuite/libgomp.c/target-11.c b/libgomp/testsuite/libgomp.c/target-11.c
new file mode 100644
index 00000000000..625c2863f4b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -0,0 +1,86 @@
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 32
+
+void test_array_section (int *p)
+{
+  #pragma omp target data map(alloc: p[0:N])
+    {
+      int ok = 1;
+      for (int i = 10; i < 10 + 4; i++)
+	p[i] = 997 * i;
+
+      #pragma omp target map(always to:p[10:4]) map(tofrom: ok)
+	for (int i = 10; i < 10 + 4; i++)
+	  if (p[i] != 997 * i)
+	    ok = 0;
+
+      assert (ok);
+
+      #pragma omp target map(always from:p[7:9])
+	for (int i = 0; i < N; i++)
+	  p[i] = i;
+    }
+}
+
+int main ()
+{
+  int aa = 0, bb = 0, cc = 0, dd = 0;
+
+  #pragma omp target data map(tofrom: aa) map(to: bb) map(from: cc, dd)
+    {
+      int ok;
+      aa = bb = cc = 1;
+
+      /* Set dd on target to 0 for the further check.  */
+      #pragma omp target map(always to: dd)
+	;
+
+      dd = 1;
+      #pragma omp target map(tofrom: aa) map(always to: bb) \
+	map(always from: cc) map(to: dd) map(from: ok)
+	{
+	  /* bb is always to, aa and dd are not.  */
+	  ok = (aa == 0) && (bb == 1) && (dd == 0);
+	  aa = bb = cc = dd = 2;
+	}
+
+      assert (ok);
+      assert (aa == 1);
+      assert (bb == 1);
+      assert (cc == 2); /* cc is always from.  */
+      assert (dd == 1);
+
+      dd = 3;
+      #pragma omp target map(from: cc) map(always to: dd) map(from: ok)
+	{
+	  ok = (dd == 3); /* dd is always to.  */
+	  cc = dd = 4;
+	}
+
+      assert (ok);
+      assert (cc == 2);
+      assert (dd == 3);
+    }
+
+  assert (aa == 2);
+  assert (bb == 1);
+  assert (cc == 4);
+  assert (dd == 4);
+
+  int *array = calloc (N, sizeof (int));
+  test_array_section (array);
+
+  for (int i = 0; i < 7; i++)
+    assert (array[i] == 0);
+  for (int i = 7; i < 7 + 9; i++)
+    assert (array[i] == i);
+  for (int i = 7 + 9; i < N; i++)
+    assert (array[i] == 0);
+
+  free (array);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-12.c b/libgomp/testsuite/libgomp.c/target-12.c
new file mode 100644
index 00000000000..e6b009463ad
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-12.c
@@ -0,0 +1,130 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int err;
+  int q[128], i;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  for (i = 0; i < 128; i++)
+    q[i] = i;
+
+  p = omp_target_alloc (130 * sizeof (int), d);
+  if (p == NULL)
+    return 0;
+
+  if (omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, NULL,
+			      d, id) < 3
+      || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
+				 NULL, id, d) < 3
+      || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
+				 NULL, id, id) < 3)
+    abort ();
+
+  if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) == 0)
+    {
+      size_t volume[3] = { 128, 0, 0 };
+      size_t dst_offsets[3] = { 0, 0, 0 };
+      size_t src_offsets[3] = { 1, 0, 0 };
+      size_t dst_dimensions[3] = { 128, 0, 0 };
+      size_t src_dimensions[3] = { 128, 0, 0 };
+
+      if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) != 0)
+	abort ();
+
+      if (omp_target_is_present (q, d) != 1
+	  || omp_target_is_present (&q[32], d) != 1
+	  || omp_target_is_present (&q[128], d) != 1)
+	abort ();
+
+      if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
+			     d, id) != 0)
+	abort ();
+
+      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
+      {
+	int j;
+	err = 0;
+	for (j = 0; j < 128; j++)
+	  if (q[j] != j)
+	    err = 1;
+	  else
+	    q[j] += 4;
+      }
+
+      if (err)
+	abort ();
+
+      if (omp_target_memcpy_rect (q, p, sizeof (int), 1, volume,
+				  dst_offsets, src_offsets, dst_dimensions,
+				  src_dimensions, id, d) != 0)
+	abort ();
+
+      for (i = 0; i < 128; i++)
+	if (q[i] != i + 4)
+	  abort ();
+
+      volume[2] = 2;
+      volume[1] = 3;
+      volume[0] = 6;
+      dst_offsets[2] = 1;
+      dst_offsets[1] = 0;
+      dst_offsets[0] = 0;
+      src_offsets[2] = 1;
+      src_offsets[1] = 0;
+      src_offsets[0] = 3;
+      dst_dimensions[2] = 2;
+      dst_dimensions[1] = 3;
+      dst_dimensions[0] = 6;
+      src_dimensions[2] = 3;
+      src_dimensions[1] = 4;
+      src_dimensions[0] = 6;
+      if (omp_target_memcpy_rect (p, q, sizeof (int), 3, volume,
+				  dst_offsets, src_offsets, dst_dimensions,
+				  src_dimensions, d, id) != 0)
+	abort ();
+
+      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
+      {
+	int j, k, l;
+	err = 0;
+	for (j = 0; j < 6; j++)
+	  for (k = 0; k < 3; k++)
+	    for (l = 0; l < 2; l++)
+	      if (q[j * 6 + k * 2 + l] != 3 * 12 + 4 + 1 + l + k * 3 + j * 12)
+		err = 1;
+      }
+
+      if (err)
+	abort ();
+
+      if (omp_target_memcpy (p, p, 10 * sizeof (int), 51 * sizeof (int),
+			     111 * sizeof (int), d, d) != 0)
+	abort ();
+
+      #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
+	{
+	  int j;
+	  err = 0;
+	  for (j = 0; j < 10; j++)
+	    if (q[50 + j] != q[110 + j])
+	      err = 1;
+	}
+
+      if (err)
+	abort ();
+
+      if (omp_target_disassociate_ptr (q, d) != 0)
+	abort ();
+    }
+
+  omp_target_free (p, d);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-13.c b/libgomp/testsuite/libgomp.c/target-13.c
new file mode 100644
index 00000000000..168850b507c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-13.c
@@ -0,0 +1,45 @@
+#ifdef __cplusplus
+extern "C"
+#else
+extern
+#endif
+void abort (void);
+struct S { int s, t; };
+
+void
+foo ()
+{
+  int x = 5, y = 6, err = 0;
+  struct S u = { 7, 8 }, v = { 9, 10 };
+  double s = 11.5, t = 12.5;
+  #pragma omp target private (x, u, s) firstprivate (y, v, t) map(from:err)
+  {
+    x = y;
+    u = v;
+    s = t;
+    err = (x != 6 || y != 6
+	   || u.s != 9 || u.t != 10 || v.s != 9 || v.t != 10
+	   || s != 12.5 || t != 12.5);
+    x += 1;
+    y += 2;
+    u.s += 3;
+    v.t += 4;
+    s += 2.5;
+    t += 3.0;
+    if (x != 7 || y != 8
+	|| u.s != 12 || u.t != 10 || v.s != 9 || v.t != 14
+	|| s != 15.0 || t != 15.5)
+      err = 1;
+  }
+  if (err || x != 5 || y != 6
+      || u.s != 7 || u.t != 8 || v.s != 9 || v.t != 10
+      || s != 11.5 || t != 12.5)
+    abort ();
+}
+
+int
+main ()
+{
+  foo ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-14.c b/libgomp/testsuite/libgomp.c/target-14.c
new file mode 100644
index 00000000000..17d383407a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-14.c
@@ -0,0 +1,38 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int err;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  p = omp_target_alloc (128 * sizeof (int), d);
+  if (p == NULL)
+    return 0;
+
+  #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0)
+  {
+    int i, *q = (int *) p;
+    for (i = 0; i < 128; i++)
+      q[i] = i + 7;
+  }
+  #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0) map(from:err)
+  {
+    int i;
+    err = 0;
+    for (i = 0; i < 128; i++)
+      if (((int *) p)[i] != i + 7)
+	err = 1;
+  }
+  if (err)
+    abort ();
+
+  omp_target_free (p, d);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-15.c b/libgomp/testsuite/libgomp.c/target-15.c
new file mode 100644
index 00000000000..fee9252ef3d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-15.c
@@ -0,0 +1,74 @@
+extern void abort (void);
+
+void
+foo (int *x)
+{
+  int a[10], b[15], err, i;
+  for (i = 0; i < 10; i++)
+    a[i] = 7 * i;
+  for (i = 0; i < 15; i++)
+    b[i] = 8 * i;
+  #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if (x[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n, int v)
+{
+  int a[n], b[n], c[n], d[n], e[n], err, i;
+  int (*x)[n] = &c;
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 4 * i;
+      a[i] = 7 * i;
+      b[i] = 8 * i;
+    }
+  #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 9 * i;
+      a[i] = 12 * i;
+      b[i] = 13 * i;
+    }
+  #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 45 + 9 * i
+	  || a[i] != 12 * i
+	  || b[5 + i] != 65 + 13 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int x[15], i;
+  for (i = 0; i < 15; i++)
+    x[i] = 4 * i;
+  foo (x);
+  bar (15, 5);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-16.c b/libgomp/testsuite/libgomp.c/target-16.c
new file mode 100644
index 00000000000..7b0919b1b00
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-16.c
@@ -0,0 +1,45 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 7 * i;
+  #pragma omp target firstprivate (a) map(from:err) private (i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n)
+{
+  int a[n], i, err;
+  #pragma omp target private (a) map(from:err)
+  {
+    #pragma omp parallel for
+    for (i = 0; i < n; i++)
+      a[i] = 7 * i;
+    err = 0;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err |= 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (7);
+  bar (7);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-17.c b/libgomp/testsuite/libgomp.c/target-17.c
new file mode 100644
index 00000000000..4a762012eaf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-17.c
@@ -0,0 +1,99 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 5 * i;
+  #pragma omp target map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 5 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 6 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  int on = n;
+  #pragma omp target firstprivate (n) map(tofrom: n)
+  {
+    n++;
+  }
+  if (on != n)
+    abort ();
+  #pragma omp target map(tofrom: n) private (n)
+  {
+    n = 25;
+  }
+  if (on != n)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 9 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(tofrom:a) map(from:err) private(a, i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      a[i] = 7;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7)
+	err |= 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    if (a[i] != 10 * i)
+      abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-18.c b/libgomp/testsuite/libgomp.c/target-18.c
new file mode 100644
index 00000000000..cbacaf6a77a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-18.c
@@ -0,0 +1,52 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[4] = { 0, 1, 2, 3 }, b[n];
+  int *p = a + 1, i, err;
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(p) map(from:err)
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < 4; i++)
+    a[i] = 23 + i;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(a) map(from:err)
+  #pragma omp target is_device_ptr(a) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (a[i] != 23 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:b)
+  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c
new file mode 100644
index 00000000000..710c5078ff6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-19.c
@@ -0,0 +1,127 @@
+extern void abort (void);
+
+void
+foo (int *p, int *q, int *r, int n, int m)
+{
+  int i, err, *s = r;
+  #pragma omp target data map(to:p[0:8])
+  {
+    /* For zero length array sections, p points to the start of
+       already mapped range, q to the end of it, and r does not point
+       to an mapped range.  */
+    #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* Implicit mapping of pointers behaves the same way.  */
+    #pragma omp target private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* And zero-length array sections, though not known at compile
+       time, behave the same.  */
+    #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* Non-zero length array sections, though not known at compile,
+       behave differently.  */
+    #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1)
+	  err = 1;
+      if (q[0] != 9 || r[0] != 10)
+	err = 1;
+    }
+    if (err)
+      abort ();
+    #pragma omp target data map(to:q[0:1])
+    {
+      /* For zero length array sections, p points to the start of
+	 already mapped range, q points to the start of another one,
+	 and r to the end of the second one.  */
+      #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* Implicit mapping of pointers behaves the same way.  */
+      #pragma omp target private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* And zero-length array sections, though not known at compile
+	 time, behave the same.  */
+      #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* Non-zero length array sections, though not known at compile,
+	 behave differently.  */
+      #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r[0] != 10)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+    }
+  }
+}
+
+int
+main ()
+{
+  int a[32], i;
+  for (i = 0; i < 32; i++)
+    a[i] = i;
+  foo (a + 1, a + 9, a + 10, 0, 1);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-2.c b/libgomp/testsuite/libgomp.c/target-2.c
index ada8dad81ad..0ba766c0a82 100644
--- a/libgomp/testsuite/libgomp.c/target-2.c
+++ b/libgomp/testsuite/libgomp.c/target-2.c
@@ -23,7 +23,7 @@ fn2 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d + x, x);
-  #pragma omp target map(to: b, c[:x], d[x:x], e)
+  #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
@@ -38,7 +38,7 @@ fn3 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d, x);
-  #pragma omp target
+  #pragma omp target map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[i];
@@ -56,7 +56,7 @@ fn4 (int x)
   #pragma omp target data map(from: b, c[:x], d[x:x], e)
     {
       #pragma omp target update to(b, c[:x], d[x:x], e)
-      #pragma omp target map(c[:x], d[x:x])
+      #pragma omp target map(c[:x], d[x:x], s)
 	#pragma omp parallel for reduction(+:s)
 	  for (i = 0; i < x; i++)
 	    {
diff --git a/libgomp/testsuite/libgomp.c/target-20.c b/libgomp/testsuite/libgomp.c/target-20.c
new file mode 100644
index 00000000000..3f4e798a755
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-20.c
@@ -0,0 +1,120 @@
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 40
+
+int sum;
+int var1 = 1;
+int var2 = 2;
+
+#pragma omp declare target
+int D[N];
+#pragma omp end declare target
+
+void enter_data (int *X)
+{
+  #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
+}
+
+void exit_data_0 (int *D)
+{
+  #pragma omp target exit data map(delete: D[:N])
+}
+
+void exit_data_1 ()
+{
+  #pragma omp target exit data map(from: var1)
+}
+
+void exit_data_2 (int *X)
+{
+  #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
+}
+
+void exit_data_3 (int *p)
+{
+  #pragma omp target exit data map(from: p[:0])
+}
+
+void test_nested ()
+{
+  int X = 0, Y = 0, Z = 0;
+
+  #pragma omp target data map(from: X, Y, Z)
+    {
+      #pragma omp target data map(from: X, Y, Z)
+	{
+	  #pragma omp target map(from: X, Y, Z)
+	    X = Y = Z = 1337;
+	  assert (X == 0);
+	  assert (Y == 0);
+	  assert (Z == 0);
+
+	  #pragma omp target exit data map(from: X) map(release: Y)
+	  assert (X == 0);
+	  assert (Y == 0);
+
+	  #pragma omp target exit data map(release: Y) map(delete: Z)
+	  assert (Y == 0);
+	  assert (Z == 0);
+	}
+      assert (X == 1337);
+      assert (Y == 0);
+      assert (Z == 0);
+
+      #pragma omp target map(from: X)
+	X = 2448;
+      assert (X == 2448);
+      assert (Y == 0);
+      assert (Z == 0);
+
+      X = 4896;
+    }
+  assert (X == 4896);
+  assert (Y == 0);
+  assert (Z == 0);
+}
+
+int main ()
+{
+  int *X = malloc (N * sizeof (int));
+  int *Y = malloc (N * sizeof (int));
+  X[10] = 10;
+  Y[20] = 20;
+  enter_data (X);
+
+  exit_data_0 (D); /* This should have no effect on D.  */
+
+  #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) map(always from: sum)
+    {
+      var1 += X[10];
+      var2 += Y[20];
+      sum = var1 + var2;
+      D[sum]++;
+    }
+
+  assert (var1 == 1);
+  assert (var2 == 2);
+  assert (sum == 33);
+
+  exit_data_1 ();
+  assert (var1 == 11);
+  assert (var2 == 2);
+
+  /* Increase refcount of already mapped X[0:N].  */
+  #pragma omp target enter data map(alloc: X[16:1])
+
+  exit_data_2 (X);
+  assert (var2 == 22);
+
+  exit_data_3 (X + 5); /* Unmap X[0:N].  */
+
+  free (X);
+  free (Y);
+
+  test_nested ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-21.c b/libgomp/testsuite/libgomp.c/target-21.c
new file mode 100644
index 00000000000..41498cf2148
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-21.c
@@ -0,0 +1,79 @@
+extern void abort (void);
+union U { int x; long long y; };
+struct T { int a; union U b; int c; };
+struct S { int s; int u; struct T v; int x[10]; union U w; int y[10]; int z[10]; };
+volatile int z;
+
+int
+main ()
+{
+  struct S s;
+  s.s = 0;
+  s.u = 1;
+  s.v.a = 2;
+  s.v.b.y = 3LL;
+  s.v.c = 19;
+  s.w.x = 4;
+  s.x[0] = 7;
+  s.x[1] = 8;
+  s.y[3] = 9;
+  s.y[4] = 10;
+  s.y[5] = 11;
+  int err = 0;
+  #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+		     map (tofrom:s.y[3:3]) \
+		     map (from: s.w, s.z[z + 1:z + 3], err)
+  {
+    err = 0;
+    if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+	|| s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
+      err = 1;
+    s.w.x = 6;
+    s.y[3] = 12;
+    s.y[4] = 13;
+    s.y[5] = 14;
+    s.z[1] = 15;
+    s.z[2] = 16;
+    s.z[3] = 17;
+  }
+  if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+      || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
+    abort ();
+  s.u++;
+  s.v.a++;
+  s.v.b.y++;
+  s.w.x++;
+  s.x[1] = 18;
+  s.z[0] = 19;
+  #pragma omp target data map (tofrom: s)
+  #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b, s.z[z:z + 1])
+  {
+    err = 0;
+    if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19)
+      err = 1;
+    s.w.x = 8;
+    s.x[1] = 20;
+    s.z[0] = 21;
+  }
+  if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
+    abort ();
+  s.u++;
+  s.v.a++;
+  s.v.b.y++;
+  s.w.x++;
+  s.x[0] = 22;
+  s.x[1] = 23;
+  #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+  #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
+  {
+    err = 0;
+    if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23)
+      err = 1;
+    s.w.x = 11;
+    s.x[0] = 24;
+    s.x[1] = 25;
+  }
+  if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-22.c b/libgomp/testsuite/libgomp.c/target-22.c
new file mode 100644
index 00000000000..aad8a0a09df
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-22.c
@@ -0,0 +1,51 @@
+extern void abort (void);
+struct T { int a; int *b; int c; };
+struct S { int *s; char *u; struct T v; short *w; };
+volatile int z;
+
+int
+main ()
+{
+  struct S s;
+  int a[32], i;
+  char b[32];
+  short c[32];
+  for (i = 0; i < 32; i++)
+    {
+      a[i] = i;
+      b[i] = 32 + i;
+      c[i] = 64 + i;
+    }
+  s.s = a;
+  s.u = b + 2;
+  s.v.b = a + 16;
+  s.w = c + 3;
+  int err = 0;
+  #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+		     map (tofrom:s.s[3:3]) \
+		     map (from: s.w[z:4], err) private (i)
+  {
+    err = 0;
+    for (i = 0; i < 7; i++)
+      if (s.v.b[i] != 16 + i)
+	err = 1;
+    for (i = 1; i < 5; i++)
+      if (s.u[i] != 34 + i)
+	err = 1;
+    for (i = 3; i < 6; i++)
+      if (s.s[i] != i)
+	err = 1;
+      else
+	s.s[i] = 128 + i;
+    for (i = 0; i < 4; i++)
+      s.w[i] = 96 + i;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < 32; i++)
+    if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+	|| b[i] != 32 + i
+	|| c[i] != ((i >= 3 && i < 7) ? 93 + i : 64 + i))
+      abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c
new file mode 100644
index 00000000000..fb1532a07b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-23.c
@@ -0,0 +1,48 @@
+extern void abort (void);
+struct S { int s; int *u; int v[5]; };
+volatile int z;
+
+int
+main ()
+{
+  int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+  struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+  int *v = u + 4;
+  #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+  s.s++;
+  u[3]++;
+  s.v[1]++;
+  #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3])
+  #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+  {
+    err = 0;
+    if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+      err = 1;
+    if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+      err = 1;
+    s.s++;
+    s.v[2] += 2;
+    v[-1] = 5;
+    v[3] = 9;
+  }
+  if (err)
+    abort ();
+  #pragma omp target map (alloc: s.u[0:5])
+  {
+    err = 0;
+    if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] != 9)
+      err = 1;
+    s.u[1] = 12;
+  }
+  #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+  if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+      || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+      || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+      || s.v[3] != 13 || s.v[4] != 14)
+    abort ();
+  #pragma omp target exit data map (release: s.s)
+  #pragma omp target exit data map (release: s.u[0:5])
+  #pragma omp target exit data map (delete: s.v[1:3])
+  #pragma omp target exit data map (release: s.s)
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-24.c b/libgomp/testsuite/libgomp.c/target-24.c
new file mode 100644
index 00000000000..e0ff29aaee8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-24.c
@@ -0,0 +1,43 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+  int *b = a;
+  int shared_mem = 0;
+  #pragma omp target map (alloc: shared_mem)
+  shared_mem = 1;
+  if (omp_target_is_present (b, d) != shared_mem)
+    abort ();
+  #pragma omp target enter data map (to: a)
+  if (omp_target_is_present (b, d) == 0)
+    abort ();
+  #pragma omp target enter data map (alloc: b[:0])
+  if (omp_target_is_present (b, d) == 0)
+    abort ();
+  #pragma omp target exit data map (release: b[:0])
+  if (omp_target_is_present (b, d) == 0)
+    abort ();
+  #pragma omp target exit data map (release: b[:0])
+  if (omp_target_is_present (b, d) != shared_mem)
+    abort ();
+  #pragma omp target enter data map (to: a)
+  if (omp_target_is_present (b, d) == 0)
+    abort ();
+  #pragma omp target enter data map (always, to: b[:0])
+  if (omp_target_is_present (b, d) == 0)
+    abort ();
+  #pragma omp target exit data map (delete: b[:0])
+  if (omp_target_is_present (b, d) != shared_mem)
+    abort ();
+  #pragma omp target exit data map (from: b[:0])
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-25.c b/libgomp/testsuite/libgomp.c/target-25.c
new file mode 100644
index 00000000000..aeb19aee510
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-25.c
@@ -0,0 +1,84 @@
+#include <stdlib.h>
+#include <unistd.h>
+
+int
+main ()
+{
+  int x = 0, y = 0, z = 0, s = 11, t = 12, u = 13, w = 7, err;
+  #pragma omp parallel
+  #pragma omp single
+  {
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      x = 1;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (6000);
+      y = 2;
+    }
+    #pragma omp task depend(out: z)
+    {
+      usleep (7000);
+      z = 3;
+    }
+    #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z)
+    err = (x != 1 || y != 2 || z != 3);
+    if (err)
+      abort ();
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      x = 4;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (4000);
+      y = 5;
+    }
+    #pragma omp task depend(in: z)
+    {
+      usleep (3000);
+      z = 6;
+    }
+    #pragma omp target enter data nowait map (to: w)
+    #pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
+    #pragma omp target map (alloc: x, y, z) map(from: err)
+    {
+      err = (x != 4 || y != 5 || z != 6);
+      x = 7;
+      y = 8;
+      z = 9;
+    }
+    if (err)
+      abort ();
+    #pragma omp taskwait
+    #pragma omp target map (alloc: w) map(from: err)
+    {
+      err = w != 7;
+      w = 17;
+    }
+    if (err)
+      abort (); 
+    #pragma omp task depend(in: x)
+    {
+      usleep (2000);
+      s = 14;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (3000);
+      t = 15;
+    }
+    #pragma omp task depend(in: z)
+    {
+      usleep (4000);
+      u = 16;
+    }
+    #pragma omp target exit data depend (inout: x, z) map (from: x, y, z, w)
+    if (x != 7 || y != 8 || z != 9 || s != 14 || t != 15 || u != 16 || w != 17)
+      abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-26.c b/libgomp/testsuite/libgomp.c/target-26.c
new file mode 100644
index 00000000000..fa6b52598da
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-26.c
@@ -0,0 +1,36 @@
+extern void abort (void);
+#pragma omp declare target
+int a[4] = { 2, 3, 4, 5 }, *b;
+#pragma omp end declare target
+
+int
+main ()
+{
+  int err;
+  int c[3] = { 6, 7, 8 };
+  b = c;
+  #pragma omp target map(to: a[0:2], b[0:2]) map(from: err)
+  err = a[0] != 2 || a[1] != 3 || a[2] != 4 || a[3] != 5 || b[0] != 6 || b[1] != 7;
+  if (err)
+    abort ();
+  a[1] = 9;
+  a[2] = 10;
+  #pragma omp target map(always,to:a[1:2]) map(from: err)
+  err = a[0] != 2 || a[1] != 9 || a[2] != 10 || a[3] != 5;
+  if (err)
+    abort ();
+  #pragma omp parallel firstprivate(a, b, c, err) num_threads (2)
+  #pragma omp single
+  {
+    b = c + 1;
+    a[0] = 11;
+    a[2] = 13;
+    c[1] = 14;
+    int d = 0;
+    #pragma omp target map(to: a[0:3], b[d:2]) map (from: err)
+    err = a[0] != 11 || a[1] != 9 || a[2] != 13 || b[0] != 14 || b[1] != 8;
+    if (err)
+      abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-27.c b/libgomp/testsuite/libgomp.c/target-27.c
new file mode 100644
index 00000000000..c86651b02e3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-27.c
@@ -0,0 +1,67 @@
+#include <stdlib.h>
+#include <unistd.h>
+
+int
+main ()
+{
+  int x = 0, y = 0, z = 0, err;
+  int shared_mem = 0;
+  #pragma omp target map(to: shared_mem)
+  shared_mem = 1;
+  #pragma omp parallel
+  #pragma omp single
+  {
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      x = 1;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (6000);
+      y = 2;
+    }
+    #pragma omp task depend(out: z)
+    {
+      usleep (7000);
+      z = 3;
+    }
+    #pragma omp target enter data map(to: x, y, z) depend(inout: x, z) nowait
+    #pragma omp task depend(inout: x, z)
+    {
+      x++; y++; z++;
+    }
+    #pragma omp target update to(x, y) depend(inout: x) nowait
+    #pragma omp target enter data map(always, to: z) depend(inout: z) nowait
+    #pragma omp target map (alloc: x, y, z) map (from: err) depend(inout: x, z)
+    {
+      err = x != 2 || y != 3 || z != 4;
+      x = 5; y = 6; z = 7;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      if (!shared_mem)
+	x = 1;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (6000);
+      if (!shared_mem)
+	y = 2;
+    }
+    #pragma omp task depend(out: z)
+    {
+      usleep (3000);
+      if (!shared_mem)
+	z = 3;
+    }
+    #pragma omp target exit data map(release: z) depend(inout: z) nowait
+    #pragma omp target exit data map(from: x, y) depend(inout: x) nowait
+    #pragma omp target exit data map(from: z) depend(inout: z) nowait
+    #pragma omp taskwait
+    if (err || x != 5 || y != 6 || z != 7)
+      abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-7.c b/libgomp/testsuite/libgomp.c/target-7.c
index 0fe6150283d..41a1332bb0c 100644
--- a/libgomp/testsuite/libgomp.c/target-7.c
+++ b/libgomp/testsuite/libgomp.c/target-7.c
@@ -37,63 +37,63 @@ foo (int f)
     abort ();
   #pragma omp target data device (d) map (to: h)
   {
-    #pragma omp target device (d)
+    #pragma omp target device (d) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
       abort ();
     #pragma omp target update device (d) from (h)
   }
   #pragma omp target data if (v > 1) map (to: h)
   {
-    #pragma omp target if (v > 1)
+    #pragma omp target if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
       abort ();
     #pragma omp target update if (v > 1) from (h)
   }
   #pragma omp target data device (d) if (v > 1) map (to: h)
   {
-    #pragma omp target device (d) if (v > 1)
+    #pragma omp target device (d) if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
       abort ();
     #pragma omp target update device (d) if (v > 1) from (h)
   }
   #pragma omp target data if (v <= 1) map (to: h)
   {
-    #pragma omp target if (v <= 1)
+    #pragma omp target if (v <= 1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 8)
       abort ();
     #pragma omp target update if (v <= 1) from (h)
   }
   #pragma omp target data device (d) if (v <= 1) map (to: h)
   {
-    #pragma omp target device (d) if (v <= 1)
+    #pragma omp target device (d) if (v <= 1) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
       abort ();
     #pragma omp target update device (d) if (v <= 1) from (h)
   }
   #pragma omp target data if (0) map (to: h)
   {
-    #pragma omp target if (0)
+    #pragma omp target if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
       abort ();
     #pragma omp target update if (0) from (h)
   }
   #pragma omp target data device (d) if (0) map (to: h)
   {
-    #pragma omp target device (d) if (0)
+    #pragma omp target device (d) if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
       abort ();
     #pragma omp target update device (d) if (0) from (h)
   }
   #pragma omp target data if (1) map (to: h)
   {
-    #pragma omp target if (1)
+    #pragma omp target if (1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 12)
       abort ();
     #pragma omp target update if (1) from (h)
   }
   #pragma omp target data device (d) if (1) map (to: h)
   {
-    #pragma omp target device (d) if (1)
+    #pragma omp target device (d) if (1) map (tofrom: h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
       abort ();
     #pragma omp target update device (d) if (1) from (h)
diff --git a/libgomp/testsuite/libgomp.c/taskloop-1.c b/libgomp/testsuite/libgomp.c/taskloop-1.c
new file mode 100644
index 00000000000..21551f2950c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/taskloop-1.c
@@ -0,0 +1,46 @@
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp -std=c99" } */
+
+int q, r, e;
+
+__attribute__((noinline, noclone)) void
+foo (long a, long b)
+{
+  #pragma omp taskloop lastprivate (q) nogroup
+    for (long d = a; d < b; d += 2)
+      {
+	q = d;
+	if (d < 2 || d > 6 || (d & 1))
+	  #pragma omp atomic
+	    e |= 1;
+      }
+}
+
+__attribute__((noinline, noclone)) int
+bar (int a, int b)
+{
+  int q = 7;
+  #pragma omp taskloop lastprivate (q)
+    for (int d = a; d < b; d++)
+      {
+	if (d < 12 || d > 17)
+	  #pragma omp atomic
+	    e |= 1;
+	q = d;
+      }
+  return q;
+}
+
+int
+main ()
+{
+  #pragma omp parallel
+    #pragma omp single
+      {
+	foo (2, 7);
+	r = bar (12, 18);
+      }
+  if (q != 6 || r != 17 || e)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/taskloop-2.c b/libgomp/testsuite/libgomp.c/taskloop-2.c
new file mode 100644
index 00000000000..be893ebf80a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/taskloop-2.c
@@ -0,0 +1,147 @@
+/* { dg-do run } */
+/* { dg-options "-O2 -std=c99" } */
+/* { dg-additional-options "-msse2" { target sse2_runtime } } */
+/* { dg-additional-options "-mavx" { target avx_runtime } } */
+
+int u[1024], v[1024], w[1024], m;
+
+__attribute__((noinline, noclone)) void
+f1 (long a, long b)
+{
+  #pragma omp taskloop simd default(none) shared(u, v, w) nogroup
+  for (long d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+}
+
+__attribute__((noinline, noclone)) int
+f2 (long a, long b, long c)
+{
+  int d, e;
+  #pragma omp taskloop simd default(none) shared(u, v, w) linear(d:1) linear(c:5) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      c = c + 5;
+      e = c + 9;
+    }
+  return d + c + e;
+}
+
+__attribute__((noinline, noclone)) int
+f3 (long a, long b)
+{
+  int d;
+  #pragma omp taskloop simd default(none) shared(u, v, w)
+  for (d = a; d < b; d++)
+    {
+      int *p = &d;
+      u[d] = v[d] + w[d];
+    }
+  return d;
+}
+
+__attribute__((noinline, noclone)) int
+f4 (long a, long b, long c, long d)
+{
+  int e, f, g;
+  #pragma omp taskloop simd default(none) shared(u, v, w) collapse(2) lastprivate(g)
+  for (e = a; e < b; e++)
+    for (f = c; f < d; f++)
+      {
+	int *p = &e;
+	int *q = &f;
+	int r = 32 * e + f;
+	u[r] = v[r] + w[r];
+	g = r;
+      }
+  return e + f + g;
+}
+
+__attribute__((noinline, noclone)) int
+f5 (long a, long b, long c, long d)
+{
+  int e, f;
+  #pragma omp taskloop simd default(none) shared(u, v, w) collapse(2)
+  for (e = a; e < b; e++)
+    for (f = c; f < d; f++)
+      {
+	int r = 32 * e + f;
+	u[r] = v[r] + w[r];
+      }
+  return e + f;
+}
+
+int
+main ()
+{
+  int i;
+  for (i = 0; i < 1024; i++)
+    {
+      v[i] = i;
+      w[i] = i + 1;
+    }
+  #pragma omp parallel
+    #pragma omp single
+      f1 (0, 1024);
+  for (i = 0; i < 1024; i++)
+    if (u[i] != 2 * i + 1)
+      __builtin_abort ();
+    else
+      {
+	v[i] = 1024 - i;
+	w[i] = 512 - i;
+      }
+  #pragma omp parallel
+    #pragma omp single
+      m = f2 (2, 1022, 17);
+  for (i = 0; i < 1024; i++)
+    if ((i < 2 || i >= 1022) ? u[i] != 2 * i + 1 : u[i] != 1536 - 2 * i)
+      __builtin_abort ();
+    else
+      {
+	v[i] = i;
+	w[i] = i + 1;
+      }
+  if (m != 1022 + 2 * (1020 * 5 + 17) + 9)
+    __builtin_abort ();
+  #pragma omp parallel
+    #pragma omp single
+      m = f3 (0, 1024);
+  for (i = 0; i < 1024; i++)
+    if (u[i] != 2 * i + 1)
+      __builtin_abort ();
+    else
+      {
+	v[i] = 1024 - i;
+	w[i] = 512 - i;
+      }
+  if (m != 1024)
+    __builtin_abort ();
+  #pragma omp parallel
+    #pragma omp single
+      m = f4 (0, 32, 0, 32);
+  for (i = 0; i < 1024; i++)
+    if (u[i] != 1536 - 2 * i)
+      __builtin_abort ();
+    else
+      {
+	v[i] = i;
+	w[i] = i + 1;
+      }
+  if (m != 32 + 32 + 1023)
+    __builtin_abort ();
+  #pragma omp parallel
+    #pragma omp single
+      m = f5 (0, 32, 0, 32);
+  for (i = 0; i < 1024; i++)
+    if (u[i] != 2 * i + 1)
+      __builtin_abort ();
+    else
+      {
+	v[i] = 1024 - i;
+	w[i] = 512 - i;
+      }
+  if (m != 32 + 32)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/taskloop-3.c b/libgomp/testsuite/libgomp.c/taskloop-3.c
new file mode 100644
index 00000000000..5356d7f0251
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/taskloop-3.c
@@ -0,0 +1,84 @@
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp -std=c99" } */
+
+int g;
+int a[1024];
+
+__attribute__((noinline, noclone)) int
+f1 (int x)
+{
+  #pragma omp taskloop firstprivate (x) lastprivate (x)
+  for (int i = 0; i < 64; i++)
+    {
+      if (x != 74)
+	__builtin_abort ();
+      if (i == 63)
+	x = i + 4;
+    }
+  return x;
+}
+
+__attribute__((noinline, noclone)) void
+f2 (void)
+{
+  #pragma omp taskloop firstprivate (g) lastprivate (g) nogroup
+  for (int i = 0; i < 64; i++)
+    {
+      if (g != 77)
+	__builtin_abort ();
+      if (i == 63)
+	g = i + 9;
+    }
+}
+
+__attribute__((noinline, noclone)) long long
+f3 (long long a, long long b, long long c)
+{
+  long long i;
+  int l;
+  #pragma omp taskloop default (none) lastprivate (i, l)
+  for (i = a; i < b; i += c)
+    l = i;
+  return l * 7 + i;
+}
+
+__attribute__((noinline, noclone)) long long
+f4 (long long a, long long b, long long c, long long d,
+    long long e, long long f, int k)
+{
+  long long i, j;
+  int l;
+  #pragma omp taskloop default (none) collapse(2) \
+	      firstprivate (k) lastprivate (i, j, k, l)
+  for (i = a; i < b; i += e)
+    for (j = c; j < d; j += f)
+      {
+	if (k != 73)
+	  __builtin_abort ();
+	if (i == 31 && j == 46)
+	  k = i;
+	l = j;
+      }
+  return i + 5 * j + 11 * k + 17 * l;
+}
+
+int
+main ()
+{
+  #pragma omp parallel
+    #pragma omp single
+      {
+	if (f1 (74) != 63 + 4)
+	  __builtin_abort ();
+	g = 77;
+	f2 ();
+	#pragma omp taskwait
+	if (g != 63 + 9)
+	  __builtin_abort ();
+	if (f3 (7, 12, 2) != 11 * 7 + 13)
+	  __builtin_abort ();
+	if (f4 (0, 32, 16, 48, 1, 2, 73) != 32 + 5 * 48 + 11 * 31 + 17 * 46)
+	  __builtin_abort ();
+      }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/taskloop-4.c b/libgomp/testsuite/libgomp.c/taskloop-4.c
new file mode 100644
index 00000000000..a69be19c9c2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/taskloop-4.c
@@ -0,0 +1,97 @@
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp" } */
+
+int u[64], v;
+
+__attribute__((noinline, noclone)) int
+test (int a, int b, int c, int d, void (*fn) (int, int, int, int),
+      int *num_tasks, int *min_iters, int *max_iters)
+{
+  int i, t = 0;
+  __builtin_memset (u, 0, sizeof u);
+  v = 0;
+  fn (a, b, c, d);
+  *min_iters = 0;
+  *max_iters = 0;
+  *num_tasks = v;
+  if (v)
+    {
+      *min_iters = u[0];
+      *max_iters = u[0];
+      t = u[0];
+      for (i = 1; i < v; i++)
+	{
+	  if (*min_iters > u[i])
+	    *min_iters = u[i];
+	  if (*max_iters < u[i])
+	    *max_iters = u[i];
+	  t += u[i];
+	}
+    }
+  return t;
+}
+
+void
+grainsize (int a, int b, int c, int d)
+{
+  int i, j = 0, k = 0;
+  #pragma omp taskloop firstprivate (j, k) grainsize(d)
+  for (i = a; i < b; i += c)
+    {
+      if (j == 0)
+	{
+	  #pragma omp atomic capture
+	    k = v++;
+	  if (k >= 64)
+	    __builtin_abort ();
+	}
+      u[k] = ++j;
+    }
+}
+
+void
+num_tasks (int a, int b, int c, int d)
+{
+  int i, j = 0, k = 0;
+  #pragma omp taskloop firstprivate (j, k) num_tasks(d)
+  for (i = a; i < b; i += c)
+    {
+      if (j == 0)
+	{
+	  #pragma omp atomic capture
+	    k = v++;
+	  if (k >= 64)
+	    __builtin_abort ();
+	}
+      u[k] = ++j;
+    }
+}
+
+int
+main ()
+{
+  #pragma omp parallel
+    #pragma omp single
+      {
+	int min_iters, max_iters, ntasks;
+	/* If grainsize is present, # of task loop iters is >= grainsize && < 2 * grainsize,
+	   unless # of loop iterations is smaller than grainsize.  */
+	if (test (0, 79, 1, 17, grainsize, &ntasks, &min_iters, &max_iters) != 79
+	    || min_iters < 17 || max_iters >= 17 * 2)
+	  __builtin_abort ();
+	if (test (-49, 2541, 7, 28, grainsize, &ntasks, &min_iters, &max_iters) != 370
+	    || min_iters < 28 || max_iters >= 28 * 2)
+	  __builtin_abort ();
+	if (test (7, 21, 2, 15, grainsize, &ntasks, &min_iters, &max_iters) != 7
+	    || ntasks != 1 || min_iters != 7 || max_iters != 7)
+	  __builtin_abort ();
+	/* If num_tasks is present, # of task loop iters is min (# of loop iters, num_tasks).  */
+	if (test (-51, 2500, 48, 9, num_tasks, &ntasks, &min_iters, &max_iters) != 54
+	    || ntasks != 9)
+	  __builtin_abort ();
+	if (test (0, 25, 2, 17, num_tasks, &ntasks, &min_iters, &max_iters) != 13
+	    || ntasks != 13)
+	  __builtin_abort ();
+      }
+  return 0;
+}