builtin-types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR, [...]): New.

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

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

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

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

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

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

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

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

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

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

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

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

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

From-SVN: r228777

1 files changed, 1110 insertions, 98 deletions

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 25a81f6..4a9f7fd 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -91,6 +91,8 @@ enum gimplify_omp_var_data
   /* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference.  */
   GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384,
 
+  GOVD_MAP_0LEN_ARRAY = 32768,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -110,7 +112,11 @@ enum omp_region_type
   /* Data region.  */
   ORT_TARGET_DATA = 16,
   /* Data region with offloading.  */
-  ORT_TARGET = 32
+  ORT_TARGET = 32,
+  ORT_COMBINED_TARGET = 33,
+  /* Dummy OpenMP region, used to disable expansion of
+     DECL_VALUE_EXPRs in taskloop pre body.  */
+  ORT_NONE = 64
 };
 
 /* Gimplify hashtable helper.  */
@@ -147,11 +153,16 @@ struct gimplify_omp_ctx
   struct gimplify_omp_ctx *outer_context;
   splay_tree variables;
   hash_set<tree> *privatized_types;
+  /* Iteration variables in an OMP_FOR.  */
+  vec<tree> loop_iter_var;
   location_t location;
   enum omp_clause_default_kind default_kind;
   enum omp_region_type region_type;
   bool combined_loop;
   bool distribute;
+  bool target_map_scalars_firstprivate;
+  bool target_map_pointers_as_0len_arrays;
+  bool target_firstprivatize_array_bases;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -382,6 +393,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
 {
   splay_tree_delete (c->variables);
   delete c->privatized_types;
+  c->loop_iter_var.release ();
   XDELETE (c);
 }
 
@@ -1070,7 +1082,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 	  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
 
 	  /* Mark variable as local.  */
-	  if (ctx && !DECL_EXTERNAL (t)
+	  if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t)
 	      && (! DECL_SEEN_IN_BIND_EXPR_P (t)
 		  || splay_tree_lookup (ctx->variables,
 					(splay_tree_key) t) == NULL))
@@ -2255,8 +2267,7 @@ maybe_fold_stmt (gimple_stmt_iterator *gsi)
 {
   struct gimplify_omp_ctx *ctx;
   for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
-    if (ctx->region_type == ORT_TARGET
-	|| (ctx->region_type & (ORT_PARALLEL | ORT_TASK)) != 0)
+    if ((ctx->region_type & (ORT_TARGET | ORT_PARALLEL | ORT_TASK)) != 0)
       return false;
   return fold_stmt (gsi);
 }
@@ -4467,6 +4478,13 @@ is_gimple_stmt (tree t)
     case OMP_ORDERED:
     case OMP_CRITICAL:
     case OMP_TASK:
+    case OMP_TARGET:
+    case OMP_TARGET_DATA:
+    case OMP_TARGET_UPDATE:
+    case OMP_TARGET_ENTER_DATA:
+    case OMP_TARGET_EXIT_DATA:
+    case OMP_TASKLOOP:
+    case OMP_TEAMS:
       /* These are always void.  */
       return true;
 
@@ -5536,7 +5554,7 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
 {
   splay_tree_node n;
 
-  if (decl == NULL || !DECL_P (decl))
+  if (decl == NULL || !DECL_P (decl) || ctx->region_type == ORT_NONE)
     return;
 
   do
@@ -5551,8 +5569,13 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
 	  else
 	    return;
 	}
-      else if (ctx->region_type == ORT_TARGET)
-	omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+      else if ((ctx->region_type & ORT_TARGET) != 0)
+	{
+	  if (ctx->target_map_scalars_firstprivate)
+	    omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
+	  else
+	    omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+	}
       else if (ctx->region_type != ORT_WORKSHARE
 	       && ctx->region_type != ORT_SIMD
 	       && ctx->region_type != ORT_TARGET_DATA)
@@ -5628,7 +5651,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
   unsigned int nflags;
   tree t;
 
-  if (error_operand_p (decl))
+  if (error_operand_p (decl) || ctx->region_type == ORT_NONE)
     return;
 
   /* Never elide decls whose type has TREE_ADDRESSABLE set.  This means
@@ -5638,7 +5661,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
     flags |= GOVD_SEEN;
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (n != NULL && n->value != GOVD_ALIGNED)
+  if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
     {
       /* We shouldn't be re-adding the decl with the same data
 	 sharing class.  */
@@ -5668,6 +5691,9 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
 	    nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
 	  else if (flags & GOVD_PRIVATE)
 	    nflags = GOVD_PRIVATE;
+	  else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+		   && (flags & GOVD_FIRSTPRIVATE))
+	    nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
 	  else
 	    nflags = GOVD_FIRSTPRIVATE;
 	  nflags |= flags & GOVD_SEEN;
@@ -5712,7 +5738,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
       if ((flags & GOVD_SHARED) == 0)
 	{
 	  t = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl)));
-	  if (TREE_CODE (t) != INTEGER_CST)
+	  if (DECL_P (t))
 	    omp_notice_variable (ctx, t, true);
 	}
     }
@@ -5736,7 +5762,7 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
   struct gimplify_omp_ctx *octx;
 
   for (octx = ctx; octx; octx = octx->outer_context)
