gcc/

2015-11-05  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* builtin-types.def
	(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove.
	(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): New.
	* cgraph.h (enum cgraph_simd_clone_arg_type): Add
	SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP,
	SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP and
	SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP.
	(struct cgraph_simd_clone_arg): Adjust comment.
	* omp-builtins.def (BUILT_IN_GOMP_TARGET): Rename GOMP_target_41
	to GOMP_target_ext.  Add num_teams and thread_limit arguments.
	(BUILT_IN_GOMP_TARGET_DATA): Rename GOMP_target_data_41
	to GOMP_target_data_ext.
	(BUILT_IN_GOMP_TARGET_UPDATE): Rename GOMP_target_update_41
	to GOMP_target_update_ext.
	(BUILT_IN_GOMP_LOOP_NONMONOTONIC_DYNAMIC_START,
	BUILT_IN_GOMP_LOOP_NONMONOTONIC_GUIDED_START,
	BUILT_IN_GOMP_LOOP_NONMONOTONIC_DYNAMIC_NEXT,
	BUILT_IN_GOMP_LOOP_NONMONOTONIC_GUIDED_NEXT,
	BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_DYNAMIC_START,
	BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_GUIDED_START,
	BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_DYNAMIC_NEXT,
	BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_GUIDED_NEXT,
	BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_DYNAMIC,
	BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_GUIDED): New built-ins.
	* tree-core.h (enum omp_clause_schedule_kind): Add
	OMP_CLAUSE_SCHEDULE_MASK, OMP_CLAUSE_SCHEDULE_MONOTONIC,
	OMP_CLAUSE_SCHEDULE_NONMONOTONIC and change
	OMP_CLAUSE_SCHEDULE_LAST value.
	* tree.def (OMP_SIMD, CILK_SIMD, CILK_FOR, OMP_DISTRIBUTE,
	OMP_TASKLOOP, OACC_LOOP): Add OMP_FOR_ORIG_DECLS argument.
	* tree.h (OMP_FOR_ORIG_DECLS): Use OMP_LOOP_CHECK instead of
	OMP_FOR_CHECK.  Remove comment.
	* tree-pretty-print.c (dump_omp_clause): Handle
	GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER.
	Simplify.  Print schedule clause modifiers.
	* tree-vect-stmts.c (vectorizable_simd_clone_call): Add
	SIMD_CLONE_ARG_TYPE_LINEAR_{REF,VAL,UVAL}_VARIABLE_STEP
	cases.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_ALWAYS_TO.
	(omp_default_clause): Tweak for
	private/firstprivate/is_device_ptr variables on target
	construct and use_device_ptr on target data.
	(omp_check_private): Likewise.
	(omp_notice_variable): For references check whether what it refers
	to has mappable type, rather than the reference itself.
	(omp_is_private): Diagnose linear iteration variables on non-simd
	constructs.
	(omp_no_lastprivate): Return true only for Fortran.
	(gimplify_scan_omp_clauses): Or in GOVD_MAP_ALWAYS_TO for
	GOMP_MAP_ALWAYS_TO or GOMP_MAP_ALWAYS_TOFROM kinds.
	Add support for GOMP_MAP_FIRSTPRIVATE_REFERENCE and
	GOMP_MAP_ALWAYS_POINTER, remove old handling of structure element
	based array sections.  Use GOMP_MAP_ALWAYS_P.  Fix up handling of
	lastprivate and linear when combined with distribute.  Gimplify
	variable low-bound for array reduction.  Look through
	POINTER_PLUS_EXPR when looking for ADDR_EXPR for array section
	reductions.
	(gimplify_adjust_omp_clauses_1): For implicit references to
	variables with reference type and when not ref to scalar or
	ref to pointer, map what they refer to using tofrom and
	use GOMP_MAP_FIRSTPRIVATE_REFERENCE for the reference.
	(gimplify_adjust_omp_clauses): Remove GOMP_MAP_ALWAYS_POINTER
	from target exit data.  Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE.
	Drop OMP_CLAUSE_MAP_PRIVATE support.  Use GOMP_MAP_ALWAYS_P.
	Diagnose the same var on both firstprivate and lastprivate on
	distribute construct.
	(gimplify_omp_for): Fix up handling of predetermined
	lastprivate or linear iter vars when combined with distribute.
	(find_omp_teams, computable_teams_clause, optimize_target_teams): New
	functions.
	(gimplify_omp_workshare): Call optimize_target_teams.
	* omp-low.c (struct omp_region): Add sched_modifiers field.
	(struct omp_for_data): Likewise.
	(omp_any_child_fn_dumped): New variable.
	(extract_omp_for_data): Fill in sched_modifiers, and mask out
	OMP_CLAUSE_SCHEDULE_KIND bits outside of OMP_CLAUSE_SCHEDULE_MASK
	from sched_kind.
	(determine_parallel_type): Use only OMP_CLAUSE_SCHEDULE_MASK
	bits of OMP_CLAUSE_SCHED_KIND.
	(scan_sharing_clauses): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE,
	drop OMP_CLAUSE_MAP_PRIVATE support.  Look through POINTER_PLUS_EXPR
	for array section reductions.
	(add_taskreg_looptemp_clauses): Add one extra _looptemp_ clause even
	for distribute parallel for, if there are lastprivate clauses on the
	for.
	(lower_rec_input_clauses): Handle non-zero low-bound on array
	section reductions.
	(lower_reduction_clauses): Likewise.
	(lower_send_clauses): Look through POINTER_PLUS_EXPR
	for array section reductions.
	(expand_parallel_call): Use nonmonotonic entrypoints for
	nonmonotonic: dynamic/guided.
	(expand_omp_taskreg): Call assign_assembler_name_if_neeeded on
	child_fn if current_function_decl has assembler name set, but child_fn
	does not.  Dump the header and IL of the child function when not in SSA
	form.
	(expand_omp_target): Likewise.  Pass num_teams and thread_limit
	arguments to BUILT_IN_GOMP_TARGET.
	(expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
	Initialize the extra _looptemp_ clause to fd->loop.n2.
	(expand_omp_for): Use nonmonotonic entrypoints for
	nonmonotonic: dynamic/guided.  Initialize region->sched_modifiers.
	(expand_omp): Clear omp_any_child_fn_dumped.  Dump function header
	again if we have dumped any child functions.
	(lower_omp_for_lastprivate): Determine the right count variable
	for distribute simd, or distribute parallel for{, simd}.
	(lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE
	and GOMP_MAP_ALWAYS_POINTER.  Drop OMP_CLAUSE_MAP_PRIVATE
	support.
	(simd_clone_clauses_extract): Handle variable step
	for references and arguments passed by reference.
	(simd_clone_mangle): Mangle ref/uval/val variable steps.
	(simd_clone_adjust_argument_types): Handle
	SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP like
	SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP and
	SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP like
	SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP.
	(simd_clone_linear_addend): New function.
	(simd_clone_adjust): Handle variable step like similarly
	to constant step, use simd_clone_linear_addend to determine
	the actual step at runtime.
gcc/c-family/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>

	* c-common.h (c_finish_omp_atomic): Add TEST argument.
	(c_omp_check_loop_iv, c_omp_check_loop_iv_exprs): New prototypes.
	* c-omp.c (c_finish_omp_atomic): Add TEST argument.  Don't call
	save_expr or create_tmp_var* if TEST is true.
	(c_finish_omp_for): Store OMP_FOR_ORIG_DECLS always.
	Don't call add_stmt here.
	(struct c_omp_check_loop_iv_data): New type.
	(c_omp_check_loop_iv_r, c_omp_check_loop_iv,
	c_omp_check_loop_iv_exprs): New functions.
	(c_omp_split_clauses): Adjust for lastprivate being allowed on
	distribute.
	(c_omp_declare_simd_clauses_to_numbers): Change
	OMP_CLAUSE_LINEAR_VARIABLE_STRIDE OMP_CLAUSE_LINEAR_STEP into numbers.
	(c_omp_declare_simd_clauses_to_decls): Similarly change those
	from numbers to PARM_DECLs.
gcc/c/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* c-parser.c: Include context.h and gimple-expr.h.
	(c_parser_omp_clause_schedule): Parse schedule modifiers, diagnose
	monotonic together with nonmonotonic.
	(c_parser_omp_for_loop): Call c_omp_check_loop_iv.  Call add_stmt here.
	(OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause.
	(c_parser_omp_target_data, c_parser_omp_target_enter_data,
	c_parser_omp_target_exit_data): Allow GOMP_MAP_ALWAYS_POINTER.
	(c_parser_omp_target): Likewise.  Evaluate num_teams and thread_limit
	expressions on combined target teams before the target.
	(c_parser_omp_declare_target): If decl has "omp declare target" or
	"omp declare target link" attribute, and cgraph or varpool node already
	exists, then set corresponding flags.  Call c_finish_omp_clauses
	in the parenthesized extended-list syntax case.
	* c-decl.c (c_decl_attributes): Don't diagnose block scope vars inside
	declare target.
	* c-typeck.c (handle_omp_array_sections_1): Allow non-zero low-bound
	on OMP_CLAUSE_REDUCTION array sections.
	(handle_omp_array_sections): Encode low-bound into the MEM_REF, either
	into the constant offset, or for variable low-bound using
	POINTER_PLUS_EXPR.  For structure element based array sections use
	GOMP_MAP_ALWAYS_POINTER instead of GOMP_MAP_FIRSTPRIVATE_POINTER.
	(c_finish_omp_clauses): Drop generic_field_head, structure
	elements are now always mapped even as array section bases,
	diagnose same var in data sharing and mapping clauses.  Diagnose if
	linear step on declare simd is neither a constant nor a uniform
	parameter.  Look through POINTER_PLUS_EXPR for array section
	reductions.  Diagnose the same var or function appearing multiple
	times on the same directive.  Fix up wording for the to clause if t
	is neither a FUNCTION_DECL nor a VAR_DECL.  Diagnose nonmonotonic
	modifier on kinds other than dynamic or guided or nonmonotonic
	modifier together with ordered clause.
gcc/cp/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* cp-tree.h (finish_omp_for): Add ORIG_INITS argument.
	(omp_privatize_field): Add SHARED argument.
	* parser.c: Include context.h.
	(cp_parser_omp_clause_schedule): Parse schedule
	modifiers, diagnose monotonic together with nonmonotonic.
	(cp_parser_omp_clause_linear): Add DECLARE_SIMD argument.  Parse
	parameter name as linear step as id-expression rather than expression.
	(cp_parser_omp_all_clauses): Adjust caller.
	(cp_parser_omp_for_loop_init): Add ORIG_INIT argument,
	initialize it.  Adjust omp_privatize_field caller.
	(cp_parser_omp_for_loop): Compute orig_inits, pass it's address
	to finish_omp_for.
	(OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause.
	(cp_parser_omp_target_data,
	cp_parser_omp_target_enter_data,
	cp_parser_omp_target_exit_data): Allow GOMP_MAP_ALWAYS_POINTER
	and GOMP_MAP_FIRSTPRIVATE_REFERENCE.
	(cp_parser_omp_target): Likewise.  Evaluate num_teams and
	thread_limit expressions on combined target teams before the target.
	(cp_parser_omp_declare_target): If decl has "omp declare target" or
	"omp declare target link" attribute, and cgraph or varpool node already
	exists, then set corresponding flags.  Call finish_omp_clauses
	in the parenthesized extended-list syntax case.  Call
	cp_parser_require_pragma_eol instead of cp_parser_skip_to_pragma_eol.
	(cp_parser_omp_end_declare_target): Call cp_parser_require_pragma_eol
	instead of cp_parser_skip_to_pragma_eol.
	* decl2.c (cplus_decl_attributes): Don't diagnose block scope vars inside
	declare target.
	* pt.c (tsubst_omp_clauses): If OMP_CLAUSE_LINEAR_VARIABLE_STRIDE,
	use tsubst_omp_clause_decl instead of tsubst_expr on
	OMP_CLAUSE_LINEAR_STEP.  Handle non-static data members in shared
	clauses.
	(tsubst_omp_for_iterator): Adjust omp_privatize_field caller.
	(tsubst_find_omp_teams): New function.
	(tsubst_expr): Evaluate num_teams and thread_limit expressions on
	combined target teams before the target.  Use OMP_FOR_ORIG_DECLS for
	all OpenMP/OpenACC/Cilk+ looping constructs.  Adjust finish_omp_for
	caller.
	* semantics.c (omp_privatize_field): Add SHARED argument, if true,
	always create artificial var and never put it into the hash table
	or vector.
	(handle_omp_array_sections_1): Adjust omp_privatize_field caller.
	Allow non-zero low-bound on OMP_CLAUSE_REDUCTION array sections.
	(handle_omp_array_sections): For structure element
	based array sections use GOMP_MAP_ALWAYS_POINTER instead of
	GOMP_MAP_FIRSTPRIVATE_POINTER.  Encode low-bound into the MEM_REF,
	either into the constant offset, or for variable low-bound using
	POINTER_PLUS_EXPR.
	(finish_omp_clauses): Adjust omp_privatize_field caller.  Drop
	generic_field_head, structure elements are now always mapped even
	as array section bases, diagnose same var in data sharing and
	mapping clauses.  For references map what they refer to using
	GOMP_MAP_ALWAYS_POINTER for structure elements and
	GOMP_MAP_FIRSTPRIVATE_REFERENCE otherwise.  Diagnose if linear step
	on declare simd is neither a constant nor a uniform parameter.
	Allow non-static data members on shared clauses.  Look through
	POINTER_PLUS_EXPR for array section reductions.  Diagnose nonmonotonic
	modifier on kinds other than dynamic or guided or nonmonotonic
	modifier together with ordered clause.  Diagnose the same var or
	function appearing multiple times on the same directive.  Fix up
	wording for the to clause if t is neither a FUNCTION_DECL nor a
	VAR_DECL, use special wording for OVERLOADs and TEMPLATE_ID_EXPR.
	(handle_omp_for_class_iterator): Add ORIG_DECLS argument.  Call
	c_omp_check_loop_iv_exprs on cond.
	(finish_omp_for): Add ORIG_INITS argument.  Call
	c_omp_check_loop_iv_exprs on ORIG_INITS elements.  Adjust
	handle_omp_for_class_iterator caller.  Call c_omp_check_loop_iv.
	Call add_stmt.
	(finish_omp_atomic): Adjust c_finish_omp_atomic caller.
gcc/fortran/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>

	* types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove.
	(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): New.
gcc/testsuite/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>

	* c-c++-common/gomp/clauses-2.c (foo): Adjust for diagnostics
	of variables in both data sharing and mapping clauses and for
	structure element based array sections being mapped rather than
	privatized.
	* c-c++-common/gomp/declare-target-2.c: Add various new tests.  Adjust
	expected diagnostics wording in one case.
	* c-c++-common/gomp/distribute-1.c: New test.
	* c-c++-common/gomp/element-1.c: New test.
	* c-c++-common/gomp/pr61486-2.c: Add #pragma omp declare target
	and #pragma omp end declare target pair around the function.
	Change s from a parameter to a file scope variable.
	* c-c++-common/gomp/pr67521.c: Add dg-error directives.
	* c-c++-common/gomp/reduction-1.c (foo): Don't expect diagnostics
	on non-zero low-bound in reduction array sections.  Add further
	tests.
	* c-c++-common/gomp/schedule-modifiers-1.c: New test.
	* c-c++-common/gomp/target-teams-1.c: New test.
	* gcc.dg/gomp/declare-simd-1.c: Add scan-assembler-times directives
	for expected mangling on x86_64/i?86.
	* gcc.dg/gomp/declare-simd-3.c: New test.
	* gcc.dg/gomp/declare-simd-4.c: New test.
	* gcc.dg/gomp/for-20.c: New test.
	* gcc.dg/gomp/for-21.c: New test.
	* gcc.dg/gomp/for-22.c: New test.
	* gcc.dg/gomp/for-23.c: New test.
	* gcc.dg/gomp/for-24.c: New test.
	* gcc.dg/gomp/linear-1.c: New test.
	* gcc.dg/gomp/loop-1.c: New test.
	* g++.dg/gomp/atomic-17.C: New test.
	* g++.dg/gomp/clause-1.C (T::test): Don't expect error on
	non-static data member in shared clause.  Add single construct.
	* g++.dg/gomp/declare-simd-1.C: Add dg-options.  Add
	scan-assembler-times directives for expected mangling on x86_64/i?86.
	* g++.dg/gomp/declare-simd-3.C: Likewise.
	* g++.dg/gomp/declare-simd-4.C: New test.
	* g++.dg/gomp/declare-simd-5.C: New test.
	* g++.dg/gomp/declare-target-1.C: New test.
	* g++.dg/gomp/linear-2.C: New test.
	* g++.dg/gomp/loop-1.C: New test.
	* g++.dg/gomp/loop-2.C: New test.
	* g++.dg/gomp/loop-3.C: New test.
	* g++.dg/gomp/member-2.C (B::m2, B::m4): Don't expect error on
	non-static data member in shared clause.
	* g++.dg/gomp/member-3.C: New test.
	* g++.dg/gomp/member-4.C: New test.
	* g++.dg/gomp/pr38639.C (foo): Adjust dg-error.
	(bar): Remove dg-message.
	* g++.dg/gomp/target-teams-1.C: New test.
include/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_2): Define.
	(GOMP_MAP_FLAG_ALWAYS): Remove.
	(enum gomp_map_kind): Use GOMP_MAP_FLAG_SPECIAL_2 instead of
	GOMP_MAP_FLAG_ALWAYS for GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM,
	GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_STRUCT, GOMP_MAP_RELEASE.
	Add GOMP_MAP_ALWAYS_POINTER and GOMP_MAP_FIRSTPRIVATE_REFERENCE.
	(GOMP_MAP_ALWAYS_P): Define.
	(GOMP_TARGET_FLAG_NOWAIT): Adjust comment.
libgomp/
2015-11-05  Jakub Jelinek  <jakub@redhat.com>
	    Ilya Verbin  <ilya.verbin@intel.com>

	* libgomp_g.h (GOMP_loop_nonmonotonic_dynamic_next,
	GOMP_loop_nonmonotonic_dynamic_start,
	GOMP_loop_nonmonotonic_guided_next,
	GOMP_loop_nonmonotonic_guided_start,
	GOMP_loop_ull_nonmonotonic_dynamic_next,
	GOMP_loop_ull_nonmonotonic_dynamic_start,
	GOMP_loop_ull_nonmonotonic_guided_next,
	GOMP_loop_ull_nonmonotonic_guided_start,
	GOMP_parallel_loop_nonmonotonic_dynamic,
	GOMP_parallel_loop_nonmonotonic_guided): New prototypes.
	(GOMP_target_41): Renamed to ...
	(GOMP_target_ext): ... this.  Add num_teams and thread_limit
	arguments.
	(GOMP_target_data_41): Renamed to ...
	(GOMP_target_data_ext): ... this.
	(GOMP_target_update_41): Renamed to ...
	(GOMP_target_update_ext): ... this.
	* libgomp.map (GOMP_4.5): Export GOMP_target_ext,
	GOMP_target_data_ext and GOMP_target_update_ext instead of
	GOMP_target_41, GOMP_target_data_41 and GOMP_target_update_41.
	Export GOMP_loop_nonmonotonic_dynamic_next,
	GOMP_loop_nonmonotonic_dynamic_start,
	GOMP_loop_nonmonotonic_guided_next,
	GOMP_loop_nonmonotonic_guided_start,
	GOMP_loop_ull_nonmonotonic_dynamic_next,
	GOMP_loop_ull_nonmonotonic_dynamic_start,
	GOMP_loop_ull_nonmonotonic_guided_next,
	GOMP_loop_ull_nonmonotonic_guided_start,
	GOMP_parallel_loop_nonmonotonic_dynamic and
	GOMP_parallel_loop_nonmonotonic_guided.
	* loop.c (GOMP_parallel_loop_nonmonotonic_dynamic,
	GOMP_parallel_loop_nonmonotonic_guided,
	GOMP_loop_nonmonotonic_dynamic_start,
	GOMP_loop_nonmonotonic_guided_start,
	GOMP_loop_nonmonotonic_dynamic_next,
	GOMP_loop_nonmonotonic_guided_next): New aliases or functions.
	* loop_ull.c (GOMP_loop_ull_nonmonotonic_dynamic_start,
	GOMP_loop_ull_nonmonotonic_guided_start,
	GOMP_loop_ull_nonmonotonic_dynamic_next,
	GOMP_loop_ull_nonmonotonic_guided_next): Likewise.
	* target.c (gomp_map_0len_lookup, gomp_map_val): New inline
	functions.
	(gomp_map_vars): Handle GOMP_MAP_ALWAYS_POINTER.  For
	GOMP_MAP_ZERO_LEN_ARRAY_SECTION use gomp_map_0len_lookup.
	Use gomp_map_val function.
	(gomp_target_fallback_firstprivate): New static function.
	(GOMP_target_41): Renamed to ...
	(GOMP_target_ext): ... this.  Add num_teams and thread_limit
	arguments.  Move firstprivate fallback handling into a new
	function.
	(GOMP_target_data_41): Renamed to ...
	(GOMP_target_data_ext): ... this.
	(GOMP_target_update_41): Renamed to ...
	(GOMP_target_update_ext): ... this.
	(gomp_exit_data): For GOMP_MAP_*ZERO_LEN* use
	gomp_map_0len_lookup instead of gomp_map_lookup.
	(omp_target_is_present): Use gomp_map_0len_lookup instead of
	gomp_map_lookup.
	* testsuite/libgomp.c/target-28.c: Likewise.
	* testsuite/libgomp.c/monotonic-1.c: New test.
	* testsuite/libgomp.c/monotonic-2.c: New test.
	* testsuite/libgomp.c/nonmonotonic-1.c: New test.
	* testsuite/libgomp.c/nonmonotonic-2.c: New test.
	* testsuite/libgomp.c/pr66199-5.c: New test.
	* testsuite/libgomp.c/pr66199-6.c: New test.
	* testsuite/libgomp.c/pr66199-7.c: New test.
	* testsuite/libgomp.c/pr66199-8.c: New test.
	* testsuite/libgomp.c/pr66199-9.c: New test.
	* testsuite/libgomp.c/reduction-11.c: New test.
	* testsuite/libgomp.c/reduction-12.c: New test.
	* testsuite/libgomp.c/reduction-13.c: New test.
	* testsuite/libgomp.c/reduction-14.c: New test.
	* testsuite/libgomp.c/reduction-15.c: New test.
	* testsuite/libgomp.c/target-12.c (main): Adjust for
	omp_target_is_present change for one-past-last element.
	* testsuite/libgomp.c/target-17.c (foo): Drop tests where
	the same var is both mapped and privatized.
	* testsuite/libgomp.c/target-19.c (foo): Adjust for different
	handling of zero-length array sections.
	* testsuite/libgomp.c/target-28.c: New test.
	* testsuite/libgomp.c/target-29.c: New test.
	* testsuite/libgomp.c/target-30.c: New test.
	* testsuite/libgomp.c/target-teams-1.c: New test.
	* testsuite/libgomp.c++/member-6.C: New test.
	* testsuite/libgomp.c++/member-7.C: New test.
	* testsuite/libgomp.c++/monotonic-1.C: New test.
	* testsuite/libgomp.c++/monotonic-2.C: New test.
	* testsuite/libgomp.c++/nonmonotonic-1.C: New test.
	* testsuite/libgomp.c++/nonmonotonic-2.C: New test.
	* testsuite/libgomp.c++/pr66199-3.C: New test.
	* testsuite/libgomp.c++/pr66199-4.C: New test.
	* testsuite/libgomp.c++/pr66199-5.C: New test.
	* testsuite/libgomp.c++/pr66199-6.C: New test.
	* testsuite/libgomp.c++/pr66199-7.C: New test.
	* testsuite/libgomp.c++/pr66199-8.C: New test.
	* testsuite/libgomp.c++/pr66199-9.C: New test.
	* testsuite/libgomp.c++/reduction-11.C: New test.
	* testsuite/libgomp.c++/reduction-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.


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

49 files changed, 3304 insertions, 139 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 46680f183dc..26377b64060 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,112 @@
+2015-11-05  Jakub Jelinek  <jakub@redhat.com>
+	    Ilya Verbin  <ilya.verbin@intel.com>
+
+	* libgomp_g.h (GOMP_loop_nonmonotonic_dynamic_next,
+	GOMP_loop_nonmonotonic_dynamic_start,
+	GOMP_loop_nonmonotonic_guided_next,
+	GOMP_loop_nonmonotonic_guided_start,
+	GOMP_loop_ull_nonmonotonic_dynamic_next,
+	GOMP_loop_ull_nonmonotonic_dynamic_start,
+	GOMP_loop_ull_nonmonotonic_guided_next,
+	GOMP_loop_ull_nonmonotonic_guided_start,
+	GOMP_parallel_loop_nonmonotonic_dynamic,
+	GOMP_parallel_loop_nonmonotonic_guided): New prototypes.
+	(GOMP_target_41): Renamed to ...
+	(GOMP_target_ext): ... this.  Add num_teams and thread_limit
+	arguments.
+	(GOMP_target_data_41): Renamed to ...
+	(GOMP_target_data_ext): ... this.
+	(GOMP_target_update_41): Renamed to ...
+	(GOMP_target_update_ext): ... this.
+	* libgomp.map (GOMP_4.5): Export GOMP_target_ext,
+	GOMP_target_data_ext and GOMP_target_update_ext instead of
+	GOMP_target_41, GOMP_target_data_41 and GOMP_target_update_41.
+	Export GOMP_loop_nonmonotonic_dynamic_next,
+	GOMP_loop_nonmonotonic_dynamic_start,
+	GOMP_loop_nonmonotonic_guided_next,
+	GOMP_loop_nonmonotonic_guided_start,
+	GOMP_loop_ull_nonmonotonic_dynamic_next,
+	GOMP_loop_ull_nonmonotonic_dynamic_start,
+	GOMP_loop_ull_nonmonotonic_guided_next,
+	GOMP_loop_ull_nonmonotonic_guided_start,
+	GOMP_parallel_loop_nonmonotonic_dynamic and
+	GOMP_parallel_loop_nonmonotonic_guided.
+	* loop.c (GOMP_parallel_loop_nonmonotonic_dynamic,
+	GOMP_parallel_loop_nonmonotonic_guided,
+	GOMP_loop_nonmonotonic_dynamic_start,
+	GOMP_loop_nonmonotonic_guided_start,
+	GOMP_loop_nonmonotonic_dynamic_next,
+	GOMP_loop_nonmonotonic_guided_next): New aliases or functions.
+	* loop_ull.c (GOMP_loop_ull_nonmonotonic_dynamic_start,
+	GOMP_loop_ull_nonmonotonic_guided_start,
+	GOMP_loop_ull_nonmonotonic_dynamic_next,
+	GOMP_loop_ull_nonmonotonic_guided_next): Likewise.
+	* target.c (gomp_map_0len_lookup, gomp_map_val): New inline
+	functions.
+	(gomp_map_vars): Handle GOMP_MAP_ALWAYS_POINTER.  For
+	GOMP_MAP_ZERO_LEN_ARRAY_SECTION use gomp_map_0len_lookup.
+	Use gomp_map_val function.
+	(gomp_target_fallback_firstprivate): New static function.
+	(GOMP_target_41): Renamed to ...
+	(GOMP_target_ext): ... this.  Add num_teams and thread_limit
+	arguments.  Move firstprivate fallback handling into a new
+	function.
+	(GOMP_target_data_41): Renamed to ...
+	(GOMP_target_data_ext): ... this.
+	(GOMP_target_update_41): Renamed to ...
+	(GOMP_target_update_ext): ... this.
+	(gomp_exit_data): For GOMP_MAP_*ZERO_LEN* use
+	gomp_map_0len_lookup instead of gomp_map_lookup.
+	(omp_target_is_present): Use gomp_map_0len_lookup instead of
+	gomp_map_lookup.
+	* testsuite/libgomp.c/target-28.c: Likewise.
+	* testsuite/libgomp.c/monotonic-1.c: New test.
+	* testsuite/libgomp.c/monotonic-2.c: New test.
+	* testsuite/libgomp.c/nonmonotonic-1.c: New test.
+	* testsuite/libgomp.c/nonmonotonic-2.c: New test.
+	* testsuite/libgomp.c/pr66199-5.c: New test.
+	* testsuite/libgomp.c/pr66199-6.c: New test.
+	* testsuite/libgomp.c/pr66199-7.c: New test.
+	* testsuite/libgomp.c/pr66199-8.c: New test.
+	* testsuite/libgomp.c/pr66199-9.c: New test.
+	* testsuite/libgomp.c/reduction-11.c: New test.
+	* testsuite/libgomp.c/reduction-12.c: New test.
+	* testsuite/libgomp.c/reduction-13.c: New test.
+	* testsuite/libgomp.c/reduction-14.c: New test.
+	* testsuite/libgomp.c/reduction-15.c: New test.
+	* testsuite/libgomp.c/target-12.c (main): Adjust for
+	omp_target_is_present change for one-past-last element.
+	* testsuite/libgomp.c/target-17.c (foo): Drop tests where
+	the same var is both mapped and privatized.
+	* testsuite/libgomp.c/target-19.c (foo): Adjust for different
+	handling of zero-length array sections.
+	* testsuite/libgomp.c/target-28.c: New test.
+	* testsuite/libgomp.c/target-29.c: New test.
+	* testsuite/libgomp.c/target-30.c: New test.
+	* testsuite/libgomp.c/target-teams-1.c: New test.
+	* testsuite/libgomp.c++/member-6.C: New test.
+	* testsuite/libgomp.c++/member-7.C: New test.
+	* testsuite/libgomp.c++/monotonic-1.C: New test.
+	* testsuite/libgomp.c++/monotonic-2.C: New test.
+	* testsuite/libgomp.c++/nonmonotonic-1.C: New test.
+	* testsuite/libgomp.c++/nonmonotonic-2.C: New test.
+	* testsuite/libgomp.c++/pr66199-3.C: New test.
+	* testsuite/libgomp.c++/pr66199-4.C: New test.
+	* testsuite/libgomp.c++/pr66199-5.C: New test.
+	* testsuite/libgomp.c++/pr66199-6.C: New test.
+	* testsuite/libgomp.c++/pr66199-7.C: New test.
+	* testsuite/libgomp.c++/pr66199-8.C: New test.
+	* testsuite/libgomp.c++/pr66199-9.C: New test.
+	* testsuite/libgomp.c++/reduction-11.C: New test.
+	* testsuite/libgomp.c++/reduction-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.
+
 2015-11-04  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* testsuite/libgomp.oacc-fortran/reduction-1.f90: Fix dimensions
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2153661ed5a..39faba95cbd 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -266,9 +266,9 @@ GOMP_4.0.1 {
 
 GOMP_4.5 {
   global:
-	GOMP_target_41;
-	GOMP_target_data_41;
-	GOMP_target_update_41;
+	GOMP_target_ext;
+	GOMP_target_data_ext;
+	GOMP_target_update_ext;
 	GOMP_target_enter_exit_data;
 	GOMP_taskloop;
 	GOMP_taskloop_ull;
@@ -286,6 +286,16 @@ GOMP_4.5 {
 	GOMP_loop_ull_doacross_static_start;
 	GOMP_doacross_ull_post;
 	GOMP_doacross_ull_wait;
+	GOMP_loop_nonmonotonic_dynamic_next;
+	GOMP_loop_nonmonotonic_dynamic_start;
+	GOMP_loop_nonmonotonic_guided_next;
+	GOMP_loop_nonmonotonic_guided_start;
+	GOMP_loop_ull_nonmonotonic_dynamic_next;
+	GOMP_loop_ull_nonmonotonic_dynamic_start;
+	GOMP_loop_ull_nonmonotonic_guided_next;
+	GOMP_loop_ull_nonmonotonic_guided_start;
+	GOMP_parallel_loop_nonmonotonic_dynamic;
+	GOMP_parallel_loop_nonmonotonic_guided;
 } GOMP_4.0.1;
 
 OACC_2.0 {
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index c28ad2116dc..c238e6a7556 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -52,6 +52,10 @@ extern bool GOMP_loop_static_start (long, long, long, long, long *, long *);
 extern bool GOMP_loop_dynamic_start (long, long, long, long, long *, long *);
 extern bool GOMP_loop_guided_start (long, long, long, long, long *, long *);
 extern bool GOMP_loop_runtime_start (long, long, long, long *, long *);
+extern bool GOMP_loop_nonmonotonic_dynamic_start (long, long, long, long,
+						  long *, long *);
+extern bool GOMP_loop_nonmonotonic_guided_start (long, long, long, long,
+						 long *, long *);
 
 extern bool GOMP_loop_ordered_static_start (long, long, long, long,
 					    long *, long *);
@@ -65,6 +69,8 @@ extern bool GOMP_loop_static_next (long *, long *);
 extern bool GOMP_loop_dynamic_next (long *, long *);
 extern bool GOMP_loop_guided_next (long *, long *);
 extern bool GOMP_loop_runtime_next (long *, long *);
+extern bool GOMP_loop_nonmonotonic_dynamic_next (long *, long *);
+extern bool GOMP_loop_nonmonotonic_guided_next (long *, long *);
 
 extern bool GOMP_loop_ordered_static_next (long *, long *);
 extern bool GOMP_loop_ordered_dynamic_next (long *, long *);
@@ -100,6 +106,12 @@ extern void GOMP_parallel_loop_guided (void (*)(void *), void *,
 extern void GOMP_parallel_loop_runtime (void (*)(void *), void *,
 					unsigned, long, long, long,
 					unsigned);
+extern void GOMP_parallel_loop_nonmonotonic_dynamic (void (*)(void *), void *,
+						     unsigned, long, long,
+						     long, long, unsigned);
+extern void GOMP_parallel_loop_nonmonotonic_guided (void (*)(void *), void *,
+						    unsigned, long, long,
+						    long, long, unsigned);
 
 extern void GOMP_loop_end (void);
 extern void GOMP_loop_end_nowait (void);
@@ -130,6 +142,18 @@ extern bool GOMP_loop_ull_runtime_start (bool, unsigned long long,
 					 unsigned long long,
 					 unsigned long long *,
 					 unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_dynamic_start (bool, unsigned long long,
+						      unsigned long long,
+						      unsigned long long,
+						      unsigned long long,
+						      unsigned long long *,
+						      unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_guided_start (bool, unsigned long long,
+						     unsigned long long,
+						     unsigned long long,
+						     unsigned long long,
+						     unsigned long long *,
+						     unsigned long long *);
 
 extern bool GOMP_loop_ull_ordered_static_start (bool, unsigned long long,
 						unsigned long long,
@@ -163,6 +187,10 @@ extern bool GOMP_loop_ull_guided_next (unsigned long long *,
 				       unsigned long long *);
 extern bool GOMP_loop_ull_runtime_next (unsigned long long *,
 					unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_dynamic_next (unsigned long long *,
+						     unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_guided_next (unsigned long long *,
+						    unsigned long long *);
 
 extern bool GOMP_loop_ull_ordered_static_next (unsigned long long *,
 					       unsigned long long *);
@@ -249,17 +277,18 @@ extern void GOMP_single_copy_end (void *);
 
 extern void GOMP_target (int, void (*) (void *), const void *,
 			 size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_41 (int, void (*) (void *), size_t, void **, size_t *,
-			  unsigned short *, unsigned int, void **);
+extern void GOMP_target_ext (int, void (*) (void *), size_t, void **, size_t *,
+			     unsigned short *, unsigned int, void **,
+			     int, int);
 extern void GOMP_target_data (int, const void *,
 			      size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_data_41 (int, size_t, void **, size_t *,
-			       unsigned short *);
+extern void GOMP_target_data_ext (int, size_t, void **, size_t *,
+				  unsigned short *);
 extern void GOMP_target_end_data (void);
 extern void GOMP_target_update (int, const void *,
 				size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_update_41 (int, size_t, void **, size_t *,
-				   unsigned short *, unsigned int, void **);
+extern void GOMP_target_update_ext (int, size_t, void **, size_t *,
+				    unsigned short *, unsigned int, void **);
 extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *,
 					 unsigned short *, unsigned int,
 					 void **);
diff --git a/libgomp/loop.c b/libgomp/loop.c
index 812f66cd725..a7b539885fd 100644
--- a/libgomp/loop.c
+++ b/libgomp/loop.c
@@ -111,6 +111,11 @@ gomp_loop_static_start (long start, long end, long incr, long chunk_size,
   return !gomp_iter_static_next (istart, iend);
 }
 
+/* The current dynamic implementation is always monotonic.  The
+   entrypoints without nonmonotonic in them have to be always monotonic,
+   but the nonmonotonic ones could be changed to use work-stealing for
+   improved scalability.  */
+
 static bool
 gomp_loop_dynamic_start (long start, long end, long incr, long chunk_size,
 			 long *istart, long *iend)
@@ -136,6 +141,9 @@ gomp_loop_dynamic_start (long start, long end, long incr, long chunk_size,
   return ret;
 }
 
+/* Similarly as for dynamic, though the question is how can the chunk sizes
+   be decreased without a central locking or atomics.  */
+
 static bool
 gomp_loop_guided_start (long start, long end, long incr, long chunk_size,
 			long *istart, long *iend)
@@ -630,6 +638,37 @@ GOMP_parallel_loop_guided (void (*fn) (void *), void *data,
   GOMP_parallel_end ();
 }
 
+#ifdef HAVE_ATTRIBUTE_ALIAS
+extern __typeof(GOMP_parallel_loop_dynamic) GOMP_parallel_loop_nonmonotonic_dynamic
+	__attribute__((alias ("GOMP_parallel_loop_dynamic")));
+extern __typeof(GOMP_parallel_loop_guided) GOMP_parallel_loop_nonmonotonic_guided
+	__attribute__((alias ("GOMP_parallel_loop_guided")));
+#else
+void
+GOMP_parallel_loop_nonmonotonic_dynamic (void (*fn) (void *), void *data,
+					 unsigned num_threads, long start,
+					 long end, long incr, long chunk_size,
+					 unsigned flags)
+{
+  gomp_parallel_loop_start (fn, data, num_threads, start, end, incr,
+			    GFS_DYNAMIC, chunk_size, flags);
+  fn (data);
+  GOMP_parallel_end ();
+}
+
+void
+GOMP_parallel_loop_nonmonotonic_guided (void (*fn) (void *), void *data,
+					unsigned num_threads, long start,
+					long end, long incr, long chunk_size,
+					unsigned flags)
+{
+  gomp_parallel_loop_start (fn, data, num_threads, start, end, incr,
+			    GFS_GUIDED, chunk_size, flags);
+  fn (data);
+  GOMP_parallel_end ();
+}
+#endif
+
 void
 GOMP_parallel_loop_runtime (void (*fn) (void *), void *data,
 			    unsigned num_threads, long start, long end,
@@ -678,6 +717,10 @@ extern __typeof(gomp_loop_dynamic_start) GOMP_loop_dynamic_start
 	__attribute__((alias ("gomp_loop_dynamic_start")));
 extern __typeof(gomp_loop_guided_start) GOMP_loop_guided_start
 	__attribute__((alias ("gomp_loop_guided_start")));
+extern __typeof(gomp_loop_dynamic_start) GOMP_loop_nonmonotonic_dynamic_start
+	__attribute__((alias ("gomp_loop_dynamic_start")));
+extern __typeof(gomp_loop_guided_start) GOMP_loop_nonmonotonic_guided_start
+	__attribute__((alias ("gomp_loop_guided_start")));
 
 extern __typeof(gomp_loop_ordered_static_start) GOMP_loop_ordered_static_start
 	__attribute__((alias ("gomp_loop_ordered_static_start")));
@@ -699,6 +742,10 @@ extern __typeof(gomp_loop_dynamic_next) GOMP_loop_dynamic_next
 	__attribute__((alias ("gomp_loop_dynamic_next")));
 extern __typeof(gomp_loop_guided_next) GOMP_loop_guided_next
 	__attribute__((alias ("gomp_loop_guided_next")));
+extern __typeof(gomp_loop_dynamic_next) GOMP_loop_nonmonotonic_dynamic_next
+	__attribute__((alias ("gomp_loop_dynamic_next")));
+extern __typeof(gomp_loop_guided_next) GOMP_loop_nonmonotonic_guided_next
+	__attribute__((alias ("gomp_loop_guided_next")));
 
 extern __typeof(gomp_loop_ordered_static_next) GOMP_loop_ordered_static_next
 	__attribute__((alias ("gomp_loop_ordered_static_next")));
@@ -729,6 +776,21 @@ GOMP_loop_guided_start (long start, long end, long incr, long chunk_size,
 }
 
 bool
+GOMP_loop_nonmonotonic_dynamic_start (long start, long end, long incr,
+				      long chunk_size, long *istart,
+				      long *iend)
+{
+  return gomp_loop_dynamic_start (start, end, incr, chunk_size, istart, iend);
+}
+
+bool
+GOMP_loop_nonmonotonic_guided_start (long start, long end, long incr,
+				     long chunk_size, long *istart, long *iend)
+{
+  return gomp_loop_guided_start (start, end, incr, chunk_size, istart, iend);
+}
+
+bool
 GOMP_loop_ordered_static_start (long start, long end, long incr,
 				long chunk_size, long *istart, long *iend)
 {
@@ -795,6 +857,18 @@ GOMP_loop_guided_next (long *istart, long *iend)
 }
 
 bool
+GOMP_loop_nonmonotonic_dynamic_next (long *istart, long *iend)
+{
+  return gomp_loop_dynamic_next (istart, iend);
+}
+
+bool
+GOMP_loop_nonmonotonic_guided_next (long *istart, long *iend)
+{
+  return gomp_loop_guided_next (istart, iend);
+}
+
+bool
 GOMP_loop_ordered_static_next (long *istart, long *iend)
 {
   return gomp_loop_ordered_static_next (istart, iend);
diff --git a/libgomp/loop_ull.c b/libgomp/loop_ull.c
index 1f2ed546024..b00fcd52bc7 100644
--- a/libgomp/loop_ull.c
+++ b/libgomp/loop_ull.c
@@ -566,6 +566,10 @@ extern __typeof(gomp_loop_ull_dynamic_start) GOMP_loop_ull_dynamic_start
 	__attribute__((alias ("gomp_loop_ull_dynamic_start")));
 extern __typeof(gomp_loop_ull_guided_start) GOMP_loop_ull_guided_start
 	__attribute__((alias ("gomp_loop_ull_guided_start")));
+extern __typeof(gomp_loop_ull_dynamic_start) GOMP_loop_ull_nonmonotonic_dynamic_start
+	__attribute__((alias ("gomp_loop_ull_dynamic_start")));
+extern __typeof(gomp_loop_ull_guided_start) GOMP_loop_ull_nonmonotonic_guided_start
+	__attribute__((alias ("gomp_loop_ull_guided_start")));
 
 extern __typeof(gomp_loop_ull_ordered_static_start) GOMP_loop_ull_ordered_static_start
 	__attribute__((alias ("gomp_loop_ull_ordered_static_start")));
@@ -587,6 +591,10 @@ extern __typeof(gomp_loop_ull_dynamic_next) GOMP_loop_ull_dynamic_next
 	__attribute__((alias ("gomp_loop_ull_dynamic_next")));
 extern __typeof(gomp_loop_ull_guided_next) GOMP_loop_ull_guided_next
 	__attribute__((alias ("gomp_loop_ull_guided_next")));
+extern __typeof(gomp_loop_ull_dynamic_next) GOMP_loop_ull_nonmonotonic_dynamic_next
+	__attribute__((alias ("gomp_loop_ull_dynamic_next")));
+extern __typeof(gomp_loop_ull_guided_next) GOMP_loop_ull_nonmonotonic_guided_next
+	__attribute__((alias ("gomp_loop_ull_guided_next")));
 
 extern __typeof(gomp_loop_ull_ordered_static_next) GOMP_loop_ull_ordered_static_next
 	__attribute__((alias ("gomp_loop_ull_ordered_static_next")));
@@ -623,6 +631,25 @@ GOMP_loop_ull_guided_start (bool up, gomp_ull start, gomp_ull end,
 }
 
 bool
+GOMP_loop_ull_nonmonotonic_dynamic_start (bool up, gomp_ull start,
+					  gomp_ull end, gomp_ull incr,
+					  gomp_ull chunk_size,
+					  gomp_ull *istart, gomp_ull *iend)
+{
+  return gomp_loop_ull_dynamic_start (up, start, end, incr, chunk_size, istart,
+				      iend);
+}
+
+bool
+GOMP_loop_ull_nonmonotonic_guided_start (bool up, gomp_ull start, gomp_ull end,
+					 gomp_ull incr, gomp_ull chunk_size,
+					 gomp_ull *istart, gomp_ull *iend)
+{
+  return gomp_loop_ull_guided_start (up, start, end, incr, chunk_size, istart,
+				     iend);
+}
+
+bool
 GOMP_loop_ull_ordered_static_start (bool up, gomp_ull start, gomp_ull end,
 				    gomp_ull incr, gomp_ull chunk_size,
 				    gomp_ull *istart, gomp_ull *iend)
@@ -695,6 +722,18 @@ GOMP_loop_ull_guided_next (gomp_ull *istart, gomp_ull *iend)
 }
 
 bool
+GOMP_loop_ull_nonmonotonic_dynamic_next (gomp_ull *istart, gomp_ull *iend)
+{
+  return gomp_loop_ull_dynamic_next (istart, iend);
+}
+
+bool
+GOMP_loop_ull_nonmonotonic_guided_next (gomp_ull *istart, gomp_ull *iend)
+{
+  return gomp_loop_ull_guided_next (istart, iend);
+}
+
+bool
 GOMP_loop_ull_ordered_static_next (gomp_ull *istart, gomp_ull *iend)
 {
   return gomp_loop_ull_ordered_static_next (istart, iend);
diff --git a/libgomp/target.c b/libgomp/target.c
index b7674100326..1bddc6ffcee 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -162,7 +162,20 @@ gomp_map_lookup (splay_tree mem_map, splay_tree_key key)
   return splay_tree_lookup (mem_map, key);
 }
 
-/* Handle the case where gomp_map_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_0len_lookup (splay_tree mem_map, splay_tree_key key)
+{
+  if (key->host_start != key->host_end)
+    return splay_tree_lookup (mem_map, key);
+
+  key->host_end++;
+  splay_tree_key n = splay_tree_lookup (mem_map, key);
+  key->host_end--;
+  return n;
+}
+
+/* Handle the case where gomp_map_lookup, splay_tree_lookup or
+   gomp_map_0len_lookup found oldn for newn.
    Helper function of gomp_map_vars.  */
 
 static inline void
@@ -306,6 +319,26 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n,
 	      (void *) cur_node.host_end);
 }
 
+static inline uintptr_t
+gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
+{
+  if (tgt->list[i].key != NULL)
+    return tgt->list[i].key->tgt->tgt_start
+	   + tgt->list[i].key->tgt_offset
+	   + tgt->list[i].offset;
+  if (tgt->list[i].offset == ~(uintptr_t) 0)
+    return (uintptr_t) hostaddrs[i];
+  if (tgt->list[i].offset == ~(uintptr_t) 1)
+    return 0;
+  if (tgt->list[i].offset == ~(uintptr_t) 2)
+    return tgt->list[i + 1].key->tgt->tgt_start
+	   + tgt->list[i + 1].key->tgt_offset
+	   + tgt->list[i + 1].offset
+	   + (uintptr_t) hostaddrs[i]
+	   - (uintptr_t) hostaddrs[i + 1];
+  return tgt->tgt_start + tgt->list[i].offset;
+}
+
 attribute_hidden struct target_mem_desc *
 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -400,6 +433,13 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	  i--;
 	  continue;
 	}
+      else if ((kind & typemask) == GOMP_MAP_ALWAYS_POINTER)
+	{
+	  tgt->list[i].key = NULL;
+	  tgt->list[i].offset = ~(uintptr_t) 1;
+	  has_firstprivate = true;
+	  continue;
+	}
       cur_node.host_start = (uintptr_t) hostaddrs[i];
       if (!GOMP_MAP_POINTER_P (kind & typemask))
 	cur_node.host_end = cur_node.host_start + sizes[i];
@@ -420,7 +460,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
       splay_tree_key n;
       if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
 	{
-	  n = gomp_map_lookup (mem_map, &cur_node);
+	  n = gomp_map_0len_lookup (mem_map, &cur_node);
 	  if (!n)
 	    {
 	      tgt->list[i].key = NULL;
@@ -558,6 +598,32 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 					    sizes, kinds);
 		i--;
 		continue;
+	      case GOMP_MAP_ALWAYS_POINTER:
+		cur_node.host_start = (uintptr_t) hostaddrs[i];
+		cur_node.host_end = cur_node.host_start + sizeof (void *);
+		n = splay_tree_lookup (mem_map, &cur_node);
+		if (n == NULL
+		    || n->host_start > cur_node.host_start
+		    || n->host_end < cur_node.host_end)
+		  {
+		    gomp_mutex_unlock (&devicep->lock);
+		    gomp_fatal ("always pointer not mapped");
+		  }
+		if ((get_kind (short_mapkind, kinds, i - 1) & typemask)
+		    != GOMP_MAP_ALWAYS_POINTER)
+		  cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i - 1);
+		if (cur_node.tgt_offset)
+		  cur_node.tgt_offset -= sizes[i];
+		devicep->host2dev_func (devicep->target_id,
+					(void *) (n->tgt->tgt_start
+						  + n->tgt_offset
+						  + cur_node.host_start
+						  - n->host_start),
+					(void *) &cur_node.tgt_offset,
+					sizeof (void *));
+		cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+				      + cur_node.host_start - n->host_start;
+		continue;
 	      default:
 		break;
 	      }
@@ -701,26 +767,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
     {
       for (i = 0; i < mapnum; i++)
 	{
-	  if (tgt->list[i].key == NULL)
-	    {
-	      if (tgt->list[i].offset == ~(uintptr_t) 0)
-		cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
-	      else if (tgt->list[i].offset == ~(uintptr_t) 1)
-		cur_node.tgt_offset = 0;
-	      else if (tgt->list[i].offset == ~(uintptr_t) 2)
-		cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start
-				      + tgt->list[i + 1].key->tgt_offset
-				      + tgt->list[i + 1].offset
-				      + (uintptr_t) hostaddrs[i]
-				      - (uintptr_t) hostaddrs[i + 1];
-	      else
-		cur_node.tgt_offset = tgt->tgt_start
-				      + tgt->list[i].offset;
-	    }
-	  else
-	    cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
-				  + tgt->list[i].key->tgt_offset
-				  + tgt->list[i].offset;
+	  cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
 	  /* FIXME: see above FIXME comment.  */
 	  devicep->host2dev_func (devicep->target_id,
 				  (void *) (tgt->tgt_start
@@ -1213,7 +1260,7 @@ gomp_fini_device (struct gomp_device_descr *devicep)
   devicep->is_initialized = false;
 }
 
-/* Host fallback for GOMP_target{,_41} routines.  */
+/* Host fallback for GOMP_target{,_ext} routines.  */
 
 static void
 gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
@@ -1231,7 +1278,45 @@ gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
   *thr = old_thr;
 }
 
-/* Helper function of GOMP_target{,_41} routines.  */
+/* Host fallback with firstprivate map-type handling.  */
+
+static void
+gomp_target_fallback_firstprivate (void (*fn) (void *), size_t mapnum,
+				   void **hostaddrs, size_t *sizes,
+				   unsigned short *kinds)
+{
+  size_t i, tgt_align = 0, tgt_size = 0;
+  char *tgt = NULL;
+  for (i = 0; i < mapnum; i++)
+    if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+      {
+	size_t align = (size_t) 1 << (kinds[i] >> 8);
+	if (tgt_align < align)
+	  tgt_align = align;
+	tgt_size = (tgt_size + align - 1) & ~(align - 1);
+	tgt_size += sizes[i];
+      }
+  if (tgt_align)
+    {
+      tgt = gomp_alloca (tgt_size + tgt_align - 1);
+      uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
+      if (al)
+	tgt += tgt_align - al;
+      tgt_size = 0;
+      for (i = 0; i < mapnum; i++)
+	if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+	  {
+	    size_t align = (size_t) 1 << (kinds[i] >> 8);
+	    tgt_size = (tgt_size + align - 1) & ~(align - 1);
+	    memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
+	    hostaddrs[i] = tgt + tgt_size;
+	    tgt_size = tgt_size + sizes[i];
+	  }
+    }
+  gomp_target_fallback (fn, hostaddrs);
+}
+
+/* Helper function of GOMP_target{,_ext} routines.  */
 
 static void *
 gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
@@ -1294,13 +1379,31 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
   gomp_unmap_vars (tgt_vars, true);
 }
 
+/* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
+   and several arguments have been added:
+   FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
+   DEPEND is array of dependencies, see GOMP_task for details.
+   NUM_TEAMS is positive if GOMP_teams will be called in the body with
+   that value, or 1 if teams construct is not present, or 0, if
+   teams construct does not have num_teams clause and so the choice is
+   implementation defined, and -1 if it can't be determined on the host
+   what value will GOMP_teams have on the device.
+   THREAD_LIMIT similarly is positive if GOMP_teams will be called in the
+   body with that value, or 0, if teams construct does not have thread_limit
+   clause or the teams construct is not present, or -1 if it can't be
+   determined on the host what value will GOMP_teams have on the device.  */
+
 void
-GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
-		void **hostaddrs, size_t *sizes, unsigned short *kinds,
-		unsigned int flags, void **depend)
+GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
+		 void **hostaddrs, size_t *sizes, unsigned short *kinds,
+		 unsigned int flags, void **depend, int num_teams,
+		 int thread_limit)
 {
   struct gomp_device_descr *devicep = resolve_device (device);
 
+  (void) num_teams;
+  (void) thread_limit;
+
   /* If there are depend clauses, but nowait is not present,
      block the parent task until the dependencies are resolved
      and then just continue with the rest of the function as if it
@@ -1315,35 +1418,7 @@ GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
     {
-      size_t i, tgt_align = 0, tgt_size = 0;
-      char *tgt = NULL;
-      for (i = 0; i < mapnum; i++)
-	if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
-	  {
-	    size_t align = (size_t) 1 << (kinds[i] >> 8);
-	    if (tgt_align < align)
-	      tgt_align = align;
-	    tgt_size = (tgt_size + align - 1) & ~(align - 1);
-	    tgt_size += sizes[i];
-	  }
-      if (tgt_align)
-	{
-	  tgt = gomp_alloca (tgt_size + tgt_align - 1);
-	  uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
-	  if (al)
-	    tgt += tgt_align - al;
-	  tgt_size = 0;
-	  for (i = 0; i < mapnum; i++)
-	    if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
-	      {
-		size_t align = (size_t) 1 << (kinds[i] >> 8);
-		tgt_size = (tgt_size + align - 1) & ~(align - 1);
-		memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
-		hostaddrs[i] = tgt + tgt_size;
-		tgt_size = tgt_size + sizes[i];
-	      }
-	}
-      gomp_target_fallback (fn, hostaddrs);
+      gomp_target_fallback_firstprivate (fn, mapnum, hostaddrs, sizes, kinds);
       return;
     }
 
@@ -1366,7 +1441,7 @@ GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
   gomp_unmap_vars (tgt_vars, true);
 }
 
-/* Host fallback for GOMP_target_data{,_41} routines.  */
+/* Host fallback for GOMP_target_data{,_ext} routines.  */
 
 static void
 gomp_target_data_fallback (void)
@@ -1405,8 +1480,8 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
 }
 
 void
-GOMP_target_data_41 (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
-		     unsigned short *kinds)
+GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
+		      size_t *sizes, unsigned short *kinds)
 {
   struct gomp_device_descr *devicep = resolve_device (device);
 
@@ -1448,9 +1523,9 @@ GOMP_target_update (int device, const void *unused, size_t mapnum,
 }
 
 void
-GOMP_target_update_41 (int device, size_t mapnum, void **hostaddrs,
-		       size_t *sizes, unsigned short *kinds,
-		       unsigned int flags, void **depend)
+GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
+			size_t *sizes, unsigned short *kinds,
+			unsigned int flags, void **depend)
 {
   struct gomp_device_descr *devicep = resolve_device (device);
 
@@ -1527,7 +1602,7 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
 	  cur_node.host_end = cur_node.host_start + sizes[i];
 	  splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
 			      || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
-	    ? gomp_map_lookup (&devicep->mem_map, &cur_node)
+	    ? gomp_map_0len_lookup (&devicep->mem_map, &cur_node)
 	    : splay_tree_lookup (&devicep->mem_map, &cur_node);
 	  if (!k)
 	    continue;
@@ -1642,7 +1717,7 @@ gomp_target_task_fn (void *data)
   struct gomp_target_task *ttask = (struct gomp_target_task *) data;
   if (ttask->fn != NULL)
     {
-      /* GOMP_target_41 */
+      /* GOMP_target_ext */
     }
   else if (ttask->devicep == NULL
 	   || !(ttask->devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
@@ -1759,7 +1834,7 @@ omp_target_is_present (void *ptr, int device_num)
 
   cur_node.host_start = (uintptr_t) ptr;
   cur_node.host_end = cur_node.host_start;
-  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
+  splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
   int ret = n != NULL;
   gomp_mutex_unlock (&devicep->lock);
   return ret;
diff --git a/libgomp/testsuite/libgomp.c++/member-6.C b/libgomp/testsuite/libgomp.c++/member-6.C
new file mode 100644
index 00000000000..5c0c22bce37
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/member-6.C
@@ -0,0 +1,114 @@
+// { dg-do run }
+
+#include <omp.h>
+
+struct R { R () {}; ~R () {}; int r; };
+struct T { T () {}; virtual ~T () {}; int t; };
+int c;
+struct A : public R, virtual public T { A () : b(c) {} int a; int &b; void m1 (); };
+
+void
+take (int &a, int &b, int &c, int &d)
+{
+  asm volatile ("" : : "g" (&a), "g" (&b), "g" (&c), "g" (&d) : "memory");
+}
+
+void
+A::m1 ()
+{
+  #pragma omp parallel private (a, T::t) shared (r, A::b) default(none)
+  {
+    int q = omp_get_thread_num (), q2;
+    a = q;
+    t = 3 * q;
+    #pragma omp single copyprivate (q2)
+    {
+      r = 2 * q;
+      b = 4 * q;
+      q2 = q;
+    }
+    take (a, r, t, b);
+    #pragma omp barrier
+    if (A::a != q || R::r != 2 * q2 || T::t != 3 * q || A::b != 4 * q2)
+      __builtin_abort ();
+  }
+  a = 7;
+  r = 8;
+  t = 9;
+  b = 10;
+  #pragma omp parallel shared (A::a) default (none) firstprivate (R::r, b) shared (t)
+  {
+    int q = omp_get_thread_num (), q2;
+    take (A::a, R::r, T::t, A::b);
+    if (a != 7 || r != 8 || t != 9 || b != 10)
+      __builtin_abort ();
+    R::r = 6 * q;
+    #pragma omp barrier
+    #pragma omp single copyprivate (q2)
+    {
+      A::a = 5 * q;
+      T::t = 7 * q;
+      q2 = q;
+    }
+    A::b = 8 * q;
+    take (a, r, t, b);
+    #pragma omp barrier
+    if (a != 5 * q2 || r != 6 * q || t != 7 * q2 || b != 8 * q)
+      __builtin_abort ();
+  }
+  a = 1;
+  b = 2;
+  R::r = 3;
+  t = 4;
+  bool f = false;
+  #pragma omp parallel private (f)
+    {
+      f = false;
+    #pragma omp single
+    #pragma omp taskloop default(none) firstprivate (r, A::a, f) shared (T::t, b)
+      for (int i = 0; i < 30; i++)
+	{
+	  int q = omp_get_thread_num ();
+	  int tv, bv;
+	  #pragma omp atomic read
+	  tv = t;
+	  #pragma omp atomic read
+	  bv = A::b;
+	  if (i == 16)
+	    {
+	      if (bv != 2 || tv != 4)
+		__builtin_abort ();
+	    }
+	  else
+	    {
+	      if ((bv != 2 && bv != 8) || (tv != 4 && tv != 9))
+		__builtin_abort ();
+	    }
+	  if (!f)
+	    {
+	      if (A::a != 1 || R::r != 3)
+		__builtin_abort ();
+	    }
+	  else if (a != 7 * q || r != 9 * q)
+	    __builtin_abort ();
+	  take (a, r, t, b);
+	  A::a = 7 * q;
+	  R::r = 9 * q;
+	  if (i == 16)
+	    {
+	      #pragma omp atomic write
+	      A::b = 8;
+	      #pragma omp atomic write
+	      T::t = 9;
+	    }
+	  f = true;
+	}
+    }
+}
+
+int
+main ()
+{
+  A a;
+  a.m1 ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/member-7.C b/libgomp/testsuite/libgomp.c++/member-7.C
new file mode 100644
index 00000000000..15848c5952e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/member-7.C
@@ -0,0 +1,119 @@
+// { dg-do run }
+
+#include <omp.h>
+
+int c, d, e;
+struct R { R () {}; ~R () {}; int r; };
+template <typename Q>
+struct T { T () : t(d) {}; virtual ~T () {}; Q t; };
+template <typename Q>
+struct A : public R, virtual public T<Q> { A () : b(c), a(e) {} Q a; int &b; void m1 (); };
+
+void
+take (int &a, int &b, int &c, int &d)
+{
+  asm volatile ("" : : "g" (&a), "g" (&b), "g" (&c), "g" (&d) : "memory");
+}
+
+template <typename Q>
+void
+A<Q>::m1 ()
+{
+  #pragma omp parallel private (a, T<Q>::t) shared (r, A::b) default(none)
+  {
+    int q = omp_get_thread_num (), q2;
+    a = q;
+    T<Q>::t = 3 * q;
+    #pragma omp single copyprivate (q2)
+    {
+      r = 2 * q;
+      b = 4 * q;
+      q2 = q;
+    }
+    take (a, r, T<Q>::t, b);
+    #pragma omp barrier
+    if (A::a != q || R::r != 2 * q2 || T<Q>::t != 3 * q || A::b != 4 * q2)
+      __builtin_abort ();
+  }
+  a = 7;
+  r = 8;
+  T<Q>::t = 9;
+  b = 10;
+  #pragma omp parallel shared (A::a) default (none) firstprivate (R::r, b) shared (T<Q>::t)
+  {
+    int q = omp_get_thread_num (), q2;
+    take (A::a, R::r, T<Q>::t, A::b);
+    if (a != 7 || r != 8 || T<Q>::t != 9 || b != 10)
+      __builtin_abort ();
+    R::r = 6 * q;
+    #pragma omp barrier
+    #pragma omp single copyprivate (q2)
+    {
+      A::a = 5 * q;
+      T<Q>::t = 7 * q;
+      q2 = q;
+    }
+    A::b = 8 * q;
+    take (a, r, T<Q>::t, b);
+    #pragma omp barrier
+    if (a != 5 * q2 || r != 6 * q || T<Q>::t != 7 * q2 || b != 8 * q)
+      __builtin_abort ();
+  }
+  a = 1;
+  b = 2;
+  R::r = 3;
+  T<Q>::t = 4;
+  bool f = false;
+  #pragma omp parallel private (f)
+    {
+      f = false;
+    #pragma omp single
+    #pragma omp taskloop default(none) firstprivate (r, A::a, f) shared (T<Q>::t, b)
+      for (int i = 0; i < 30; i++)
+	{
+	  int q = omp_get_thread_num ();
+	  int tv, bv;
+	  #pragma omp atomic read
+	  tv = T<Q>::t;
+	  #pragma omp atomic read
+	  bv = A::b;
+	  if (i == 16)
+	    {
+	      if (bv != 2 || tv != 4)
+		__builtin_abort ();
+	    }
+	  else
+	    {
+	      if ((bv != 2 && bv != 8) || (tv != 4 && tv != 9))
+		__builtin_abort ();
+	    }
+	  if (!f)
+	    {
+	      if (A::a != 1 || R::r != 3)
+		__builtin_abort ();
+	    }
+	  else if (a != 7 * q || r != 9 * q)
+	    __builtin_abort ();
+	  take (a, r, T<Q>::t, b);
+	  A::a = 7 * q;
+	  R::r = 9 * q;
+	  if (i == 16)
+	    {
+	      #pragma omp atomic write
+	      A::b = 8;
+	      #pragma omp atomic write
+	      T<Q>::t = 9;
+	    }
+	  f = true;
+	}
+    }
+}
+
+int
+main ()
+{
+  A<int> a;
+  a.m1 ();
+  A<int &> b;
+  b.m1 ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/monotonic-1.C b/libgomp/testsuite/libgomp.c++/monotonic-1.C
new file mode 100644
index 00000000000..464bed6b6a5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/monotonic-1.C
@@ -0,0 +1,3 @@
+// { dg-do run }
+
+#include "../libgomp.c/monotonic-1.c"
diff --git a/libgomp/testsuite/libgomp.c++/monotonic-2.C b/libgomp/testsuite/libgomp.c++/monotonic-2.C
new file mode 100644
index 00000000000..41eef0318ea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/monotonic-2.C
@@ -0,0 +1,3 @@
+// { dg-do run }
+
+#include "../libgomp.c/monotonic-2.c"
diff --git a/libgomp/testsuite/libgomp.c++/nonmonotonic-1.C b/libgomp/testsuite/libgomp.c++/nonmonotonic-1.C
new file mode 100644
index 00000000000..678e7654dfb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/nonmonotonic-1.C
@@ -0,0 +1,3 @@
+// { dg-do run }
+
+#include "../libgomp.c/nonmonotonic-1.c"
diff --git a/libgomp/testsuite/libgomp.c++/nonmonotonic-2.C b/libgomp/testsuite/libgomp.c++/nonmonotonic-2.C
new file mode 100644
index 00000000000..e230875f382
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/nonmonotonic-2.C
@@ -0,0 +1,3 @@
+// { dg-do run }
+
+#include "../libgomp.c/nonmonotonic-2.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-3.C b/libgomp/testsuite/libgomp.c++/pr66199-3.C
new file mode 100644
index 00000000000..871dd44d450
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-3.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-3.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-4.C b/libgomp/testsuite/libgomp.c++/pr66199-4.C
new file mode 100644
index 00000000000..bc71a6da043
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-4.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-4.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-5.C b/libgomp/testsuite/libgomp.c++/pr66199-5.C
new file mode 100644
index 00000000000..c7549e8f424
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-5.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-5.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-6.C b/libgomp/testsuite/libgomp.c++/pr66199-6.C
new file mode 100644
index 00000000000..46adbd721b7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-6.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-6.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-7.C b/libgomp/testsuite/libgomp.c++/pr66199-7.C
new file mode 100644
index 00000000000..b03d284ec3d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-7.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-7.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-8.C b/libgomp/testsuite/libgomp.c++/pr66199-8.C
new file mode 100644
index 00000000000..b998ef622c2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-8.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-8.c"
diff --git a/libgomp/testsuite/libgomp.c++/pr66199-9.C b/libgomp/testsuite/libgomp.c++/pr66199-9.C
new file mode 100644
index 00000000000..1f8dfa302cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66199-9.C
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-9.c"
diff --git a/libgomp/testsuite/libgomp.c++/reduction-11.C b/libgomp/testsuite/libgomp.c++/reduction-11.C
new file mode 100644
index 00000000000..588cc5fd3c3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/reduction-11.C
@@ -0,0 +1,121 @@
+// { dg-do run { xfail *-*-* } }
+
+char z[10] = { 0 };
+
+__attribute__((noinline, noclone)) void
+foo (int (*&x)[3][2], int *y, long (&w)[1][2], int s, int t)
+{
+  unsigned long long a[9] = {};
+  short b[5] = {};
+  #pragma omp parallel for reduction(+:x[-1:2][:][0:2], z[t + 2:4]) \
+			   reduction(*:y[-s:3]) reduction(|:a[s + 3:4]) \
+			   reduction(&:w[s + 1:][t:2]) reduction(max:b[2:])
+  for (int i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[1] *= 3;
+      if ((i & 31) == 2)
+	y[2] *= 7;
+      if ((i & 63) == 3)
+	y[3] *= 17;
+      z[i / 32 + 2] += (i & 3);
+      if (i < 4)
+	z[i + 2] += i;
+      a[i / 32 + 2] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      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 (int i = 0; i < 9; i++)
+    if (a[i] != ((i < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 0 || b[1] != 0 || b[2] != 22 || b[3] != 84 || b[4] != 127)
+    __builtin_abort ();
+}
+
+int a3[4][3][2];
+int (*p3)[3][2] = &a3[2];
+int y3[5] = { 0, 1, 1, 1, 0 };
+long w3[1][2] = { ~0L, ~0L };
+short bb[5];
+
+struct S
+{
+  int (*&x)[3][2];
+  int *y;
+  long (&w)[1][2];
+  char z[10];
+  short (&b)[5];
+  unsigned long long a[9];
+  S() : x(p3), y(y3), w(w3), z(), a(), b(bb) {}
+  __attribute__((noinline, noclone)) void foo (int s, int t);
+};
+
+void
+S::foo (int s, int t)
+{
+  #pragma omp parallel for reduction(+:x[-1:2][:][0:2], z[t + 2:4]) \
+			   reduction(*:y[-s:3]) reduction(|:a[s + 3:4]) \
+			   reduction(&:w[s + 1:][t:2]) reduction(max:b[2:])
+  for (int i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[1] *= 3;
+      if ((i & 31) == 2)
+	y[2] *= 7;
+      if ((i & 63) == 3)
+	y[3] *= 17;
+      z[i / 32 + 2] += (i & 3);
+      if (i < 4)
+	z[i + 2] += i;
+      a[i / 32 + 2] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      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;
+    }
+}
+
+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 (*p)[3][2] = &a[2];
+  int y[5] = { 0, 1, 1, 1, 0 };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  long w[1][2] = { ~0L, ~0L };
+  foo (p, y, w, -1, 0);
+  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 ();
+  S s;
+  s.foo (-1, 0);
+  for (int i = 0; i < 9; i++)
+    if (s.a[i] != ((i < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (__builtin_memcmp (a3, a2, sizeof (a3))
+      || __builtin_memcmp (y3, y2, sizeof (y3))
+      || __builtin_memcmp (s.z, z2, sizeof (s.z))
+      || w3[0][0] != ~0x249249L
+      || w3[0][1] != ~0x249249L)
+    __builtin_abort ();
+  if (bb[0] != 0 || bb[1] != 0 || bb[2] != 22 || bb[3] != 84 || bb[4] != 127)
+    __builtin_abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/reduction-12.C b/libgomp/testsuite/libgomp.c++/reduction-12.C
new file mode 100644
index 00000000000..744798a91e8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/reduction-12.C
@@ -0,0 +1,195 @@
+// { dg-do run { xfail *-*-* } }
+
+template <typename T>
+struct A
+{
+  A () { t = 0; }
+  A (T x) { t = x; }
+  A (const A &x) { t = x.t; }
+  ~A () {}
+  T t;
+};
+template <typename T>
+struct M
+{
+  M () { t = 1; }
+  M (T x) { t = x; }
+  M (const M &x) { t = x.t; }
+  ~M () {}
+  T t;
+};
+template <typename T>
+struct B
+{
+  B () { t = ~(T) 0; }
+  B (T x) { t = x; }
+  B (const B &x) { t = x.t; }
+  ~B () {}
+  T t;
+};
+template <typename T>
+void
+add (T &x, T &y)
+{
+  x.t += y.t;
+}
+template <typename T>
+void
+zero (T &x)
+{
+  x.t = 0;
+}
+template <typename T>
+void
+orit (T *x, T *y)
+{
+  y->t |= x->t;
+}
+B<long> bb;
+#pragma omp declare reduction(+:A<int>:omp_out.t += omp_in.t)
+#pragma omp declare reduction(+:A<char>:add (omp_out, omp_in)) initializer(zero (omp_priv))
+#pragma omp declare reduction(*:M<int>:omp_out.t *= omp_in.t) initializer(omp_priv = 1)
+#pragma omp declare reduction(|:A<unsigned long long>:orit (&omp_in, &omp_out))
+#pragma omp declare reduction(&:B<long>:omp_out.t = omp_out.t & omp_in.t) initializer(orit (&omp_priv, &omp_orig))
+#pragma omp declare reduction(maxb:short:omp_out = omp_in > omp_out ? omp_in : omp_out) initializer(omp_priv = -6)
+
+A<char> z[10];
+
+template <int N>
+__attribute__((noinline, noclone)) void
+foo (A<int> (*&x)[3][N], M<int> *y, B<long> (&w)[1][N], int p1, long p2, long p3, int p4,
+     int p5, long p6, short p7, int s, int t)
+{
+  A<unsigned long long> a[p7 + 4];
+  short bb[p7];
+  short (&b)[p7] = bb;
+  for (int i = 0; i < p7; i++)
+    bb[i] = -6;
+  #pragma omp parallel for reduction(+:x[-1:p1 + 1][:p2 + N - 2], z[t + N:p3]) \
+			   reduction(*:y[-s:p4]) reduction(|:a[s + 3:p5 - N + 2]) \
+			   reduction(&:w[s + 1:p6 - 3 + N][t:p6]) reduction(maxb:b[N:])
+  for (int i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[1].t *= 3;
+      if ((i & 31) == N)
+	y[2].t *= 7;
+      if ((i & 63) == 3)
+	y[N + 1].t *= 17;
+      z[i / 32 + 2].t += (i & 3);
+      if (i < 4)
+	z[i + N].t += i;
+      a[i / 32 + 2].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      if ((i % 23) > b[N])
+	b[N] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+  for (int i = 0; i < 9; i++)
+    if (a[i].t != ((i < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (bb[0] != -6 || bb[1] != -6 || bb[N] != 22 || bb[3] != 84 || bb[4] != 127)
+    __builtin_abort ();
+}
+
+A<int> a3[4][3][2];
+A<int> (*p3)[3][2] = &a3[2];
+M<int> y3[5] = { 0, 1, 1, 1, 0 };
+B<long> w3[1][2];
+
+template <int N>
+struct S
+{
+  A<int> (*&x)[3][N];
+  M<int> *y;
+  B<long> (&w)[1][N];
+  A<char> z[10];
+  short b[5];
+  A<unsigned long long> a[9];
+  S() : x(p3), y(y3), w(w3), z(), a(), b() {}
+  __attribute__((noinline, noclone)) void foo (int, long, long, int, int, long, short, int, int);
+};
+
+template <int N>
+void
+S<N>::foo (int p1, long p2, long p3, int p4, int p5, long p6, short p7, int s, int t)
+{
+  #pragma omp parallel for reduction(+:x[-1:p1 + 1][:p2][0:N], z[t + N:p3 + N - 2]) \
+			   reduction(*:y[-s:p4]) reduction(|:a[s + 3:p5]) \
+			   reduction(&:w[s + 1:p6 - 3 + N][t:p6]) reduction(maxb:b[N:])
+  for (int i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[1].t *= 3;
+      if ((i & 31) == N)
+	y[2].t *= 7;
+      if ((i & 63) == 3)
+	y[N + 1].t *= 17;
+      z[i / 32 + 2].t += (i & 3);
+      if (i < 4)
+	z[i + N].t += i;
+      a[i / 32 + 2].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      if ((i % 23) > b[N])
+	b[N] = i % 23;
+      if ((i % 85) > b[3])
+	b[3] = i % 85;
+      if ((i % 192) > b[4])
+	b[4] = i % 192;
+    }
+}
+
+int
+main ()
+{
+  A<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 }}};
+  A<int> (*p)[3][2] = &a[2];
+  M<int> y[5] = { 0, 1, 1, 1, 0 };
+  int y2[5] = { 0, 6561, 2401, 289, 0 };
+  char z2[10] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  B<long> w[1][2];
+  foo<2> (p, y, w, 1, 3L, 4L, 3, 4, 2L, 5, -1, 0);
+  for (int i = 0; i < 4; i++)
+    for (int j = 0; j < 3; j++)
+      for (int k = 0; k < 2; k++)
+	if (a[i][j][k].t != a2[i][j][k])
+	  __builtin_abort ();
+  for (int i = 0; i < 5; i++)
+    if (y[i].t != y2[i])
+      __builtin_abort ();
+  for (int 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 ();
+  S<2> s;
+  s.foo (1, 3L, 4L, 3, 4, 2L, 5, -1, 0);
+  for (int i = 0; i < 9; i++)
+    if (s.a[i].t != ((i < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  for (int i = 0; i < 4; i++)
+    for (int j = 0; j < 3; j++)
+      for (int k = 0; k < 2; k++)
+	if (a3[i][j][k].t != a2[i][j][k])
+	  __builtin_abort ();
+  for (int i = 0; i < 5; i++)
+    if (y3[i].t != y2[i])
+      __builtin_abort ();
+  for (int i = 0; i < 10; i++)
+    if (s.z[i].t != z2[i])
+      __builtin_abort ();
+  if (w3[0][0].t != ~0x249249L || w3[0][1].t != ~0x249249L)
+    __builtin_abort ();
+  if (s.b[0] != 0 || s.b[1] != 0 || s.b[2] != 22
+      || s.b[3] != 84 || s.b[4] != 127)
+    __builtin_abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-13.C b/libgomp/testsuite/libgomp.c++/target-13.C
new file mode 100644
index 00000000000..5279ac02684
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-13.C
@@ -0,0 +1,36 @@
+extern "C" void abort (void);
+
+int g;
+#pragma omp declare target (g)
+
+#pragma omp declare target
+int
+foo (void)
+{
+  static int s;
+  return ++s + g;
+}
+#pragma omp end declare target
+
+int
+bar (void)
+{
+  static int s;
+  #pragma omp declare target to (s)
+  return ++s;
+}
+#pragma omp declare target (bar)
+
+int
+main ()
+{
+  int r;
+  #pragma omp target map(from:r)
+  {
+    r = (foo () == 1) + (bar () == 1);
+    r += (foo () == 2) + (bar () == 2);
+  }
+  if (r != 4)
+    abort ();
+  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..fa0695737e2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-14.C
@@ -0,0 +1,110 @@
+extern "C" void abort ();
+int x;
+
+__attribute__((noinline, noclone)) void
+foo (int &a, int (&b)[10], short &c, long (&d)[5], int n)
+{
+  int err;
+  int &t = x;
+  int y[n + 1];
+  int (&z)[n + 1] = y;
+  for (int i = 0; i < n + 1; i++)
+    z[i] = i + 27;
+  #pragma omp target enter data map (to: z, c) map (alloc: b, t)
+  #pragma omp target update to (b, t)
+  #pragma omp target map (tofrom: a, d) map (from: b, c) map (alloc: t, z) map (from: err)
+  {
+    err = a++ != 7;
+    for (int i = 0; i < 10; i++)
+      {
+	err |= b[i] != 10 - i;
+	b[i] = i - 16;
+	if (i >= 6) continue;
+	err |= z[i] != i + 27;
+	z[i] = 2 * i + 9;
+	if (i == 5) continue;
+	err |= d[i] != 12L + i;
+	d[i] = i + 7;
+      }
+    err |= c != 25;
+    c = 142;
+    err |= t != 8;
+    t = 19;
+  }
+  if (err) abort ();
+  #pragma omp target update from (z, c)
+  #pragma omp target exit data map (from: b, t) map (release: z, c)
+  if (a != 8 || c != 142 || t != 19)
+    abort ();
+  a = 29;
+  c = 149;
+  t = 15;
+  for (int i = 0; i < 10; i++)
+    {
+      if (b[i] != i - 16) abort ();
+      b[i] = i ^ 1;
+      if (i >= 6) continue;
+      if (z[i] != 2 * i + 9) abort ();
+      z[i]++;
+      if (i == 5) continue;
+      if (d[i] != i + 7) abort ();
+      d[i] = 7 - i;
+    }
+  #pragma omp target defaultmap(tofrom: scalar)
+  {
+    err = a++ != 29;
+    for (int i = 0; i < 10; i++)
+      {
+	err |= b[i] != i ^ 1;
+	b[i] = i + 5;
+	if (i >= 6) continue;
+	err |= z[i] != 2 * i + 10;
+	z[i] = 9 - 3 * i;
+	if (i == 5) continue;
+	err |= d[i] != 7L - i;
+	d[i] = i;
+      }
+    err |= c != 149;
+    c = -2;
+    err |= t != 15;
+    t = 155;
+  }
+  if (err || a != 30 || c != -2 || t != 155)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    {
+      if (b[i] != i + 5) abort ();
+      if (i >= 6) continue;
+      if (z[i] != 9 - 3 * i) abort ();
+      z[i]++;
+      if (i == 5) continue;
+      if (d[i] != i) abort ();
+    }
+  #pragma omp target data map (alloc: z)
+  {
+    #pragma omp target update to (z)
+    #pragma omp target map(from: err)
+    {
+      err = 0;
+      for (int i = 0; i < 6; i++)
+	if (z[i] != 10 - 3 * i) err = 1;
+	else z[i] = i;
+    }
+    if (err) abort ();
+    #pragma omp target update from (z)
+  }
+  for (int i = 0; i < 6; i++)
+    if (z[i] != i)
+      abort ();
+}
+
+int
+main ()
+{
+  int a = 7;
+  int b[10] = { 10, 9, 8, 7, 6, 5, 4, 3, 2, 1 };
+  short c = 25;
+  long d[5] = { 12, 13, 14, 15, 16 };
+  x = 8;
+  foo (a, b, c, d, 5);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C
new file mode 100644
index 00000000000..4b320c31229
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-15.C
@@ -0,0 +1,168 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S { char p[64]; int a; int b[2]; long c[4]; int *d; unsigned char &e; char (&f)[2]; short (&g)[4]; int *&h; char q[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (S s)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int sep = 1;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  int err;
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  {
+    err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+    err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+    err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+    err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+    s.a = 35; s.b[0] = 36; s.b[1] = 37;
+    s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+    s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+    s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+    sep = 0;
+  }
+  if (err) abort ();
+  err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+  err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+  err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+  err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+  if (err) abort ();
+  s.a = 50; s.b[0] = 49; s.b[1] = 48;
+  s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+  s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+  s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  {
+    if (!omp_target_is_present (&s.a, d)
+	|| !omp_target_is_present (s.b, d)
+	|| !omp_target_is_present (&s.c[1], d)
+	|| !omp_target_is_present (s.d, d)
+	|| !omp_target_is_present (&s.d[-2], d)
+	|| !omp_target_is_present (&s.e, d)
+	|| !omp_target_is_present (s.f, d)
+	|| !omp_target_is_present (&s.g[1], d)
+	|| !omp_target_is_present (&s.h, d)
+	|| !omp_target_is_present (&s.h[2], d))
+      abort ();
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    {
+      err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+      err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+      err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+      err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+      s.a = 17; s.b[0] = 18; s.b[1] = 19;
+      s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+      s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+      s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+    }
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  }
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+  err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+  err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+  err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+  if (err) abort ();
+  s.a = 33; s.b[0] = 34; s.b[1] = 35;
+  s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+  s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+  s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  {
+    err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+    err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+    err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+    err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+    s.a = 49; s.b[0] = 48; s.b[1] = 47;
+    s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+    s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+    s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+  }
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+  err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+  err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+  err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+  if (err) abort ();
+}
+
+int
+main ()
+{
+  int d[3] = { 18, 19, 20 };
+  unsigned char e = 21;
+  char f[2] = { 22, 23 };
+  short g[4] = { 24, 25, 26, 27 };
+  int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+  int *h = hb + 1;
+  S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+  foo (s);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C
new file mode 100644
index 00000000000..cd102d90594
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-16.C
@@ -0,0 +1,170 @@
+#include <omp.h>
+#include <stdlib.h>
+
+template <typename C, typename I, typename L, typename UC, typename SH>
+struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UC &e; C (&f)[2]; SH (&g)[4]; I *&h; C q[64]; };
+
+template <typename C, typename I, typename L, typename UC, typename SH>
+__attribute__((noinline, noclone)) void
+foo (S<C, I, L, UC, SH> s)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int sep = 1;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  int err;
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  {
+    err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+    err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+    err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+    err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+    s.a = 35; s.b[0] = 36; s.b[1] = 37;
+    s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+    s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+    s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+    sep = 0;
+  }
+  if (err) abort ();
+  err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+  err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+  err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+  err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+  if (err) abort ();
+  s.a = 50; s.b[0] = 49; s.b[1] = 48;
+  s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+  s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+  s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  {
+    if (!omp_target_is_present (&s.a, d)
+	|| !omp_target_is_present (s.b, d)
+	|| !omp_target_is_present (&s.c[1], d)
+	|| !omp_target_is_present (s.d, d)
+	|| !omp_target_is_present (&s.d[-2], d)
+	|| !omp_target_is_present (&s.e, d)
+	|| !omp_target_is_present (s.f, d)
+	|| !omp_target_is_present (&s.g[1], d)
+	|| !omp_target_is_present (&s.h, d)
+	|| !omp_target_is_present (&s.h[2], d))
+      abort ();
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    {
+      err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+      err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+      err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+      err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+      s.a = 17; s.b[0] = 18; s.b[1] = 19;
+      s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+      s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+      s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+    }
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  }
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+  err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+  err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+  err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+  if (err) abort ();
+  s.a = 33; s.b[0] = 34; s.b[1] = 35;
+  s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+  s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+  s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  {
+    err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+    err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+    err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+    err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+    s.a = 49; s.b[0] = 48; s.b[1] = 47;
+    s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+    s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+    s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+  }
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+  err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+  err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+  err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+  if (err) abort ();
+}
+
+int
+main ()
+{
+  int d[3] = { 18, 19, 20 };
+  unsigned char e = 21;
+  char f[2] = { 22, 23 };
+  short g[4] = { 24, 25, 26, 27 };
+  int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+  int *h = hb + 1;
+  S<char, int, long, unsigned char, short> s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+  foo (s);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C
new file mode 100644
index 00000000000..d81ff19a411
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-17.C
@@ -0,0 +1,173 @@
+#include <omp.h>
+#include <stdlib.h>
+
+template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR>
+struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UCR e; CAR f; SH g; IPR h; C q[64]; };
+
+template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR>
+__attribute__((noinline, noclone)) void
+foo (S<C, I, L, UCR, CAR, SH, IPR> s)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int sep = 1;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  int err;
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  {
+    err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+    err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+    err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+    err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+    s.a = 35; s.b[0] = 36; s.b[1] = 37;
+    s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+    s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+    s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+    sep = 0;
+  }
+  if (err) abort ();
+  err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+  err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+  err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+  err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+  if (err) abort ();
+  s.a = 50; s.b[0] = 49; s.b[1] = 48;
+  s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+  s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+  s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  {
+    if (!omp_target_is_present (&s.a, d)
+	|| !omp_target_is_present (s.b, d)
+	|| !omp_target_is_present (&s.c[1], d)
+	|| !omp_target_is_present (s.d, d)
+	|| !omp_target_is_present (&s.d[-2], d)
+	|| !omp_target_is_present (&s.e, d)
+	|| !omp_target_is_present (s.f, d)
+	|| !omp_target_is_present (&s.g[1], d)
+	|| !omp_target_is_present (&s.h, d)
+	|| !omp_target_is_present (&s.h[2], d))
+      abort ();
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    {
+      err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+      err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+      err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+      err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+      s.a = 17; s.b[0] = 18; s.b[1] = 19;
+      s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+      s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+      s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+    }
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  }
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+  err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+  err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+  err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+  if (err) abort ();
+  s.a = 33; s.b[0] = 34; s.b[1] = 35;
+  s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+  s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+  s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  {
+    err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+    err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+    err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+    err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+    s.a = 49; s.b[0] = 48; s.b[1] = 47;
+    s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+    s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+    s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+  }
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d)
+      || !omp_target_is_present (&s.e, d)
+      || !omp_target_is_present (s.f, d)
+      || !omp_target_is_present (&s.g[1], d)
+      || !omp_target_is_present (&s.h, d)
+      || !omp_target_is_present (&s.h[2], d))
+    abort ();
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)
+	  || omp_target_is_present (&s.e, d)
+	  || omp_target_is_present (s.f, d)
+	  || omp_target_is_present (&s.g[1], d)
+	  || omp_target_is_present (&s.h, d)
+	  || omp_target_is_present (&s.h[2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+  err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+  err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+  err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+  if (err) abort ();
+}
+
+int
+main ()
+{
+  int d[3] = { 18, 19, 20 };
+  unsigned char e = 21;
+  char f[2] = { 22, 23 };
+  short g[4] = { 24, 25, 26, 27 };
+  int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+  int *h = hb + 1;
+  typedef char (&CAR)[2];
+  typedef short (&SH)[4];
+  S<char, int, long, unsigned char &, CAR, SH, int *&> s
+    = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+  foo (s);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-18.C b/libgomp/testsuite/libgomp.c++/target-18.C
new file mode 100644
index 00000000000..f1085b14022
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-18.C
@@ -0,0 +1,167 @@
+extern "C" void abort ();
+
+__attribute__((noinline, noclone)) void
+foo (int *&p, int *&q, int *&r, int n, int m)
+{
+  int i, err, *s = r;
+  int sep = 1;
+  #pragma omp target map(to:sep)
+  sep = 0;
+  #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 (with nothing mapped
+       after 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)
+	  err = 1;
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || 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)
+	  err = 1;
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || 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)
+	  err = 1;
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || 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)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (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)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (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)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (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;
+  int *p = a + 1, *q = a + 9, *r = a + 10;
+  foo (p, q, r, 0, 1);
+  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..afa6e68d5cc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-19.C
@@ -0,0 +1,59 @@
+extern "C" void abort ();
+struct S { char a[64]; int (&r)[2]; char b[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (S s, int (&t)[3], int z)
+{
+  int err, sep = 1;
+  // Test that implicit mapping of reference to array does NOT
+  // behave like zero length array sections.  s.r can't be used
+  // implicitly, as that means implicit mapping of the whole s
+  // and trying to dereference the references in there is unspecified.
+  #pragma omp target map(from: err) map(to: sep)
+  {
+    err = t[0] != 1 || t[1] != 2 || t[2] != 3;
+    sep = 0;
+  }
+  if (err) abort ();
+  // But explicit zero length array section mapping does.
+  #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0])
+  {
+    if (sep)
+      err = s.r != (int *) 0 || t != (int *) 0;
+    else
+      err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+  }
+  if (err) abort ();
+  // Similarly zero length array section, but unknown at compile time.
+  #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z])
+  {
+    if (sep)
+      err = s.r != (int *) 0 || t != (int *) 0;
+    else
+      err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+  }
+  if (err) abort ();
+  #pragma omp target enter data map (to: s.r, t)
+  // But when already mapped, it binds to existing mappings.
+  #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0])
+  {
+    err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+    sep = 0;
+  }
+  if (err) abort ();
+  #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z])
+  {
+    err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+    sep = 0;
+  }
+  if (err) abort ();
+}
+
+int
+main ()
+{
+  int t[3] = { 1, 2, 3 };
+  int r[2] = { 6, 7 };
+  S s = { {}, r, {} };
+  foo (s, t, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c/monotonic-1.c b/libgomp/testsuite/libgomp.c/monotonic-1.c
new file mode 100644
index 00000000000..c3c0100f052
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/monotonic-1.c
@@ -0,0 +1,303 @@
+/* { dg-do run } */
+
+#ifndef MONOTONIC_TYPE
+#include <omp.h>
+#include <stdlib.h>
+#define MONOTONIC_TYPE int
+#define MONOTONIC_UNDEF -1
+#define MONOTONIC_END(n) n
+#endif
+
+int
+main ()
+{
+  MONOTONIC_TYPE i;
+  #pragma omp parallel
+  {
+    int cnt = omp_get_num_threads ();
+    int thr = omp_get_thread_num ();
+    MONOTONIC_TYPE l = MONOTONIC_UNDEF;
+    int c = 0;
+    int n = 0;
+    #pragma omp for nowait schedule(static, 5)
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if (cnt == 1)
+	      abort ();
+	    if (n != 5)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	if (n == 1)
+	  {
+	    if ((i % 5) != 0)
+	      abort ();
+	    if ((i / 5) % cnt != thr)
+	      abort ();
+	  }
+	l = i;
+      }
+    if (cnt == 1)
+      {
+	if (n != 73 || l != 73 - 1 || c != 1)
+	  abort ();
+      }
+    else if (thr > 73 / 5)
+      {
+	if (l != MONOTONIC_UNDEF || c != 0 || n != 0)
+	  abort ();
+      }
+    else if (thr == 73 / 5)
+      {
+	if (l != 73 - 1 || c != 1 || n != 73 % 5)
+	  abort ();
+      }
+    else if (c == 0)
+      abort ();
+    else if (l == 73 - 1)
+      {
+	if (thr != (73 / 5) % cnt || n != 73 % 5)
+	  abort ();
+      }
+    else if ((n % 5) != 0)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for schedule( monotonic: static, 7) nowait
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if (cnt == 1)
+	      abort ();
+	    if (n != 7)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	if (n == 1)
+	  {
+	    if ((i % 7) != 0)
+	      abort ();
+	    if ((i / 7) % cnt != thr)
+	      abort ();
+	  }
+	l = i;
+      }
+    if (cnt == 1)
+      {
+	if (n != 73 || l != 73 - 1 || c != 1)
+	  abort ();
+      }
+    else if (thr > 73 / 7)
+      {
+	if (l != MONOTONIC_UNDEF || c != 0 || n != 0)
+	  abort ();
+      }
+    else if (thr == 73 / 7)
+      {
+	if (l != 73 - 1 || c != 1 || n != 73 % 7)
+	  abort ();
+      }
+    else if (c == 0)
+      abort ();
+    else if (l == 73 - 1)
+      {
+	if (thr != (73 / 7) % cnt || n != 73 % 7)
+	  abort ();
+      }
+    else if ((n % 7) != 0)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for nowait schedule(static)
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  abort ();
+	l = i;
+      }
+    if (c > 1)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for nowait schedule(monotonic,simd:static)
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  abort ();
+	l = i;
+      }
+    if (c > 1)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for schedule(monotonic : dynamic, 5) nowait
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if ((n % 5) != 0 || n == 0)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	l = i;
+      }
+    if (l == 73 - 1)
+      {
+	if (n % 5 != 73 % 5)
+	  abort ();
+      }
+    else if (l == MONOTONIC_UNDEF)
+      {
+	if (n != 0 || c != 0)
+	  abort ();
+      }
+    else if ((n % 5) != 0 || n == 0)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for nowait schedule(dynamic, 7) ordered(1)
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if ((n % 7) != 0 || n == 0)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	#pragma omp ordered depend(source)
+	if (MONOTONIC_UNDEF > 0)
+	  {
+	    #pragma omp ordered depend(sink: i)
+	  }
+	else
+	  {
+	    #pragma omp ordered depend(sink: i - 1)
+	  }
+	l = i;
+      }
+    if (l == 73 - 1)
+      {
+	if (n % 7 != 73 % 7)
+	  abort ();
+      }
+    else if (l == MONOTONIC_UNDEF)
+      {
+	if (n != 0 || c != 0)
+	  abort ();
+      }
+    else if ((n % 7) != 0 || n == 0)
+      abort ();
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for schedule (monotonic :guided , 7) nowait
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if (n < 7)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	l = i;
+      }
+    l = MONOTONIC_UNDEF;
+    c = 0;
+    n = 0;
+    #pragma omp for nowait schedule(guided, 7) ordered
+    for (i = 0; i < MONOTONIC_END (73); i++)
+      {
+	if (l == MONOTONIC_UNDEF)
+	  {
+	    n = 1;
+	    c++;
+	  }
+	else if (l == i - 1)
+	  n++;
+	else
+	  {
+	    if (l >= i)
+	      abort ();
+	    if (n < 7)
+	      abort ();
+	    n = 1;
+	    c++;
+	  }
+	#pragma omp ordered
+	  l = i;
+      }
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/monotonic-2.c b/libgomp/testsuite/libgomp.c/monotonic-2.c
new file mode 100644
index 00000000000..b886a24c99a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/monotonic-2.c
@@ -0,0 +1,11 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+#define MONOTONIC_TYPE unsigned long long
+#define MONOTONIC_UNDEF -1ULL
+#define MONOTONIC_END(n) n + v
+
+volatile int v;
+
+#include "monotonic-1.c"
diff --git a/libgomp/testsuite/libgomp.c/nonmonotonic-1.c b/libgomp/testsuite/libgomp.c/nonmonotonic-1.c
new file mode 100644
index 00000000000..fc5e93fe53a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/nonmonotonic-1.c
@@ -0,0 +1,53 @@
+/* { dg-do run } */
+
+#ifndef NONMONOTONIC_TYPE
+#include <omp.h>
+#include <stdlib.h>
+#define NONMONOTONIC_TYPE int
+#define NONMONOTONIC_END(n) n
+#endif
+
+int a[73];
+
+int
+main ()
+{
+  NONMONOTONIC_TYPE i;
+  #pragma omp parallel for schedule(nonmonotonic: dynamic)
+  for (i = 0; i < NONMONOTONIC_END (73); i++)
+    a[i]++;
+  #pragma omp parallel for schedule(nonmonotonic: dynamic, 5)
+  for (i = 0; i < NONMONOTONIC_END (73); i++)
+    a[i]++;
+  #pragma omp parallel for schedule(nonmonotonic: guided)
+  for (i = 0; i < NONMONOTONIC_END (73); i++)
+    a[i]++;
+  #pragma omp parallel for schedule(nonmonotonic: guided, 7)
+  for (i = 0; i < NONMONOTONIC_END (73); i++)
+    a[i]++;
+  #pragma omp parallel
+  {
+    int cnt = omp_get_num_threads ();
+    int thr = omp_get_thread_num ();
+    if (thr < 73)
+      a[thr]++;
+    #pragma omp barrier
+    #pragma omp for schedule(nonmonotonic: dynamic)
+    for (i = 0; i < NONMONOTONIC_END (73); i++)
+      a[i]++;
+    #pragma omp for schedule(nonmonotonic: dynamic, 7)
+    for (i = 0; i < NONMONOTONIC_END (73); i++)
+      a[i]++;
+    #pragma omp for schedule(nonmonotonic: guided)
+    for (i = 0; i < NONMONOTONIC_END (73); i++)
+      a[i]++;
+    #pragma omp for schedule(nonmonotonic: guided, 5)
+    for (i = 0; i < NONMONOTONIC_END (73); i++)
+      a[i]++;
+    #pragma omp single private (i)
+    for (i = 0; i < 73; i++)
+      if (a[i] != 8 + (i < cnt))
+	abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/nonmonotonic-2.c b/libgomp/testsuite/libgomp.c/nonmonotonic-2.c
new file mode 100644
index 00000000000..c56fa1dc42a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/nonmonotonic-2.c
@@ -0,0 +1,10 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+#define NONMONOTONIC_TYPE unsigned long long
+#define NONMONOTONIC_END(n) n + v
+
+volatile int v;
+
+#include "nonmonotonic-1.c"
diff --git a/libgomp/testsuite/libgomp.c/pr66199-5.c b/libgomp/testsuite/libgomp.c/pr66199-5.c
new file mode 100644
index 00000000000..7ade43c4c9f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-5.c
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams distribute parallel for simd 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 target map(from: d, e)
+  #pragma omp 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];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute parallel for simd 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;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute parallel for simd 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];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-6.c b/libgomp/testsuite/libgomp.c/pr66199-6.c
new file mode 100644
index 00000000000..6790f7dc270
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-6.c
@@ -0,0 +1,42 @@
+/* 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)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp 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;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp 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];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-7.c b/libgomp/testsuite/libgomp.c/pr66199-7.c
new file mode 100644
index 00000000000..752367e3ac5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-7.c
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams distribute simd 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 target map(from: d, e)
+  #pragma omp teams distribute 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];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute simd 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;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute simd 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];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-8.c b/libgomp/testsuite/libgomp.c/pr66199-8.c
new file mode 100644
index 00000000000..102af0691ec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-8.c
@@ -0,0 +1,70 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams default(none) shared(a, b, d, u, v, w)
+  #pragma omp distribute simd firstprivate (a, b)
+  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 target map(from: d, e)
+  #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  #pragma omp distribute simd linear(d) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  #pragma omp distribute simd firstprivate (a1, b1, a2, b2) 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;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) firstprivate (a1, b1, a2, b2) shared(d1, d2, u, v, w)
+  #pragma omp distribute simd 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 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr66199-9.c b/libgomp/testsuite/libgomp.c/pr66199-9.c
new file mode 100644
index 00000000000..554a4765b4e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66199-9.c
@@ -0,0 +1,43 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  #pragma omp distribute lastprivate(d, e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  #pragma omp distribute firstprivate (a1, b1, a2, b2) 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 (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/reduction-11.c b/libgomp/testsuite/libgomp.c/reduction-11.c
new file mode 100644
index 00000000000..594dc673e03
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-11.c
@@ -0,0 +1,62 @@
+/* { dg-do run { xfail *-*-* } } */
+
+char z[10] = { 0 };
+
+__attribute__((noinline, noclone)) void
+foo (int (*x)[3][2], int *y, long w[1][2], int s, int t)
+{
+  unsigned long long a[9] = {};
+  short b[5] = {};
+  int i;
+  #pragma omp parallel for reduction(+:x[-1:2][:][0:2], z[t + 2:4]) \
+			   reduction(*:y[-s:3]) reduction(|:a[s + 3:4]) \
+			   reduction(&:w[s + 1:1][t:2]) reduction(max:b[2:])
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[1] *= 3;
+      if ((i & 31) == 2)
+	y[2] *= 7;
+      if ((i & 63) == 3)
+	y[3] *= 17;
+      z[i / 32 + 2] += (i & 3);
+      if (i < 4)
+	z[i + 2] += i;
+      a[i / 32 + 2] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      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 < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 0 || b[1] != 0 || 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] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  long w[1][2] = { ~0L, ~0L };
+  foo (&a[2], y, w, -1, 0);
+  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-12.c b/libgomp/testsuite/libgomp.c/reduction-12.c
new file mode 100644
index 00000000000..1c9c4f3e3dd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-12.c
@@ -0,0 +1,96 @@
+/* { dg-do run { xfail *-*-* } } */
+
+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 s, int t)
+{
+  struct C a[9] = {};
+  short b[5] = {};
+  int i;
+  #pragma omp parallel for reduction(+:x[-1:2][:][0:2], z[t + 2:4]) \
+			   reduction(*:y[-s:3]) reduction(|:a[s + 3:4]) \
+			   reduction(&:w[s + 1:1][t:2]) reduction(maxb:b[2:])
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[1].t *= 3;
+      if ((i & 31) == 2)
+	y[2].t *= 7;
+      if ((i & 63) == 3)
+	y[3].t *= 17;
+      z[i / 32 + 2].t += (i & 3);
+      if (i < 4)
+	z[i + 2].t += i;
+      a[i / 32 + 2].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      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 < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != 0 || b[1] != 0 || 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] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  struct D w[1][2] = { { { ~0L }, { ~0L } } };
+  foo (&a[2], y, w, -1, 0);
+  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-13.c b/libgomp/testsuite/libgomp.c/reduction-13.c
new file mode 100644
index 00000000000..99468b2e344
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-13.c
@@ -0,0 +1,67 @@
+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, int s, int t)
+{
+  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[-1:p1 + 1][:p2], z[t + 2:p3]) \
+			   reduction(*:y[-s:p4]) reduction(|:a[s + 3:p5]) \
+			   reduction(&:w[s + 1:p6 - 1][t:p6]) reduction(max:b[2:])
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1] += i;
+      if ((i & 15) == 1)
+	y[1] *= 3;
+      if ((i & 31) == 2)
+	y[2] *= 7;
+      if ((i & 63) == 3)
+	y[3] *= 17;
+      z[i / 32 + 2] += (i & 3);
+      if (i < 4)
+	z[i + 2] += i;
+      a[i / 32 + 2] |= 1ULL << (i & 30);
+      w[0][i & 1] &= ~(1L << (i / 17 * 3));
+      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 < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != -6 || b[1] != -6 || 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] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  long w[1][2] = { ~0L, ~0L };
+  foo (&a[2], y, w, 1, 3L, 4L, 3, 4, 2L, 5, -1, 0);
+  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-14.c b/libgomp/testsuite/libgomp.c/reduction-14.c
new file mode 100644
index 00000000000..4c92ff0fb6f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-14.c
@@ -0,0 +1,101 @@
+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, int s, int t)
+{
+  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[-1:p1 + 1][:p2], z[t + 2:p3]) \
+			   reduction(*:y[-s:p4]) reduction(|:a[s + 3:p5]) \
+			   reduction(&:w[s + 1:p6 - 1][t:p6]) reduction(maxb:b[2:])
+  for (i = 0; i < 128; i++)
+    {
+      x[i / 64 - 1][i % 3][(i / 4) & 1].t += i;
+      if ((i & 15) == 1)
+	y[1].t *= 3;
+      if ((i & 31) == 2)
+	y[2].t *= 7;
+      if ((i & 63) == 3)
+	y[3].t *= 17;
+      z[i / 32 + 2].t += (i & 3);
+      if (i < 4)
+	z[i + 2].t += i;
+      a[i / 32 + 2].t |= 1ULL << (i & 30);
+      w[0][i & 1].t &= ~(1L << (i / 17 * 3));
+      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 < 6 && i >= 2) ? 0x55555555ULL : 0))
+      __builtin_abort ();
+  if (b[0] != -6 || b[1] != -6 || 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] = { 0, 0, 48, 49, 50, 51, 0, 0, 0, 0 };
+  struct D w[1][2] = { { { ~0L }, { ~0L } } };
+  foo (&a[2], y, w, 1, 3L, 4L, 3, 4, 2L, 5, -1, 0);
+  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-15.c b/libgomp/testsuite/libgomp.c/reduction-15.c
new file mode 100644
index 00000000000..50f7e2422ca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reduction-15.c
@@ -0,0 +1,56 @@
+extern void abort (void);
+int a[16], b[16], c[16], d[5][2];
+
+__attribute__((noinline, noclone)) void
+foo (int x, int y)
+{
+  int i;
+  #pragma omp for schedule (static, 1) reduction (+:a[:3])
+  for (i = 0; i < 64; i++)
+    {
+      a[0] += i;
+      a[1] += 2 * i;
+      a[2] += 3 * i;
+    }
+  #pragma omp for schedule (guided) reduction (+:b[4:3])
+  for (i = 0; i < 64; i++)
+    {
+      b[4] += i;
+      b[5] += 2 * i;
+      b[6] += 3 * i;
+    }
+  #pragma omp for schedule (static) reduction (+:c[x:4])
+  for (i = 0; i < 64; i++)
+    {
+      c[9] += i;
+      c[10] += 2 * i;
+      c[11] += 3 * i;
+      c[12] += 4 * i;
+    }
+  #pragma omp for reduction (+:d[x - 8:2][y:])
+  for (i = 0; i < 64; i++)
+    {
+      d[1][0] += i;
+      d[1][1] += 2 * i;
+      d[2][0] += 3 * i;
+      d[2][1] += 4 * i;
+    }
+}
+
+int
+main ()
+{
+  int i;
+  #pragma omp parallel
+  foo (9, 0);
+  for (i = 0; i < 16; i++)
+    if (a[i] != (i < 3 ? 64 * 63 / 2 * (i + 1) : 0)
+	|| b[i] != ((i >= 4 && i < 7) ? 64 * 63 / 2 * (i - 3) : 0)
+	|| c[i] != ((i >= 9 && i < 13) ? 64 * 63 / 2 * (i - 8) : 0))
+      abort ();
+  for (i = 0; i < 5; i++)
+    if (d[i][0] != ((i && i <= 2) ? 64 * 63 / 2 * (2 * i - 1) : 0)
+	|| d[i][1] != ((i && i <= 2) ? 64 * 63 / 2 * (2 * i) : 0))
+      abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-12.c b/libgomp/testsuite/libgomp.c/target-12.c
index e6b009463ad..b439e56577c 100644
--- a/libgomp/testsuite/libgomp.c/target-12.c
+++ b/libgomp/testsuite/libgomp.c/target-12.c
@@ -41,7 +41,7 @@ main ()
 
       if (omp_target_is_present (q, d) != 1
 	  || omp_target_is_present (&q[32], d) != 1
-	  || omp_target_is_present (&q[128], d) != 1)
+	  || omp_target_is_present (&q[127], d) != 1)
 	abort ();
 
       if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
diff --git a/libgomp/testsuite/libgomp.c/target-17.c b/libgomp/testsuite/libgomp.c/target-17.c
index 4a762012eaf..4990826ab46 100644
--- a/libgomp/testsuite/libgomp.c/target-17.c
+++ b/libgomp/testsuite/libgomp.c/target-17.c
@@ -37,58 +37,6 @@ foo (int n)
   }
   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
diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c
index 710c5078ff6..2505cafca9f 100644
--- a/libgomp/testsuite/libgomp.c/target-19.c
+++ b/libgomp/testsuite/libgomp.c/target-19.c
@@ -1,21 +1,29 @@
 extern void abort (void);
 
-void
+__attribute__((noinline, noclone)) void
 foo (int *p, int *q, int *r, int n, int m)
 {
   int i, err, *s = r;
+  int sep = 1;
+  #pragma omp target map(to:sep)
+  sep = 0;
   #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.  */
+       already mapped range, q to the end of it (with nothing mapped
+       after 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)
+	if (p[i] != i + 1)
 	  err = 1;
-      if (p + 8 != q || (r != (int *) 0 && r != s))
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || r != s)
 	err = 1;
     }
     if (err)
@@ -25,9 +33,14 @@ foo (int *p, int *q, int *r, int n, int m)
     {
       err = 0;
       for (i = 0; i < 8; i++)
-	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	if (p[i] != i + 1)
 	  err = 1;
-      if (p + 8 != q || (r != (int *) 0 && r != s))
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || r != s)
 	err = 1;
     }
     if (err)
@@ -38,9 +51,14 @@ foo (int *p, int *q, int *r, int n, int m)
     {
       err = 0;
       for (i = 0; i < 8; i++)
-	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	if (p[i] != i + 1)
 	  err = 1;
-      if (p + 8 != q || (r != (int *) 0 && r != s))
+      if (sep)
+	{
+	  if (q != (int *) 0 || r != (int *) 0)
+	    err = 1;
+	}
+      else if (p + 8 != q || r != s)
 	err = 1;
     }
     if (err)
@@ -69,7 +87,14 @@ foo (int *p, int *q, int *r, int n, int m)
 	for (i = 0; i < 8; i++)
 	  if (p[i] != i + 1)
 	    err = 1;
-	if (q[0] != 9 || r != q + 1)
+	if (q[0] != 9)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (r != q + 1)
 	  err = 1;
       }
       if (err)
@@ -81,7 +106,14 @@ foo (int *p, int *q, int *r, int n, int m)
 	for (i = 0; i < 8; i++)
 	  if (p[i] != i + 1)
 	    err = 1;
-	if (q[0] != 9 || r != q + 1)
+	if (q[0] != 9)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (r != q + 1)
 	  err = 1;
       }
       if (err)
@@ -94,7 +126,14 @@ foo (int *p, int *q, int *r, int n, int m)
 	for (i = 0; i < 8; i++)
 	  if (p[i] != i + 1)
 	    err = 1;
-	if (q[0] != 9 || r != q + 1)
+	if (q[0] != 9)
+	  err = 1;
+	else if (sep)
+	  {
+	    if (r != (int *) 0)
+	      err = 1;
+	  }
+	else if (r != q + 1)
 	  err = 1;
       }
       if (err)
diff --git a/libgomp/testsuite/libgomp.c/target-28.c b/libgomp/testsuite/libgomp.c/target-28.c
new file mode 100644
index 00000000000..96e9e058a21
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-28.c
@@ -0,0 +1,36 @@
+extern void abort (void);
+
+int g;
+#pragma omp declare target (g)
+
+#pragma omp declare target
+int
+foo (void)
+{
+  static int s;
+  return ++s + g;
+}
+#pragma omp end declare target
+
+int
+bar (void)
+{
+  static int s;
+  #pragma omp declare target to (s)
+  return ++s;
+}
+#pragma omp declare target (bar)
+
+int
+main ()
+{
+  int r;
+  #pragma omp target map(from:r)
+  {
+    r = (foo () == 1) + (bar () == 1);
+    r += (foo () == 2) + (bar () == 2);
+  }
+  if (r != 4)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c
new file mode 100644
index 00000000000..e5095a1b6b8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-29.c
@@ -0,0 +1,112 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S { char p[64]; int a; int b[2]; long c[4]; int *d; char q[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (struct S s)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int sep = 1;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  int err;
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
+  {
+    err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+    err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+    s.a = 35; s.b[0] = 36; s.b[1] = 37;
+    s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+    sep = 0;
+  }
+  if (err) abort ();
+  err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+  err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+  if (err) abort ();
+  s.a = 50; s.b[0] = 49; s.b[1] = 48;
+  s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)))
+    abort ();
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  {
+    if (!omp_target_is_present (&s.a, d)
+	|| !omp_target_is_present (s.b, d)
+	|| !omp_target_is_present (&s.c[1], d)
+	|| !omp_target_is_present (s.d, d)
+	|| !omp_target_is_present (&s.d[-2], d))
+      abort ();
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+    {
+      err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+      err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+      s.a = 17; s.b[0] = 18; s.b[1] = 19;
+      s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+    }
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
+  }
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+  err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+  if (err) abort ();
+  s.a = 33; s.b[0] = 34; s.b[1] = 35;
+  s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d))
+    abort ();
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+  {
+    err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+    err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+    s.a = 49; s.b[0] = 48; s.b[1] = 47;
+    s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+  }
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
+  if (!omp_target_is_present (&s.a, d)
+      || !omp_target_is_present (s.b, d)
+      || !omp_target_is_present (&s.c[1], d)
+      || !omp_target_is_present (s.d, d)
+      || !omp_target_is_present (&s.d[-2], d))
+    abort ();
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
+  if (sep
+      && (omp_target_is_present (&s.a, d)
+	  || omp_target_is_present (s.b, d)
+	  || omp_target_is_present (&s.c[1], d)
+	  || omp_target_is_present (s.d, d)
+	  || omp_target_is_present (&s.d[-2], d)))
+    abort ();
+  if (err) abort ();
+  err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+  err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+  if (err) abort ();
+}
+
+int
+main ()
+{
+  int d[3] = { 18, 19, 20 };
+  struct S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, {} };
+  foo (s);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-30.c b/libgomp/testsuite/libgomp.c/target-30.c
new file mode 100644
index 00000000000..a309f7d1a21
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-30.c
@@ -0,0 +1,24 @@
+extern void abort (void);
+
+#pragma omp declare target
+int v = 6;
+#pragma omp end declare target
+
+int
+main ()
+{
+  #pragma omp target /* predetermined map(tofrom: v) */
+  v++;
+  #pragma omp target update from (v)
+  if (v != 7)
+    abort ();
+  #pragma omp parallel private (v) num_threads (1)
+  {
+    #pragma omp target /* predetermined firstprivate(v) */
+    v++;
+  }
+  #pragma omp target update from (v)
+  if (v != 7)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-teams-1.c b/libgomp/testsuite/libgomp.c/target-teams-1.c
new file mode 100644
index 00000000000..8f591e02c48
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-teams-1.c
@@ -0,0 +1,152 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int v = 6;
+
+void
+bar (long *x, long *y)
+{
+  *x += 2;
+  *y += 3;
+}
+
+int
+baz (void)
+{
+  return 5;
+}
+
+#pragma omp declare target to (bar, baz, v)
+
+__attribute__((noinline, noclone)) void
+foo (int a, int b, long c, long d)
+{
+  int err;
+  if (omp_get_num_teams () != 1)
+    abort ();
+  /* The OpenMP 4.5 spec says that these expressions are evaluated before
+     target region on combined target teams, so those cases are always
+     fine.  */
+  #pragma omp target map(from: err)
+  err = omp_get_num_teams () != 1;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
+  if (err)
+    abort ();
+  #pragma omp target teams map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (4)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > 4;
+  if (err)
+    abort ();
+  #pragma omp target teams num_teams (4) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > 4;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams thread_limit (7)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_thread_limit () > 7;
+  if (err)
+    abort ();
+  #pragma omp target teams thread_limit (7) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_thread_limit () > 7;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (4) thread_limit (8)
+  {
+    {
+      err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	    || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target teams num_teams (4) thread_limit (8) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (a) thread_limit (b)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > a || omp_get_thread_limit () > b;
+  if (err)
+    abort ();
+  #pragma omp target teams num_teams (a) thread_limit (b) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > a || omp_get_thread_limit () > b;
+  if (err)
+    abort ();
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+  if (err)
+    abort ();
+  #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+  if (err)
+    abort ();
+  #pragma omp target map (always, to: c, d) map(from: err)
+  #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+  if (err)
+    abort ();
+  #pragma omp target data map (to: c, d)
+  {
+    #pragma omp target defaultmap (tofrom: scalar)
+    bar (&c, &d);
+    /* This is one of the cases which can't be generally optimized,
+       the c and d are (or could be) already mapped and whether
+       their device and original values match is unclear.  */
+    #pragma omp target map (to: c, d) map(from: err)
+    #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+    err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	  || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+    if (err)
+      abort ();
+  }
+  /* This can't be optimized, there are function calls inside of
+     target involved.  */
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
+  if (err)
+    abort ();
+  #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
+  if (err)
+    abort ();
+  /* This one can't be optimized, as v might have different value between
+     host and target.  */
+  #pragma omp target map(from: err)
+  #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
+  err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+	|| omp_get_num_teams () > v + 1 || omp_get_thread_limit () > v - 1;
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (3, 5, 7, 9);
+  return 0;
+}