-    if (octx->region_type == ORT_TARGET)
+    if ((octx->region_type & ORT_TARGET) != 0)
       {
 	n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
 	if (n == NULL)
@@ -5871,6 +5897,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   if (error_operand_p (decl))
     return false;
 
+  if (ctx->region_type == ORT_NONE)
+    return lang_hooks.decls.omp_disregard_value_expr (decl, false);
+
   /* Threadprivate variables are predetermined.  */
   if (is_global_var (decl))
     {
@@ -5887,19 +5916,66 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
     }
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (ctx->region_type == ORT_TARGET)
+  if ((ctx->region_type & ORT_TARGET) != 0)
     {
       ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
       if (n == NULL)
 	{
-	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+	  unsigned nflags = flags;
+	  if (ctx->target_map_pointers_as_0len_arrays
+	      || ctx->target_map_scalars_firstprivate)
+	    {
+	      bool is_declare_target = false;
+	      bool is_scalar = false;
+	      if (is_global_var (decl)
+		  && varpool_node::get_create (decl)->offloadable)
+		{
+		  struct gimplify_omp_ctx *octx;
+		  for (octx = ctx->outer_context;
+		       octx; octx = octx->outer_context)
+		    {
+		      n = splay_tree_lookup (octx->variables,
+					     (splay_tree_key)decl);
+		      if (n
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
+			break;
+		    }
+		  is_declare_target = octx == NULL;
+		}
+	      if (!is_declare_target && ctx->target_map_scalars_firstprivate)
+		{
+		  tree type = TREE_TYPE (decl);
+		  if (TREE_CODE (type) == REFERENCE_TYPE)
+		    type = TREE_TYPE (type);
+		  if (TREE_CODE (type) == COMPLEX_TYPE)
+		    type = TREE_TYPE (type);
+		  if (INTEGRAL_TYPE_P (type)
+		      || SCALAR_FLOAT_TYPE_P (type)
+		      || TREE_CODE (type) == POINTER_TYPE)
+		    is_scalar = true;
+		}
+	      if (is_declare_target)
+		;
+	      else if (ctx->target_map_pointers_as_0len_arrays
+		       && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+			   || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+			       && TREE_CODE (TREE_TYPE (TREE_TYPE (decl)))
+				  == POINTER_TYPE)))
+		nflags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
+	      else if (is_scalar)
+		nflags |= GOVD_FIRSTPRIVATE;
+	    }
+	  if (nflags == flags
+	      && !lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
 	    {
 	      error ("%qD referenced in target region does not have "
 		     "a mappable type", decl);
-	      omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
+	      nflags |= GOVD_MAP | GOVD_EXPLICIT;
 	    }
-	  else
-	    omp_add_variable (ctx, decl, GOVD_MAP | flags);
+	  else if (nflags == flags)
+	    nflags |= GOVD_MAP;
+	  omp_add_variable (ctx, decl, nflags);
 	}
       else
 	{
@@ -6046,19 +6122,38 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
     {
       ctx = ctx->outer_context;
       if (ctx == NULL)
-	return !(is_global_var (decl)
-		 /* References might be private, but might be shared too,
-		    when checking for copyprivate, assume they might be
-		    private, otherwise assume they might be shared.  */
-		 || (!copyprivate
-		     && lang_hooks.decls.omp_privatize_by_reference (decl)));
+	{
+	  if (is_global_var (decl))
+	    return false;
+
+	  /* References might be private, but might be shared too,
+	     when checking for copyprivate, assume they might be
+	     private, otherwise assume they might be shared.  */
+	  if (copyprivate)
+	    return true;
+
+	  if (lang_hooks.decls.omp_privatize_by_reference (decl))
+	    return false;
+
+	  /* Treat C++ privatized non-static data members outside
+	     of the privatization the same.  */
+	  if (omp_member_access_dummy_var (decl))
+	    return false;
+
+	  return true;
+	}
 
       if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
 	continue;
 
       n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
       if (n != NULL)
-	return (n->value & GOVD_SHARED) == 0;
+	{
+	  if ((n->value & GOVD_LOCAL) != 0
+	      && omp_member_access_dummy_var (decl))
+	    return false;
+	  return (n->value & GOVD_SHARED) == 0;
+	}
     }
   while (ctx->region_type == ORT_WORKSHARE
 	 || ctx->region_type == ORT_SIMD);
@@ -6095,18 +6190,55 @@ omp_no_lastprivate (struct gimplify_omp_ctx *ctx)
   while (1);
 }
 
+/* Callback for walk_tree to find a DECL_EXPR for the given DECL.  */
+
+static tree
+find_decl_expr (tree *tp, int *walk_subtrees, void *data)
+{
+  tree t = *tp;
+
+  /* If this node has been visited, unmark it and keep looking.  */
+  if (TREE_CODE (t) == DECL_EXPR && DECL_EXPR_DECL (t) == (tree) data)
+    return t;
+
+  if (IS_TYPE_OR_DECL_P (t))
+    *walk_subtrees = 0;
+  return NULL_TREE;
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
 static void
 gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
-			   enum omp_region_type region_type)
+			   enum omp_region_type region_type,
+			   enum tree_code code)
 {
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
+  hash_map<tree, tree> *struct_map_to_clause = NULL;
+  tree *orig_list_p = list_p;
 
   ctx = new_omp_context (region_type);
   outer_ctx = ctx->outer_context;
+  if (code == OMP_TARGET && !lang_GNU_Fortran ())
+    {
+      ctx->target_map_pointers_as_0len_arrays = true;
+      /* FIXME: For Fortran we want to set this too, when
+	 the Fortran FE is updated to OpenMP 4.5.  */
+      ctx->target_map_scalars_firstprivate = true;
+    }
+  if (!lang_GNU_Fortran ())
+    switch (code)
+      {
+      case OMP_TARGET:
+      case OMP_TARGET_DATA:
+      case OMP_TARGET_ENTER_DATA:
+      case OMP_TARGET_EXIT_DATA:
+	ctx->target_firstprivatize_array_bases = true;
+      default:
+	break;
+      }
 
   while ((c = *list_p) != NULL)
     {
@@ -6152,6 +6284,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 					 (splay_tree_key) decl) == NULL)
 	    omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
 	  else if (outer_ctx
+		   && (outer_ctx->region_type & ORT_TASK) != 0
+		   && outer_ctx->combined_loop
+		   && splay_tree_lookup (outer_ctx->variables,
+					 (splay_tree_key) decl) == NULL)
+	    omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+	  else if (outer_ctx
 		   && outer_ctx->region_type == ORT_WORKSHARE
 		   && outer_ctx->combined_loop
 		   && splay_tree_lookup (outer_ctx->variables,
@@ -6171,7 +6309,28 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_REDUCTION:
 	  flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
 	  check_non_private = "reduction";
-	  goto do_add;
+	  decl = OMP_CLAUSE_DECL (c);
+	  if (TREE_CODE (decl) == MEM_REF)
+	    {
+	      tree type = TREE_TYPE (decl);
+	      if (gimplify_expr (&TYPE_MAX_VALUE (TYPE_DOMAIN (type)), pre_p,
+				 NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+		{
+		  remove = true;
+		  break;
+		}
+	      tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
+	      if (DECL_P (v))
+		{
+		  omp_firstprivatize_variable (ctx, v);
+		  omp_notice_variable (ctx, v, true);
+		}
+	      decl = TREE_OPERAND (decl, 0);
+	      if (TREE_CODE (decl) == ADDR_EXPR
+		  || TREE_CODE (decl) == INDIRECT_REF)
+		decl = TREE_OPERAND (decl, 0);
+	    }
+	  goto do_add_decl;
 	case OMP_CLAUSE_LINEAR:
 	  if (gimplify_expr (&OMP_CLAUSE_LINEAR_STEP (c), pre_p, NULL,
 			     is_gimple_val, fb_rvalue) == GS_ERROR)
@@ -6181,6 +6340,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  else
 	    {
+	      if (code == OMP_SIMD
+		  && !OMP_CLAUSE_LINEAR_NO_COPYIN (c))
+		{
+		  struct gimplify_omp_ctx *octx = outer_ctx;
+		  if (octx
+		      && octx->region_type == ORT_WORKSHARE
+		      && octx->combined_loop
+		      && !octx->distribute)
+		    {
+		      if (octx->outer_context
+			  && (octx->outer_context->region_type
+			      == ORT_COMBINED_PARALLEL))
+			octx = octx->outer_context->outer_context;
+		      else
+			octx = octx->outer_context;
+		    }
+		  if (octx
+		      && octx->region_type == ORT_WORKSHARE
+		      && octx->combined_loop
+		      && octx->distribute
+		      && !lang_GNU_Fortran ())
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"%<linear%> clause for variable other than "
+				"loop iterator specified on construct "
+				"combined with %<distribute%>");
+		      remove = true;
+		      break;
+		    }
+		}
 	      /* For combined #pragma omp parallel for simd, need to put
 		 lastprivate and perhaps firstprivate too on the
 		 parallel.  Similarly for #pragma omp for simd.  */
@@ -6199,6 +6388,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      decl = NULL_TREE;
 		      break;
 		    }
+		  flags = GOVD_SEEN;
+		  if (!OMP_CLAUSE_LINEAR_NO_COPYIN (c))
+		    flags |= GOVD_FIRSTPRIVATE;
+		  if (!OMP_CLAUSE_LINEAR_NO_COPYOUT (c))
+		    flags |= GOVD_LASTPRIVATE;
 		  if (octx
 		      && octx->region_type == ORT_WORKSHARE
 		      && octx->combined_loop)
@@ -6212,19 +6406,28 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      else if (omp_check_private (octx, decl, false))
 			break;
 		    }
+		  else if (octx
+			   && (octx->region_type & ORT_TASK) != 0
+			   && octx->combined_loop)
+		    ;
+		  else if (octx
+			   && octx->region_type == ORT_COMBINED_PARALLEL
+			   && ctx->region_type == ORT_WORKSHARE
+			   && octx == outer_ctx)
+		    flags = GOVD_SEEN | GOVD_SHARED;
+		  else if (octx
+			   && octx->region_type == ORT_COMBINED_TARGET)
+		    flags &= ~GOVD_LASTPRIVATE;
 		  else
 		    break;
-		  if (splay_tree_lookup (octx->variables,
-					 (splay_tree_key) decl) != NULL)
+		  splay_tree_node on
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  if (on && (on->value & GOVD_DATA_SHARE_CLASS) != 0)
 		    {
 		      octx = NULL;
 		      break;
 		    }
-		  flags = GOVD_SEEN;
-		  if (!OMP_CLAUSE_LINEAR_NO_COPYIN (c))
-		    flags |= GOVD_FIRSTPRIVATE;
-		  if (!OMP_CLAUSE_LINEAR_NO_COPYOUT (c))
-		    flags |= GOVD_LASTPRIVATE;
 		  omp_add_variable (octx, decl, flags);
 		  if (octx->outer_context == NULL)
 		    break;
@@ -6249,10 +6452,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_MAP:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (error_operand_p (decl))
+	    remove = true;
+	  switch (code)
 	    {
-	      remove = true;
+	    case OMP_TARGET:
+	      break;
+	    case OMP_TARGET_DATA:
+	    case OMP_TARGET_ENTER_DATA:
+	    case OMP_TARGET_EXIT_DATA:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		/* For target {,enter ,exit }data only the array slice is
+		   mapped, but not the pointer to it.  */
+		remove = true;
+	      break;
+	    default:
 	      break;
 	    }
+	  if (remove)
+	    break;
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -6262,21 +6479,254 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      remove = true;
 	      break;
 	    }
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+	    {
+	      OMP_CLAUSE_SIZE (c)
+		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL);
+	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	    }
 	  if (!DECL_P (decl))
 	    {
-	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
-				 NULL, is_gimple_lvalue, fb_lvalue)
+	      tree d = decl, *pd;
+	      if (TREE_CODE (d) == ARRAY_REF)
+		{
+		  while (TREE_CODE (d) == ARRAY_REF)
+		    d = TREE_OPERAND (d, 0);
+		  if (TREE_CODE (d) == COMPONENT_REF
+		      && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+		    decl = d;
+		}
+	      pd = &OMP_CLAUSE_DECL (c);
+	      if (d == decl
+		  && TREE_CODE (decl) == INDIRECT_REF
+		  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+		  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+		      == REFERENCE_TYPE))
+		{
+		  pd = &TREE_OPERAND (decl, 0);
+		  decl = TREE_OPERAND (decl, 0);
+		}
+	      if (TREE_CODE (decl) == COMPONENT_REF)
+		{
+		  while (TREE_CODE (decl) == COMPONENT_REF)
+		    decl = TREE_OPERAND (decl, 0);
+		}
+	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
 		  == GS_ERROR)
 		{
 		  remove = true;
 		  break;
 		}
+	      if (DECL_P (decl))
+		{
+		  if (error_operand_p (decl))
+		    {
+		      remove = true;
+		      break;
+		    }
+
+		  if (TYPE_SIZE_UNIT (TREE_TYPE (decl)) == NULL
+		      || (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (decl)))
+			  != INTEGER_CST))
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"mapping field %qE of variable length "
+				"structure", OMP_CLAUSE_DECL (c));
+		      remove = true;
+		      break;
+		    }
+
+		  tree offset;
+		  HOST_WIDE_INT bitsize, bitpos;
+		  machine_mode mode;
+		  int unsignedp, volatilep = 0;
+		  tree base = OMP_CLAUSE_DECL (c);
+		  while (TREE_CODE (base) == ARRAY_REF)
+		    base = TREE_OPERAND (base, 0);
+		  if (TREE_CODE (base) == INDIRECT_REF)
+		    base = TREE_OPERAND (base, 0);
+		  base = get_inner_reference (base, &bitsize, &bitpos, &offset,
+					      &mode, &unsignedp,
+					      &volatilep, false);
+		  gcc_assert (base == decl
+			      && (offset == NULL_TREE
+				  || TREE_CODE (offset) == INTEGER_CST));
+
+		  splay_tree_node n
+		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)
+			      == GOMP_MAP_FIRSTPRIVATE_POINTER);
+		  if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
+						    : GOVD_MAP)) == 0)
+		    {
+		      if (ptr)
+			{
+			  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						      OMP_CLAUSE_PRIVATE);
+			  OMP_CLAUSE_DECL (c2) = decl;
+			  OMP_CLAUSE_CHAIN (c2) = *orig_list_p;
+			  *orig_list_p = c2;
+			  if (struct_map_to_clause == NULL)
+			    struct_map_to_clause = new hash_map<tree, tree>;
+			  tree *osc;
+			  if (n == NULL || (n->value & GOVD_MAP) == 0)
+			    osc = NULL;
+			  else
+			    osc = struct_map_to_clause->get (decl);
+			  if (osc == NULL)
+			    struct_map_to_clause->put (decl,
+						       tree_cons (NULL_TREE,
+								  c,
+								  NULL_TREE));
+			  else
+			    *osc = tree_cons (*osc, c, NULL_TREE);
+			  flags = GOVD_PRIVATE | GOVD_EXPLICIT;
+			  goto do_add_decl;
+			}
+		      *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
+		      OMP_CLAUSE_DECL (*list_p) = decl;
+		      OMP_CLAUSE_SIZE (*list_p) = size_int (1);
+		      OMP_CLAUSE_CHAIN (*list_p) = c;
+		      if (struct_map_to_clause == NULL)
+			struct_map_to_clause = new hash_map<tree, tree>;
+		      struct_map_to_clause->put (decl, *list_p);
+		      list_p = &OMP_CLAUSE_CHAIN (*list_p);
+		      flags = GOVD_MAP | GOVD_EXPLICIT;
+		      if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+			flags |= GOVD_SEEN;
+		      goto do_add_decl;
+		    }
+		  else
+		    {
+		      tree *osc = struct_map_to_clause->get (decl);
+		      tree *sc = NULL, *pt = NULL;
+		      if (!ptr && TREE_CODE (*osc) == TREE_LIST)
+			osc = &TREE_PURPOSE (*osc);
+		      if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+			n->value |= GOVD_SEEN;
+		      offset_int o1, o2;
+		      if (offset)
+			o1 = wi::to_offset (offset);
+		      else
+			o1 = 0;
+		      if (bitpos)
+			o1 = o1 + bitpos / BITS_PER_UNIT;
+		      if (ptr)
+			pt = osc;
+		      else
+			sc = &OMP_CLAUSE_CHAIN (*osc);
+		      for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt)))
+				 : *sc != c;
+			   ptr ? (pt = &TREE_CHAIN (*pt))
+			       : (sc = &OMP_CLAUSE_CHAIN (*sc)))
+			if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
+			    && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+				!= INDIRECT_REF)
+			    && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF)
+			  break;
+			else
+			  {
+			    tree offset2;
+			    HOST_WIDE_INT bitsize2, bitpos2;
+			    base = OMP_CLAUSE_DECL (*sc);
+			    if (TREE_CODE (base) == ARRAY_REF)
+			      {
+				while (TREE_CODE (base) == ARRAY_REF)
+				  base = TREE_OPERAND (base, 0);
+				if (TREE_CODE (base) != COMPONENT_REF
+				    || (TREE_CODE (TREE_TYPE (base))
+					!= ARRAY_TYPE))
+				  break;
+			      }
+			    else if (TREE_CODE (base) == INDIRECT_REF
+				     && (TREE_CODE (TREE_OPERAND (base, 0))
+					 == COMPONENT_REF)
+				     && (TREE_CODE (TREE_TYPE
+						     (TREE_OPERAND (base, 0)))
+					 == REFERENCE_TYPE))
+			      base = TREE_OPERAND (base, 0);
+			    base = get_inner_reference (base, &bitsize2,
+							&bitpos2, &offset2,
+							&mode, &unsignedp,
+							&volatilep, false);
+			    if (base != decl)
+			      break;
+			    gcc_assert (offset == NULL_TREE
+					|| TREE_CODE (offset) == INTEGER_CST);
+			    tree d1 = OMP_CLAUSE_DECL (*sc);
+			    tree d2 = OMP_CLAUSE_DECL (c);
+			    while (TREE_CODE (d1) == ARRAY_REF)
+			      d1 = TREE_OPERAND (d1, 0);
+			    while (TREE_CODE (d2) == ARRAY_REF)
+			      d2 = TREE_OPERAND (d2, 0);
+			    if (TREE_CODE (d1) == INDIRECT_REF)
+			      d1 = TREE_OPERAND (d1, 0);
+			    if (TREE_CODE (d2) == INDIRECT_REF)
+			      d2 = TREE_OPERAND (d2, 0);
+			    while (TREE_CODE (d1) == COMPONENT_REF)
+			      if (TREE_CODE (d2) == COMPONENT_REF
+				  && TREE_OPERAND (d1, 1)
+				     == TREE_OPERAND (d2, 1))
+				{
+				  d1 = TREE_OPERAND (d1, 0);
+				  d2 = TREE_OPERAND (d2, 0);
+				}
+			      else
+				break;
+			    if (d1 == d2)
+			      {
+				error_at (OMP_CLAUSE_LOCATION (c),
+					  "%qE appears more than once in map "
+					  "clauses", OMP_CLAUSE_DECL (c));
+				remove = true;
+				break;
+			      }
+			    if (offset2)
+			      o2 = wi::to_offset (offset2);
+			    else
+			      o2 = 0;
+			    if (bitpos2)
+			      o2 = o2 + bitpos2 / BITS_PER_UNIT;
+			    if (wi::ltu_p (o1, o2)
+				|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
+			      break;
+			  }
+		      if (ptr)
+			{
+			  if (!remove)
+			    *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
+			  break;
+			}
+		      if (!remove)
+			OMP_CLAUSE_SIZE (*osc)
+			  = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+					size_one_node);
+		      if (!remove && *sc != c)
+			{
+			  *list_p = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = *sc;
+			  *sc = c;
+			  continue;
+			}
+		    }
+		}
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
+	  if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK
+	      || OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE)
+	    {
+	      /* Nothing to do.  OMP_CLAUSE_DECL will be lowered in
+		 omp-low.c.  */
+	      break;
+	    }
 	  if (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPOUND_EXPR)
 	    {
 	      gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (c), 0), pre_p,
@@ -6328,19 +6778,46 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  goto do_notice;
 
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
+	  goto do_add;
+	case OMP_CLAUSE_IS_DEVICE_PTR:
+	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
+	  goto do_add;
+
 	do_add:
 	  decl = OMP_CLAUSE_DECL (c);
+	do_add_decl:
 	  if (error_operand_p (decl))
 	    {
 	      remove = true;
 	      break;
 	    }
+	  if (DECL_NAME (decl) == NULL_TREE && (flags & GOVD_SHARED) == 0)
+	    {
+	      tree t = omp_member_access_dummy_var (decl);
+	      if (t)
+		{
+		  tree v = DECL_VALUE_EXPR (decl);
+		  DECL_NAME (decl) = DECL_NAME (TREE_OPERAND (v, 1));
+		  if (outer_ctx)
+		    omp_notice_variable (outer_ctx, t, true);
+		}
+	    }
 	  omp_add_variable (ctx, decl, flags);
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
 	      && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
 	    {
 	      omp_add_variable (ctx, OMP_CLAUSE_REDUCTION_PLACEHOLDER (c),
 				GOVD_LOCAL | GOVD_SEEN);
+	      if (OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c)
+		  && walk_tree (&OMP_CLAUSE_REDUCTION_INIT (c),
+				find_decl_expr,
+				OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c),
+				NULL) == NULL_TREE)
+		omp_add_variable (ctx,
+				  OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c),
+				  GOVD_LOCAL | GOVD_SEEN);
 	      gimplify_omp_ctxp = ctx;
 	      push_gimplify_context ();
 
@@ -6444,6 +6921,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    omp_notice_variable (outer_ctx, decl, true);
 	  if (check_non_private
 	      && region_type == ORT_WORKSHARE
+	      && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION
+		  || decl == OMP_CLAUSE_DECL (c)
+		  || (TREE_CODE (OMP_CLAUSE_DECL (c)) == MEM_REF
+		      && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
+			  == ADDR_EXPR)))
 	      && omp_check_private (ctx, decl, false))
 	    {
 	      error ("%s variable %qE is private in outer context",
@@ -6452,8 +6934,33 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  break;
 
-	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_IF:
+	  if (OMP_CLAUSE_IF_MODIFIER (c) != ERROR_MARK
+	      && OMP_CLAUSE_IF_MODIFIER (c) != code)
+	    {
+	      const char *p[2];
+	      for (int i = 0; i < 2; i++)
+		switch (i ? OMP_CLAUSE_IF_MODIFIER (c) : code)
+		  {
+		  case OMP_PARALLEL: p[i] = "parallel"; break;
+		  case OMP_TASK: p[i] = "task"; break;
+		  case OMP_TASKLOOP: p[i] = "taskloop"; break;
+		  case OMP_TARGET_DATA: p[i] = "target data"; break;
+		  case OMP_TARGET: p[i] = "target"; break;
+		  case OMP_TARGET_UPDATE: p[i] = "target update"; break;
+		  case OMP_TARGET_ENTER_DATA:
+		    p[i] = "target enter data"; break;
+		  case OMP_TARGET_EXIT_DATA: p[i] = "target exit data"; break;
+		  default: gcc_unreachable ();
+		  }
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"expected %qs %<if%> clause modifier rather than %qs",
+			p[0], p[1]);
+	      remove = true;
+	    }
+	  /* Fall through.  */
+
+	case OMP_CLAUSE_FINAL:
 	  OMP_CLAUSE_OPERAND (c, 0)
 	    = gimple_boolify (OMP_CLAUSE_OPERAND (c, 0));
 	  /* Fall through.  */
@@ -6464,21 +6971,29 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_THREAD_LIMIT:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEVICE:
+	case OMP_CLAUSE_PRIORITY:
+	case OMP_CLAUSE_GRAINSIZE:
+	case OMP_CLAUSE_NUM_TASKS:
+	case OMP_CLAUSE_HINT:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
-	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
 	  if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
 			     is_gimple_val, fb_rvalue) == GS_ERROR)
 	    remove = true;
-	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG
-	      && gimplify_expr (&OMP_CLAUSE_OPERAND (c, 1), pre_p, NULL,
-				is_gimple_val, fb_rvalue) == GS_ERROR)
+	  break;
+
+	case OMP_CLAUSE_GANG:
+	  if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
+			     is_gimple_val, fb_rvalue) == GS_ERROR)
+	    remove = true;
+	  if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 1), pre_p, NULL,
+			     is_gimple_val, fb_rvalue) == GS_ERROR)
 	    remove = true;
 	  break;
 
@@ -6497,6 +7012,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
+	case OMP_CLAUSE_SIMDLEN:
+	case OMP_CLAUSE_NOGROUP:
+	case OMP_CLAUSE_THREADS:
+	case OMP_CLAUSE_SIMD:
+	  break;
+
+	case OMP_CLAUSE_DEFAULTMAP:
+	  ctx->target_map_scalars_firstprivate = false;
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -6532,6 +7055,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
     }
 
   gimplify_omp_ctxp = ctx;
+  if (struct_map_to_clause)
+    delete struct_map_to_clause;
 }
 
 struct gimplify_adjust_omp_clauses_data
@@ -6612,6 +7137,30 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
     OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+  else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
+    {
+      tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_DECL (nc) = decl;
+      if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+	  && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE)
+	OMP_CLAUSE_DECL (clause)
+	  = build_simple_mem_ref_loc (input_location, decl);
+      OMP_CLAUSE_DECL (clause)
+	= build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause),
+		  build_int_cst (build_pointer_type (char_type_node), 0));
+      OMP_CLAUSE_SIZE (clause) = size_zero_node;
+      OMP_CLAUSE_SIZE (nc) = size_zero_node;
+      OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
+      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1;
+      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      OMP_CLAUSE_CHAIN (nc) = *list_p;
+      OMP_CLAUSE_CHAIN (clause) = nc;
+      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = ctx->outer_context;
+      gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
+		     pre_p, NULL, is_gimple_val, fb_rvalue);
+      gimplify_omp_ctxp = ctx;
+    }
   else if (code == OMP_CLAUSE_MAP)
     {
       OMP_CLAUSE_SET_MAP_KIND (clause,
@@ -6638,7 +7187,10 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 				      OMP_CLAUSE_MAP);
 	  OMP_CLAUSE_DECL (nc) = decl;
 	  OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	  OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+	  if (gimplify_omp_ctxp->target_firstprivatize_array_bases)
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  else
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
 	  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
 	  OMP_CLAUSE_CHAIN (clause) = nc;
 	}
@@ -6666,7 +7218,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 }
 
 static void
-gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
+gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p,
+			     enum tree_code code)
 {
   struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
   tree c, decl;
@@ -6761,13 +7314,56 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 	case OMP_CLAUSE_MAP:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (!DECL_P (decl))
-	    break;
+	    {
+	      if ((ctx->region_type & ORT_TARGET) != 0
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		{
+		  if (TREE_CODE (decl) == INDIRECT_REF
+		      && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+		      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			  == REFERENCE_TYPE))
+		    decl = TREE_OPERAND (decl, 0);
+		  if (TREE_CODE (decl) == COMPONENT_REF)
+		    {
+		      while (TREE_CODE (decl) == COMPONENT_REF)
+			decl = TREE_OPERAND (decl, 0);
+		      if (DECL_P (decl))
+			{
+			  n = splay_tree_lookup (ctx->variables,
+						 (splay_tree_key) decl);
+			  if (!(n->value & GOVD_SEEN))
+			    remove = true;
+			}
+		    }
+		}
+	      break;
+	    }
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
-	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
+	  if ((ctx->region_type & ORT_TARGET) != 0
+	      && !(n->value & GOVD_SEEN)
+	      && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
+	    {
+	      remove = true;
+	      /* For struct element mapping, if struct is never referenced
+		 in target block and none of the mapping has always modifier,
+		 remove all the struct element mappings, which immediately
+		 follow the GOMP_MAP_STRUCT map clause.  */
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+		{
+		  HOST_WIDE_INT cnt = tree_to_shwi (OMP_CLAUSE_SIZE (c));
+		  while (cnt--)
+		    OMP_CLAUSE_CHAIN (c)
+		      = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
+		}
+	    }
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+		   && code == OMP_TARGET_EXIT_DATA)
 	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
-		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
 	      /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
 		 for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -6787,17 +7383,33 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 		  omp_notice_variable (ctx->outer_context,
 				       OMP_CLAUSE_SIZE (c), true);
 		}
-	      tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-					  OMP_CLAUSE_MAP);
-	      OMP_CLAUSE_DECL (nc) = decl;
-	      OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
-	      OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
-	      OMP_CLAUSE_CHAIN (c) = nc;
-	      c = nc;
+	      if (((ctx->region_type & ORT_TARGET) != 0
+		   || !ctx->target_firstprivatize_array_bases)
+		  && ((n->value & GOVD_SEEN) == 0
+		      || (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0))
+		{
+		  tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					      OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_DECL (nc) = decl;
+		  OMP_CLAUSE_SIZE (nc) = size_zero_node;
+		  if (ctx->target_firstprivatize_array_bases)
+		    OMP_CLAUSE_SET_MAP_KIND (nc,
+					     GOMP_MAP_FIRSTPRIVATE_POINTER);
+		  else
+		    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+		  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+		  OMP_CLAUSE_CHAIN (c) = nc;
+		  c = nc;
+		}
+	    }
+	  else
+	    {
+	      if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+		OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
+	      if ((n->value & GOVD_SEEN)
+		  && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
+		OMP_CLAUSE_MAP_PRIVATE (c) = 1;
 	    }
-	  else if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
-	    OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
 	  break;
 
 	case OMP_CLAUSE_TO:
@@ -6846,7 +7458,18 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
+	case OMP_CLAUSE_SIMDLEN:
 	case OMP_CLAUSE_DEPEND:
+	case OMP_CLAUSE_PRIORITY:
+	case OMP_CLAUSE_GRAINSIZE:
+	case OMP_CLAUSE_NUM_TASKS:
+	case OMP_CLAUSE_NOGROUP:
+	case OMP_CLAUSE_THREADS:
+	case OMP_CLAUSE_SIMD:
+	case OMP_CLAUSE_HINT:
+	case OMP_CLAUSE_DEFAULTMAP:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
+	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
@@ -6890,8 +7513,9 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
 {
   tree expr = *expr_p;
 
-  gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE);
-  gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr));
+  gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
+			     OACC_CACHE);
+  gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
 
   /* TODO: Do something sensible with this information.  */
 
@@ -6913,7 +7537,7 @@ gimplify_omp_parallel (tree *expr_p, gimple_seq *pre_p)
   gimplify_scan_omp_clauses (&OMP_PARALLEL_CLAUSES (expr), pre_p,
 			     OMP_PARALLEL_COMBINED (expr)
 			     ? ORT_COMBINED_PARALLEL
-			     : ORT_PARALLEL);
+			     : ORT_PARALLEL, OMP_PARALLEL);
 
   push_gimplify_context ();
 
@@ -6923,7 +7547,8 @@ gimplify_omp_parallel (tree *expr_p, gimple_seq *pre_p)
   else
     pop_gimplify_context (NULL);
 
-  gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr),
+			       OMP_PARALLEL);
 
   g = gimple_build_omp_parallel (body,
 				 OMP_PARALLEL_CLAUSES (expr),
@@ -6949,7 +7574,7 @@ gimplify_omp_task (tree *expr_p, gimple_seq *pre_p)
   gimplify_scan_omp_clauses (&OMP_TASK_CLAUSES (expr), pre_p,
 			     find_omp_clause (OMP_TASK_CLAUSES (expr),
 					      OMP_CLAUSE_UNTIED)
-			     ? ORT_UNTIED_TASK : ORT_TASK);
+			     ? ORT_UNTIED_TASK : ORT_TASK, OMP_TASK);
 
   push_gimplify_context ();
 
@@ -6959,7 +7584,7 @@ gimplify_omp_task (tree *expr_p, gimple_seq *pre_p)
   else
     pop_gimplify_context (NULL);
 
-  gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr), OMP_TASK);
 
   g = gimple_build_omp_task (body,
 			     OMP_TASK_CLAUSES (expr),
@@ -7007,8 +7632,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
   gomp_for *gfor;
   gimple_seq for_body, for_pre_body;
   int i;
-  bool simd;
   bitmap has_decl_expr = NULL;
+  enum omp_region_type ort = ORT_WORKSHARE;
 
   orig_for_stmt = for_stmt = *expr_p;
 
@@ -7018,11 +7643,16 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     case CILK_FOR:
     case OMP_DISTRIBUTE:
     case OACC_LOOP:
-      simd = false;
+      break;
+    case OMP_TASKLOOP:
+      if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
+	ort = ORT_UNTIED_TASK;
+      else
+	ort = ORT_TASK;
       break;
     case OMP_SIMD:
     case CILK_SIMD:
-      simd = true;
+      ort = ORT_SIMD;
       break;
     default:
       gcc_unreachable ();
@@ -7030,7 +7660,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 
   /* Set OMP_CLAUSE_LINEAR_NO_COPYIN flag on explicit linear
      clause for the IV.  */
-  if (simd && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) == 1)
+  if (ort == ORT_SIMD && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) == 1)
     {
       t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), 0);
       gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
@@ -7057,14 +7687,16 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	}
     }
 
-  gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
-			     simd ? ORT_SIMD : ORT_WORKSHARE);
+  if (TREE_CODE (for_stmt) != OMP_TASKLOOP)
+    gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort,
+			       TREE_CODE (for_stmt));
+
   if (TREE_CODE (for_stmt) == OMP_DISTRIBUTE)
     gimplify_omp_ctxp->distribute = true;
 
   /* Handle OMP_FOR_INIT.  */
   for_pre_body = NULL;
-  if (simd && OMP_FOR_PRE_BODY (for_stmt))
+  if (ort == ORT_SIMD && OMP_FOR_PRE_BODY (for_stmt))
     {
       has_decl_expr = BITMAP_ALLOC (NULL);
       if (TREE_CODE (OMP_FOR_PRE_BODY (for_stmt)) == DECL_EXPR
@@ -7087,20 +7719,109 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	    }
 	}
     }
-  gimplify_and_add (OMP_FOR_PRE_BODY (for_stmt), &for_pre_body);
+  if (OMP_FOR_PRE_BODY (for_stmt))
+    {
+      if (TREE_CODE (for_stmt) != OMP_TASKLOOP || gimplify_omp_ctxp)
+	gimplify_and_add (OMP_FOR_PRE_BODY (for_stmt), &for_pre_body);
+      else
+	{
+	  struct gimplify_omp_ctx ctx;
+	  memset (&ctx, 0, sizeof (ctx));
+	  ctx.region_type = ORT_NONE;
+	  gimplify_omp_ctxp = &ctx;
+	  gimplify_and_add (OMP_FOR_PRE_BODY (for_stmt), &for_pre_body);
+	  gimplify_omp_ctxp = NULL;
+	}
+    }
   OMP_FOR_PRE_BODY (for_stmt) = NULL_TREE;
 
   if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
+    for_stmt = inner_for_stmt;
+
+  /* For taskloop, need to gimplify the start, end and step before the
+     taskloop, outside of the taskloop omp context.  */
+  if (TREE_CODE (orig_for_stmt) == OMP_TASKLOOP)
     {
-      for_stmt = inner_for_stmt;
-      gimplify_omp_ctxp->combined_loop = true;
+      for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
+	{
+	  t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
+	  if (!is_gimple_constant (TREE_OPERAND (t, 1)))
+	    {
+	      TREE_OPERAND (t, 1)
+		= get_initialized_tmp_var (TREE_OPERAND (t, 1),
+					   pre_p, NULL);
+	      tree c = build_omp_clause (input_location,
+					 OMP_CLAUSE_FIRSTPRIVATE);
+	      OMP_CLAUSE_DECL (c) = TREE_OPERAND (t, 1);
+	      OMP_CLAUSE_CHAIN (c) = OMP_FOR_CLAUSES (orig_for_stmt);
+	      OMP_FOR_CLAUSES (orig_for_stmt) = c;
+	    }
+
+	  /* Handle OMP_FOR_COND.  */
+	  t = TREE_VEC_ELT (OMP_FOR_COND (for_stmt), i);
+	  if (!is_gimple_constant (TREE_OPERAND (t, 1)))
+	    {
+	      TREE_OPERAND (t, 1)
+		= get_initialized_tmp_var (TREE_OPERAND (t, 1),
+					   gimple_seq_empty_p (for_pre_body)
+					   ? pre_p : &for_pre_body, NULL);
+	      tree c = build_omp_clause (input_location,
+					 OMP_CLAUSE_FIRSTPRIVATE);
+	      OMP_CLAUSE_DECL (c) = TREE_OPERAND (t, 1);
+	      OMP_CLAUSE_CHAIN (c) = OMP_FOR_CLAUSES (orig_for_stmt);
+	      OMP_FOR_CLAUSES (orig_for_stmt) = c;
+	    }
+
+	  /* Handle OMP_FOR_INCR.  */
+	  t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
+	  if (TREE_CODE (t) == MODIFY_EXPR)
+	    {
+	      decl = TREE_OPERAND (t, 0);
+	      t = TREE_OPERAND (t, 1);
+	      tree *tp = &TREE_OPERAND (t, 1);
+	      if (TREE_CODE (t) == PLUS_EXPR && *tp == decl)
+		tp = &TREE_OPERAND (t, 0);
+
+	      if (!is_gimple_constant (*tp))
+		{
+		  gimple_seq *seq = gimple_seq_empty_p (for_pre_body)
+				    ? pre_p : &for_pre_body;
+		  *tp = get_initialized_tmp_var (*tp, seq, NULL);
+		  tree c = build_omp_clause (input_location,
+					     OMP_CLAUSE_FIRSTPRIVATE);
+		  OMP_CLAUSE_DECL (c) = *tp;
+		  OMP_CLAUSE_CHAIN (c) = OMP_FOR_CLAUSES (orig_for_stmt);
+		  OMP_FOR_CLAUSES (orig_for_stmt) = c;
+		}
+	    }
+	}
+
+      gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (orig_for_stmt), pre_p, ort,
+				 OMP_TASKLOOP);
     }
 
+  if (orig_for_stmt != for_stmt)
+    gimplify_omp_ctxp->combined_loop = true;
+
   for_body = NULL;
   gcc_assert (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt))
 	      == TREE_VEC_LENGTH (OMP_FOR_COND (for_stmt)));
   gcc_assert (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt))
 	      == TREE_VEC_LENGTH (OMP_FOR_INCR (for_stmt)));
+
+  tree c = find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_ORDERED);
+  bool is_doacross = false;
+  if (c && OMP_CLAUSE_ORDERED_EXPR (c))
+    {
+      is_doacross = true;
+      gimplify_omp_ctxp->loop_iter_var.create (TREE_VEC_LENGTH
+						 (OMP_FOR_INIT (for_stmt))
+					       * 2);
+    }
+  int collapse = 1;
+  c = find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_COLLAPSE);
+  if (c)
+    collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (c));
   for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
     {
       t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
@@ -7109,16 +7830,25 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       gcc_assert (DECL_P (decl));
       gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
 		  || POINTER_TYPE_P (TREE_TYPE (decl)));
+      if (is_doacross)
+	{
+	  if (TREE_CODE (for_stmt) == OMP_FOR && OMP_FOR_ORIG_DECLS (for_stmt))
+	    gimplify_omp_ctxp->loop_iter_var.quick_push
+	      (TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (for_stmt), i));
+	  else
+	    gimplify_omp_ctxp->loop_iter_var.quick_push (decl);
+	  gimplify_omp_ctxp->loop_iter_var.quick_push (decl);
+	}
 
       /* Make sure the iteration variable is private.  */
       tree c = NULL_TREE;
       tree c2 = NULL_TREE;
       if (orig_for_stmt != for_stmt)
 	/* Do this only on innermost construct for combined ones.  */;
-      else if (simd)
+      else if (ort == ORT_SIMD)
 	{
 	  splay_tree_node n = splay_tree_lookup (gimplify_omp_ctxp->variables,
-						 (splay_tree_key)decl);
+						 (splay_tree_key) decl);
 	  omp_is_private (gimplify_omp_ctxp, decl,
 			  1 + (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt))
 			       != 1));
@@ -7169,6 +7899,11 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 		      else if (omp_check_private (outer, decl, false))
 			outer = NULL;
 		    }
+		  else if (((outer->region_type & ORT_TASK) != 0)
+			   && outer->combined_loop
+			   && !omp_check_private (gimplify_omp_ctxp,
+						  decl, false))
+		    ;
 		  else if (outer->region_type != ORT_COMBINED_PARALLEL)
 		    outer = NULL;
 		  if (outer)
@@ -7213,6 +7948,11 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 		      else if (omp_check_private (outer, decl, false))
 			outer = NULL;
 		    }
+		  else if (((outer->region_type & ORT_TASK) != 0)
+			   && outer->combined_loop
+			   && !omp_check_private (gimplify_omp_ctxp,
+						  decl, false))
+		    ;
 		  else if (outer->region_type != ORT_COMBINED_PARALLEL)
 		    outer = NULL;
 		  if (outer)
@@ -7255,14 +7995,16 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       if (orig_for_stmt != for_stmt)
 	var = decl;
       else if (!is_gimple_reg (decl)
-	       || (simd && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) > 1))
+	       || (ort == ORT_SIMD
+		   && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) > 1))
 	{
 	  var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
 	  TREE_OPERAND (t, 0) = var;
 
 	  gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
 
-	  if (simd && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) == 1)
+	  if (ort == ORT_SIMD
+	      && TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) == 1)
 	    {
 	      c2 = build_omp_clause (input_location, OMP_CLAUSE_LINEAR);
 	      OMP_CLAUSE_LINEAR_NO_COPYIN (c2) = 1;
@@ -7396,8 +8138,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	  OMP_CLAUSE_LINEAR_STEP (c2) = OMP_CLAUSE_LINEAR_STEP (c);
 	}
 
-      if ((var != decl || TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)) > 1)
-	  && orig_for_stmt == for_stmt)
+      if ((var != decl || collapse > 1) && orig_for_stmt == for_stmt)
 	{
 	  for (c = OMP_FOR_CLAUSES (for_stmt); c ; c = OMP_CLAUSE_CHAIN (c))
 	    if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
@@ -7407,16 +8148,22 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 		     && OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c) == NULL))
 		&& OMP_CLAUSE_DECL (c) == decl)
 	      {
-		t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
-		gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
-		gcc_assert (TREE_OPERAND (t, 0) == var);
-		t = TREE_OPERAND (t, 1);
-		gcc_assert (TREE_CODE (t) == PLUS_EXPR
-			    || TREE_CODE (t) == MINUS_EXPR
-			    || TREE_CODE (t) == POINTER_PLUS_EXPR);
-		gcc_assert (TREE_OPERAND (t, 0) == var);
-		t = build2 (TREE_CODE (t), TREE_TYPE (decl), decl,
-			    TREE_OPERAND (t, 1));
+		if (is_doacross && (collapse == 1 || i >= collapse))
+		  t = var;
+		else
+		  {
+		    t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
+		    gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
+		    gcc_assert (TREE_OPERAND (t, 0) == var);
+		    t = TREE_OPERAND (t, 1);
+		    gcc_assert (TREE_CODE (t) == PLUS_EXPR
+				|| TREE_CODE (t) == MINUS_EXPR
+				|| TREE_CODE (t) == POINTER_PLUS_EXPR);
+		    gcc_assert (TREE_OPERAND (t, 0) == var);
+		    t = build2 (TREE_CODE (t), TREE_TYPE (decl),
+				is_doacross ? var : decl,
+				TREE_OPERAND (t, 1));
+		  }
 		gimple_seq *seq;
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE)
 		  seq = &OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c);
@@ -7429,14 +8176,39 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 
   BITMAP_FREE (has_decl_expr);
 
-  gimplify_and_add (OMP_FOR_BODY (orig_for_stmt), &for_body);
+  if (TREE_CODE (orig_for_stmt) == OMP_TASKLOOP)
+    {
+      push_gimplify_context ();
+      if (TREE_CODE (OMP_FOR_BODY (orig_for_stmt)) != BIND_EXPR)
+	{
+	  OMP_FOR_BODY (orig_for_stmt)
+	    = build3 (BIND_EXPR, void_type_node, NULL,
+		      OMP_FOR_BODY (orig_for_stmt), NULL);
+	  TREE_SIDE_EFFECTS (OMP_FOR_BODY (orig_for_stmt)) = 1;
+	}
+    }
+
+  gimple *g = gimplify_and_return_first (OMP_FOR_BODY (orig_for_stmt),
+					 &for_body);
+
+  if (TREE_CODE (orig_for_stmt) == OMP_TASKLOOP)
+    {
+      if (gimple_code (g) == GIMPLE_BIND)
+	pop_gimplify_context (g);
+      else
+	pop_gimplify_context (NULL);
+    }
 
   if (orig_for_stmt != for_stmt)
     for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
       {
 	t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
 	decl = TREE_OPERAND (t, 0);
+	struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+	if (TREE_CODE (orig_for_stmt) == OMP_TASKLOOP)
+	  gimplify_omp_ctxp = ctx->outer_context;
 	var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
+	gimplify_omp_ctxp = ctx;
 	omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
 	TREE_OPERAND (t, 0) = var;
 	t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
@@ -7444,7 +8216,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	TREE_OPERAND (TREE_OPERAND (t, 1), 0) = var;
       }
 
-  gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt),
+			       TREE_CODE (orig_for_stmt));
 
   int kind;
   switch (TREE_CODE (orig_for_stmt))
@@ -7454,6 +8227,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
     case CILK_FOR: kind = GF_OMP_FOR_KIND_CILKFOR; break;
     case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
+    case OMP_TASKLOOP: kind = GF_OMP_FOR_KIND_TASKLOOP; break;
     case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
     default:
       gcc_unreachable ();
@@ -7488,7 +8262,139 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       gimple_omp_for_set_incr (gfor, i, TREE_OPERAND (t, 1));
     }
 
-  gimplify_seq_add_stmt (pre_p, gfor);
+  /* OMP_TASKLOOP is gimplified as two GIMPLE_OMP_FOR taskloop
+     constructs with GIMPLE_OMP_TASK sandwiched in between them.
+     The outer taskloop stands for computing the number of iterations,
+     counts for collapsed loops and holding taskloop specific clauses.
+     The task construct stands for the effect of data sharing on the
+     explicit task it creates and the inner taskloop stands for expansion
+     of the static loop inside of the explicit task construct.  */
+  if (TREE_CODE (orig_for_stmt) == OMP_TASKLOOP)
+    {
+      tree *gfor_clauses_ptr = gimple_omp_for_clauses_ptr (gfor);
+      tree task_clauses = NULL_TREE;
+      tree c = *gfor_clauses_ptr;
+      tree *gtask_clauses_ptr = &task_clauses;
+      tree outer_for_clauses = NULL_TREE;
+      tree *gforo_clauses_ptr = &outer_for_clauses;
+      for (; c; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	  /* These clauses are allowed on task, move them there.  */
+	  case OMP_CLAUSE_SHARED:
+	  case OMP_CLAUSE_FIRSTPRIVATE:
+	  case OMP_CLAUSE_DEFAULT:
+	  case OMP_CLAUSE_IF:
+	  case OMP_CLAUSE_UNTIED:
+	  case OMP_CLAUSE_FINAL:
+	  case OMP_CLAUSE_MERGEABLE:
+	  case OMP_CLAUSE_PRIORITY:
+	    *gtask_clauses_ptr = c;
+	    gtask_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	    break;
+	  case OMP_CLAUSE_PRIVATE:
+	    if (OMP_CLAUSE_PRIVATE_TASKLOOP_IV (c))
+	      {
+		/* We want private on outer for and firstprivate
+		   on task.  */
+		*gtask_clauses_ptr
+		  = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+				      OMP_CLAUSE_FIRSTPRIVATE);
+		OMP_CLAUSE_DECL (*gtask_clauses_ptr) = OMP_CLAUSE_DECL (c);
+		lang_hooks.decls.omp_finish_clause (*gtask_clauses_ptr, NULL);
+		gtask_clauses_ptr = &OMP_CLAUSE_CHAIN (*gtask_clauses_ptr);
+		*gforo_clauses_ptr = c;
+		gforo_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	      }
+	    else
+	      {
+		*gtask_clauses_ptr = c;
+		gtask_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	      }
+	    break;
+	  /* These clauses go into outer taskloop clauses.  */
+	  case OMP_CLAUSE_GRAINSIZE:
+	  case OMP_CLAUSE_NUM_TASKS:
+	  case OMP_CLAUSE_NOGROUP:
+	    *gforo_clauses_ptr = c;
+	    gforo_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	    break;
+	  /* Taskloop clause we duplicate on both taskloops.  */
+	  case OMP_CLAUSE_COLLAPSE:
+	    *gfor_clauses_ptr = c;
+	    gfor_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	    *gforo_clauses_ptr = copy_node (c);
+	    gforo_clauses_ptr = &OMP_CLAUSE_CHAIN (*gforo_clauses_ptr);
+	    break;
+	  /* For lastprivate, keep the clause on inner taskloop, and add
+	     a shared clause on task.  If the same decl is also firstprivate,
+	     add also firstprivate clause on the inner taskloop.  */
+	  case OMP_CLAUSE_LASTPRIVATE:
+	    if (OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV (c))
+	      {
+		/* For taskloop C++ lastprivate IVs, we want:
+		   1) private on outer taskloop
+		   2) firstprivate and shared on task
+		   3) lastprivate on inner taskloop  */
+		*gtask_clauses_ptr
+		  = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+				      OMP_CLAUSE_FIRSTPRIVATE);
+		OMP_CLAUSE_DECL (*gtask_clauses_ptr) = OMP_CLAUSE_DECL (c);
+		lang_hooks.decls.omp_finish_clause (*gtask_clauses_ptr, NULL);
+		gtask_clauses_ptr = &OMP_CLAUSE_CHAIN (*gtask_clauses_ptr);
+		OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c) = 1;
+		*gforo_clauses_ptr = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						       OMP_CLAUSE_PRIVATE);
+		OMP_CLAUSE_DECL (*gforo_clauses_ptr) = OMP_CLAUSE_DECL (c);
+		OMP_CLAUSE_PRIVATE_TASKLOOP_IV (*gforo_clauses_ptr) = 1;
+		TREE_TYPE (*gforo_clauses_ptr) = TREE_TYPE (c);
+		gforo_clauses_ptr = &OMP_CLAUSE_CHAIN (*gforo_clauses_ptr);
+	      }
+	    *gfor_clauses_ptr = c;
+	    gfor_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+	    *gtask_clauses_ptr
+	      = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_SHARED);
+	    OMP_CLAUSE_DECL (*gtask_clauses_ptr) = OMP_CLAUSE_DECL (c);
+	    if (OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
+	      OMP_CLAUSE_SHARED_FIRSTPRIVATE (*gtask_clauses_ptr) = 1;
+	    gtask_clauses_ptr
+	      = &OMP_CLAUSE_CHAIN (*gtask_clauses_ptr);
+	    break;
+	  default:
+	    gcc_unreachable ();
+	  }
+      *gfor_clauses_ptr = NULL_TREE;
+      *gtask_clauses_ptr = NULL_TREE;
+      *gforo_clauses_ptr = NULL_TREE;
+      g = gimple_build_bind (NULL_TREE, gfor, NULL_TREE);
+      g = gimple_build_omp_task (g, task_clauses, NULL_TREE, NULL_TREE,
+				 NULL_TREE, NULL_TREE, NULL_TREE);
+      gimple_omp_task_set_taskloop_p (g, true);
+      g = gimple_build_bind (NULL_TREE, g, NULL_TREE);
+      gomp_for *gforo
+	= gimple_build_omp_for (g, GF_OMP_FOR_KIND_TASKLOOP, outer_for_clauses,
+				gimple_omp_for_collapse (gfor),
+				gimple_omp_for_pre_body (gfor));
+      gimple_omp_for_set_pre_body (gfor, NULL);
+      gimple_omp_for_set_combined_p (gforo, true);
+      gimple_omp_for_set_combined_into_p (gfor, true);
+      for (i = 0; i < (int) gimple_omp_for_collapse (gfor); i++)
+	{
+	  t = unshare_expr (gimple_omp_for_index (gfor, i));
+	  gimple_omp_for_set_index (gforo, i, t);
+	  t = unshare_expr (gimple_omp_for_initial (gfor, i));
+	  gimple_omp_for_set_initial (gforo, i, t);
+	  gimple_omp_for_set_cond (gforo, i,
+				   gimple_omp_for_cond (gfor, i));
+	  t = unshare_expr (gimple_omp_for_final (gfor, i));
+	  gimple_omp_for_set_final (gforo, i, t);
+	  t = unshare_expr (gimple_omp_for_incr (gfor, i));
+	  gimple_omp_for_set_incr (gforo, i, t);
+	}
+      gimplify_seq_add_stmt (pre_p, gforo);
+    }
+  else
+    gimplify_seq_add_stmt (pre_p, gfor);
   if (ret != GS_ALL_DONE)
     return GS_ERROR;
   *expr_p = NULL_TREE;
@@ -7511,9 +8417,11 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     case OMP_SINGLE:
       ort = ORT_WORKSHARE;
       break;
+    case OMP_TARGET:
+      ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
+      break;
     case OACC_KERNELS:
     case OACC_PARALLEL:
-    case OMP_TARGET:
       ort = ORT_TARGET;
       break;
     case OACC_DATA:
@@ -7526,8 +8434,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     default:
       gcc_unreachable ();
     }
-  gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
-  if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+  gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+			     TREE_CODE (expr));
+  if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
     {
       push_gimplify_context ();
       gimple *g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -7560,7 +8469,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     }
   else
     gimplify_and_add (OMP_BODY (expr), &body);
-  gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr), TREE_CODE (expr));
 
   switch (TREE_CODE (expr))
     {
@@ -7625,12 +8534,19 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
     case OMP_TARGET_UPDATE:
       kind = GF_OMP_TARGET_KIND_UPDATE;
       break;
+    case OMP_TARGET_ENTER_DATA:
+      kind = GF_OMP_TARGET_KIND_ENTER_DATA;
+      break;
+    case OMP_TARGET_EXIT_DATA:
+      kind = GF_OMP_TARGET_KIND_EXIT_DATA;
+      break;
     default:
       gcc_unreachable ();
     }
   gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
-			     ORT_WORKSHARE);
-  gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr));
+			     ORT_WORKSHARE, TREE_CODE (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
+			       TREE_CODE (expr));
   stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
 
   gimplify_seq_add_stmt (pre_p, stmt);
@@ -7845,6 +8761,93 @@ gimplify_transaction (tree *expr_p, gimple_seq *pre_p)
   return GS_ALL_DONE;
 }
 
+/* Gimplify an OMP_ORDERED construct.  EXPR is the tree version.  BODY
+   is the OMP_BODY of the original EXPR (which has already been
+   gimplified so it's not present in the EXPR).
+
+   Return the gimplified GIMPLE_OMP_ORDERED tuple.  */
+
+static gimple *
+gimplify_omp_ordered (tree expr, gimple_seq body)
+{
+  tree c, decls;
+  int failures = 0;
+  unsigned int i;
+  tree source_c = NULL_TREE;
+  tree sink_c = NULL_TREE;
+
+  if (gimplify_omp_ctxp)
+    for (c = OMP_ORDERED_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+	  && gimplify_omp_ctxp->loop_iter_var.is_empty ()
+	  && (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK
+	      || OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE))
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c),
+		    "%<depend%> clause must be closely nested "
+		    "inside a loop with %<ordered%> clause with "
+		    "a parameter");
+	  failures++;
+	}
+      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+	       && OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
+	{
+	  bool fail = false;
+	  for (decls = OMP_CLAUSE_DECL (c), i = 0;
+	       decls && TREE_CODE (decls) == TREE_LIST;
+	       decls = TREE_CHAIN (decls), ++i)
+	    if (i >= gimplify_omp_ctxp->loop_iter_var.length () / 2)
+	      continue;
+	    else if (TREE_VALUE (decls)
+		     != gimplify_omp_ctxp->loop_iter_var[2 * i])
+	      {
+		error_at (OMP_CLAUSE_LOCATION (c),
+			  "variable %qE is not an iteration "
+			  "of outermost loop %d, expected %qE",
+			  TREE_VALUE (decls), i + 1,
+			  gimplify_omp_ctxp->loop_iter_var[2 * i]);
+		fail = true;
+		failures++;
+	      }
+	    else
+	      TREE_VALUE (decls)
+		= gimplify_omp_ctxp->loop_iter_var[2 * i + 1];
+	  if (!fail && i != gimplify_omp_ctxp->loop_iter_var.length () / 2)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"number of variables in %<depend(sink)%> "
+			"clause does not match number of "
+			"iteration variables");
+	      failures++;
+	    }
+	  sink_c = c;
+	}
+      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+	       && OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE)
+	{
+	  if (source_c)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"more than one %<depend(source)%> clause on an "
+			"%<ordered%> construct");
+	      failures++;
+	    }
+	  else
+	    source_c = c;
+	}
+  if (source_c && sink_c)
+    {
+      error_at (OMP_CLAUSE_LOCATION (source_c),
+		"%<depend(source)%> clause specified together with "
+		"%<depend(sink:)%> clauses on the same construct");
+      failures++;
+    }
+
+  if (failures)
+    return gimple_build_nop ();
+  return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
+}
+
 /* Convert the GENERIC expression tree *EXPR_P to GIMPLE.  If the
    expression produces a value to be used as an operand inside a GIMPLE
    statement, the value will be stored back in *EXPR_P.  This value will
@@ -8574,6 +9577,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case CILK_SIMD:
 	case CILK_FOR:
 	case OMP_DISTRIBUTE:
+	case OMP_TASKLOOP:
 	case OACC_LOOP:
 	  ret = gimplify_omp_for (expr_p, pre_p);
 	  break;
@@ -8619,6 +9623,8 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case OACC_EXIT_DATA:
 	case OACC_UPDATE:
 	case OMP_TARGET_UPDATE:
+	case OMP_TARGET_ENTER_DATA:
+	case OMP_TARGET_EXIT_DATA:
 	  gimplify_omp_target_update (expr_p, pre_p);
 	  ret = GS_ALL_DONE;
 	  break;
@@ -8655,11 +9661,17 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 		}
 		break;
 	      case OMP_ORDERED:
-		g = gimple_build_omp_ordered (body);
+		g = gimplify_omp_ordered (*expr_p, body);
 		break;
 	      case OMP_CRITICAL:
+		gimplify_scan_omp_clauses (&OMP_CRITICAL_CLAUSES (*expr_p),
+					   pre_p, ORT_WORKSHARE, OMP_CRITICAL);
+		gimplify_adjust_omp_clauses (pre_p,
+					     &OMP_CRITICAL_CLAUSES (*expr_p),
+					     OMP_CRITICAL);
 		g = gimple_build_omp_critical (body,
-		    			       OMP_CRITICAL_NAME (*expr_p));
+		    			       OMP_CRITICAL_NAME (*expr_p),
+		    			       OMP_CRITICAL_CLAUSES (*expr_p));
 		break;
 	      default:
 		gcc_unreachable ();