[gomp4] Accelerator constructs omp lowering and expansion

Hi!

This patch implements #pragma omp {target{, data, update},teams} lowering
and expansion, and adds stub calls into libgomp, so that (for now
unconditionally) we can at least always fall back to host execution.

I still need to add testcases and test reference based array section
handling and look at mapping of VLAs.

Any issues you can spot in the patch?  If not, I'll commit it tomorrow, so
that I can continue poking in these areas without too much pending stuff.

Known issues that don't affect host fallback, but will be a problem
for offloading (Richard or others, any thoughts about that?):
1) right now we pass a host fn address and a function name string
   to GOMP_target.  For host execution we of course only need the fn
   address, for offloading I guess libgomp will need to dlinfo
   that fn address to find out what shared library (or binary) contains
   it and somehow look for the special data section in it.
   But, right now the *._omp_fn.N functions are always static, so there
   is nothing to look up by name for Intel MIC, and if we wanted to make
   the symbol global, we'd need to give it a shlib resp. binary unique
   name; but if the containing function is not global, how can we do that?
   Append get_file_function_name to the name?
2) much bigger problem seems to be global #pragma omp declare target
   variables.  Those are supposed to be mapped from the beginning,
   if they are just copied into the target LTO subset streaming, they will
   be emitted normally as data variables.  But, unfortunately the
   runtime must be aware of those mappings, because you can do stuff like:
#pragma omp declare target
int v = 1;
#pragma omp end declare target
void bar (int *p)
{
  #pragma omp target map(to:p[:1])
  *p++;
}
void foo () { bar (&v); }
   where the runtime should assign target's p copy the value of target's
   v variable.  Or even for say #pragma omp target update.
   So, on the host side, we need to prepare the triplets of host var
   address, var size and during linking somehow supply it info on how
   to create the target address, and have some function (one per
   shared library resp. binary) that would locate the data section
   within the shared library/binary for the requested accelerator
   and with another argument call some libgomp function to initialize
   the device data environment for the given shared library/binary
   and accelerator.  Perhaps we could pass the address of such (.hidden)
   function as yet another argument to GOMP_target call and
   --as-needed link it from some *.a library?

2013-09-04  Jakub Jelinek  <jakub@redhat.com>

	* tree-cfg.c (make_edges): For GIMPLE_OMP_TARGET
	with GF_OMP_TARGET_KIND_UPDATE, don't look for
	GIMPLE_OMP_RETURN and immediately restore previous
	region.
	* langhooks.c (lhd_omp_mappable_type): New function.
	* omp-low.c (scan_sharing_clauses): Ignore OMP_CLAUSE_SHARED
	in GIMPLE_OMP_TEAMS constructs.  Handle OMP_CLAUSE_NUM_TEAMS,
	OMP_CLAUSE_THREAD_LIMIT, OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP,
	OMP_CLAUSE_TO and OMP_CLAUSE_FROM.
	(create_omp_child_function): If current function has
	"omp declare target" attribute or if current region
	is OMP_TARGET or lexically nested in it, add that
	attribute to the omp child function.
	(scan_omp_target, scan_omp_teams): New functions.
	(check_omp_nesting_restrictions): Fix a typo in TEAMS nesting
	check.
	(scan_omp_1_stmt): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	(lower_rec_input_clauses): Ignore OMP_CLAUSE_SHARED
	in GIMPLE_OMP_TEAMS constructs.
	(expand_omp_synch): Handle GIMPLE_OMP_TEAMS.
	(expand_omp_target): New function.
	(expand_omp): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	(build_omp_regions_1): For GIMPLE_OMP_TARGET with
	GF_OMP_TARGET_KIND_UPDATE, don't look for GIMPLE_OMP_RETURN and
	immediately restore previous region.
	(lower_omp_single): Emit a CLOBBER stmt after GIMPLE_OMP_RETURN.
	(lower_omp_taskreg): Likewise.
	(lower_omp_target, lower_omp_teams): New functions.
	(lower_omp_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* tree.h (enum omp_clause_map_kind): Improve description of
	OMP_CLAUSE_MAP_POINTER.
	(OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION): Define.
	* gimple.def (GIMPLE_OMP_TARGET): Use GSS_OMP_PARALLEL instead
	of GSS_OMP_SINGLE.
	* langhooks.h (struct lang_hooks_for_types): Add
	omp_mappable_type hook.
	* langhooks-def.h (lhd_omp_mappable_type): New prototype.
	(LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
	(LANG_HOOKS_FOR_TYPES_INITIALIZER): Use it.
	* gimple.h (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
	gimple_omp_target_set_clauses): Use gimple_omp_parallel
	instead of gimple_omp_single.
	(gimple_omp_target_child_fn, gimple_omp_target_child_fn_ptr,
	gimple_omp_target_set_child_fn, gimple_omp_target_data_arg,
	gimple_omp_target_data_arg_ptr, gimple_omp_target_set_data_arg): New
	inlines.
	* omp-builtins.def (BUILT_IN_GOMP_TARGET, BUILT_IN_GOMP_TARGET_DATA,
	BUILT_IN_GOMP_TARGET_END_DATA, BUILT_IN_GOMP_TARGET_UPDATE,
	BUILT_IN_GOMP_TEAMS): New builtins.
	* gimple-pretty-print.c (dump_gimple_omp_target): Print child_fn
	for #pragma omp target.
	* gimplify.c (omp_notice_variable): Diagnose if implicitly mapped
	decl doesn't have mappable type.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_MAP don't clear
	notice_outer.
	(gimplify_adjust_omp_clauses_1): Add OMP_CLAUSE_MAP even if the same
	decl is already referenced in target data construct surrounding it.
	(gimplify_adjust_omp_clauses): Likewise.  Handle
	OMP_CLAUSE_THREAD_LIMIT.
	(gimplify_omp_workshare): Fix up gimplification of target or target
	data construct body.  For target data add GOMP_target_end_data
	call in a try/finally cleanup.
	* builtin-types.def (BT_FN_VOID_UINT_UINT,
	BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
c/
	* c-typeck.c (handle_omp_array_sections_1): Remove pointer_based_p
	argument.  Do c_save_expr on low bound unconditionally.
	(handle_omp_array_sections): Adjust caller.  Call c_fully_fold when
	needed.  Add OMP_CLAUSE_MAP_POINTER map clause even for array based
	array sections.
	(c_finish_omp_clauses): Don't complain about non-mappable types
	for OMP_CLAUSE_MAP_POINTER.  Check for duplicates even for map
	clauses.  Handle OMP_CLAUSE_NUM_TEAMS.
	* c-parser.c (c_parser_omp_target_data, c_parser_omp_target): Call
	keep_next_level ().
cp/
	* parser.c (cp_parser_omp_target_data, cp_parser_omp_target): Call
	keep_next_level (true).
	* cp-objcp-common.h (LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
	* semantics.c (handle_omp_array_sections_1): Remove pointer_based_p
	argument.  Do cp_save_expr on low bound unconditionally.
	(handle_omp_array_sections): Adjust caller.  Add OMP_CLAUSE_MAP_POINTER
	map clause even for array based array sections.
	(finish_omp_clauses): Don't complain about non-mappable types
	for OMP_CLAUSE_MAP_POINTER.  Check for duplicates even for map
	clauses.
fortran/
	* types.def (BT_FN_VOID_UINT_UINT, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
	BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
testsuite/
	* c-c++-common/gomp/map-1.c (foo): Add dg-error for implicitly
	mapped non-mappable var.
libgomp/
	* libgomp.map (GOMP_4.0): Add GOMP_target, GOMP_target_data,
	GOMP_target_end_data, GOMP_target_update and GOMP_teams.
	* Makefile.am (libgomp_la_SOURCES): Add target.c.
	* Makefile.in: Regenerated.
	* target.c: New file.
	* libgomp_g.h (GOMP_target, GOMP_target_data,
	GOMP_target_end_data, GOMP_target_update, GOMP_teams): New prototypes.
	* testsuite/libgomp.c++/for-11.C (main): Uncomment
	#pragma omp target teams directive.
	* testsuite/libgomp.c/for-3.c (main): Likewise.
	* testsuite/libgomp.c++/target-1.C: New test.
	* testsuite/libgomp.c/target-1.c: New test.


	Jakub

On Wed, Sep 04, 2013 at 08:54:47PM +0200, Jakub Jelinek wrote:
> Any issues you can spot in the patch?  If not, I'll commit it tomorrow, so
> that I can continue poking in these areas without too much pending stuff.

Committed now.

> Known issues that don't affect host fallback, but will be a problem
> for offloading (Richard or others, any thoughts about that?):
> 1) right now we pass a host fn address and a function name string
>    to GOMP_target.  For host execution we of course only need the fn
>    address, for offloading I guess libgomp will need to dlinfo
>    that fn address to find out what shared library (or binary) contains
>    it and somehow look for the special data section in it.
>    But, right now the *._omp_fn.N functions are always static, so there
>    is nothing to look up by name for Intel MIC, and if we wanted to make
>    the symbol global, we'd need to give it a shlib resp. binary unique
>    name; but if the containing function is not global, how can we do that?
>    Append get_file_function_name to the name?
> 2) much bigger problem seems to be global #pragma omp declare target
>    variables.  Those are supposed to be mapped from the beginning,
>    if they are just copied into the target LTO subset streaming, they will
>    be emitted normally as data variables.  But, unfortunately the
>    runtime must be aware of those mappings, because you can do stuff like:
> #pragma omp declare target
> int v = 1;
> #pragma omp end declare target
> void bar (int *p)
> {
>   #pragma omp target map(to:p[:1])
>   *p++;
> }
> void foo () { bar (&v); }
>    where the runtime should assign target's p copy the value of target's
>    v variable.  Or even for say #pragma omp target update.
>    So, on the host side, we need to prepare the triplets of host var
>    address, var size and during linking somehow supply it info on how
>    to create the target address, and have some function (one per
>    shared library resp. binary) that would locate the data section
>    within the shared library/binary for the requested accelerator
>    and with another argument call some libgomp function to initialize
>    the device data environment for the given shared library/binary
>    and accelerator.  Perhaps we could pass the address of such (.hidden)
>    function as yet another argument to GOMP_target call and
>    --as-needed link it from some *.a library?

Perhaps both of the above could be solved by (though, dunno about the
feasibility of doing it in the linker plugin):

1) GCC passes
extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden"), weak));
   &__OPENMP_TARGET__[0]
as the 3rd argument of GOMP_target call, rather than the name of the
function (and NULL on targets that don't support weak symbols or named
sections or visibility).  NULL will mean always fall back to host execution,
no support for offloading.
2) GCC emits struct { void *addr; uintptr_t size; } pairs into
   .gnu.openmp_target
   section for any non-extern vars declared in the TU with "omp declare target"
   attribute and for any *.omp_fn* functions passed to GOMP_target calls
   (with size 0).  For TREE_PUBLIC vars, perhaps we should emit !TREE_PUBLIC
   aliases and store those into the section instead.
3) during .gnu.target_lto* streaming, either those .gnu.openmp_target
   additions would be streamed too as is, or modified (only the addresses,
   leaving out sizes, which don't make sense)
4) when compiling target code and linking it, it would be important not to
   reorder the individual .gnu.openmp_target section additions in any way,
   so there is a 1:1 correspondence between the host array and target array
5) the lto linker plugin, when special handling the .gnu.target_lto*
   sections would also create the __OPENMP_TARGET__ symbol at the start
   of some magic section, fill in the section header at that point
   (likely some version number, number of supported targets, and for each
   of them address of the data blob in memory (or filename if it can't be
   in memory?), blob size, and pointer to the array of target addresses
   (unrelocated or something), plus the host .gnu.openmp_target content
   as array)
Then GOMP_target, if it sees NULL, would just do host fallback execution,
if it sees non-NULL in that argument and the device doesn't have that
loaded yet, it would read the info there to see if the target is supported,
find out how to load it (COI dlopen from memory, compile and ??? for
HSAIL/PTX), and initialize the mapping data structure from the host address,
size pairs in that array plus target array (after relocation?).
Then, for the function address, GOMP_target would just look up the function
address passed to it in the target mapping data structure.

	Jakub

On Wed, Sep 4, 2013 at 8:54 PM, Jakub Jelinek <jakub@redhat.com> wrote:
> Hi!
>
> This patch implements #pragma omp {target{, data, update},teams} lowering
> and expansion, and adds stub calls into libgomp, so that (for now
> unconditionally) we can at least always fall back to host execution.
>
> I still need to add testcases and test reference based array section
> handling and look at mapping of VLAs.
>
> Any issues you can spot in the patch?  If not, I'll commit it tomorrow, so
> that I can continue poking in these areas without too much pending stuff.
>
> Known issues that don't affect host fallback, but will be a problem
> for offloading (Richard or others, any thoughts about that?):
> 1) right now we pass a host fn address and a function name string
>    to GOMP_target.  For host execution we of course only need the fn
>    address, for offloading I guess libgomp will need to dlinfo
>    that fn address to find out what shared library (or binary) contains
>    it and somehow look for the special data section in it.
>    But, right now the *._omp_fn.N functions are always static, so there
>    is nothing to look up by name for Intel MIC, and if we wanted to make
>    the symbol global, we'd need to give it a shlib resp. binary unique
>    name; but if the containing function is not global, how can we do that?
>    Append get_file_function_name to the name?

I'd say we should pass the address of a descriptor instead that we
name specially.  Like

struct gomp_descr {
   void (*host_fn) ();
   void (*target_fns[]) ();
} __gomp_descr_N;

with N simply incrementing and the descriptor being local to the TU.
We initialize that with a reference to the outlined host function that
binds locally and external references to some private name (randomized
like we do for anonymous namespaces?).  The linker plugin then
needs to look for the descriptor symbols and from the host_fn address
lookup the function section with the LTO data for it.  From the extern
references it will know the symbol names to be used for the target(s)
implementation.

The descriptor above of course lacks metadata to identify the target fn
architecture.

> 2) much bigger problem seems to be global #pragma omp declare target
>    variables.  Those are supposed to be mapped from the beginning,
>    if they are just copied into the target LTO subset streaming, they will
>    be emitted normally as data variables.  But, unfortunately the
>    runtime must be aware of those mappings, because you can do stuff like:
> #pragma omp declare target
> int v = 1;
> #pragma omp end declare target
> void bar (int *p)
> {
>   #pragma omp target map(to:p[:1])
>   *p++;
> }
> void foo () { bar (&v); }
>    where the runtime should assign target's p copy the value of target's
>    v variable.  Or even for say #pragma omp target update.
>    So, on the host side, we need to prepare the triplets of host var
>    address, var size and during linking somehow supply it info on how
>    to create the target address, and have some function (one per
>    shared library resp. binary) that would locate the data section
>    within the shared library/binary for the requested accelerator
>    and with another argument call some libgomp function to initialize
>    the device data environment for the given shared library/binary
>    and accelerator.

A similar scheme as proposed above with the descriptor should work, no?
We can of course drop the descriptor as real data object and just stream
the symbol association in a special section.

Richard.

>    Perhaps we could pass the address of such (.hidden)
>    function as yet another argument to GOMP_target call and
>    --as-needed link it from some *.a library?
>
> 2013-09-04  Jakub Jelinek  <jakub@redhat.com>
>
>         * tree-cfg.c (make_edges): For GIMPLE_OMP_TARGET
>         with GF_OMP_TARGET_KIND_UPDATE, don't look for
>         GIMPLE_OMP_RETURN and immediately restore previous
>         region.
>         * langhooks.c (lhd_omp_mappable_type): New function.
>         * omp-low.c (scan_sharing_clauses): Ignore OMP_CLAUSE_SHARED
>         in GIMPLE_OMP_TEAMS constructs.  Handle OMP_CLAUSE_NUM_TEAMS,
>         OMP_CLAUSE_THREAD_LIMIT, OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP,
>         OMP_CLAUSE_TO and OMP_CLAUSE_FROM.
>         (create_omp_child_function): If current function has
>         "omp declare target" attribute or if current region
>         is OMP_TARGET or lexically nested in it, add that
>         attribute to the omp child function.
>         (scan_omp_target, scan_omp_teams): New functions.
>         (check_omp_nesting_restrictions): Fix a typo in TEAMS nesting
>         check.
>         (scan_omp_1_stmt): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
>         (lower_rec_input_clauses): Ignore OMP_CLAUSE_SHARED
>         in GIMPLE_OMP_TEAMS constructs.
>         (expand_omp_synch): Handle GIMPLE_OMP_TEAMS.
>         (expand_omp_target): New function.
>         (expand_omp): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
>         (build_omp_regions_1): For GIMPLE_OMP_TARGET with
>         GF_OMP_TARGET_KIND_UPDATE, don't look for GIMPLE_OMP_RETURN and
>         immediately restore previous region.
>         (lower_omp_single): Emit a CLOBBER stmt after GIMPLE_OMP_RETURN.
>         (lower_omp_taskreg): Likewise.
>         (lower_omp_target, lower_omp_teams): New functions.
>         (lower_omp_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
>         * tree.h (enum omp_clause_map_kind): Improve description of
>         OMP_CLAUSE_MAP_POINTER.
>         (OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION): Define.
>         * gimple.def (GIMPLE_OMP_TARGET): Use GSS_OMP_PARALLEL instead
>         of GSS_OMP_SINGLE.
>         * langhooks.h (struct lang_hooks_for_types): Add
>         omp_mappable_type hook.
>         * langhooks-def.h (lhd_omp_mappable_type): New prototype.
>         (LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
>         (LANG_HOOKS_FOR_TYPES_INITIALIZER): Use it.
>         * gimple.h (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
>         gimple_omp_target_set_clauses): Use gimple_omp_parallel
>         instead of gimple_omp_single.
>         (gimple_omp_target_child_fn, gimple_omp_target_child_fn_ptr,
>         gimple_omp_target_set_child_fn, gimple_omp_target_data_arg,
>         gimple_omp_target_data_arg_ptr, gimple_omp_target_set_data_arg): New
>         inlines.
>         * omp-builtins.def (BUILT_IN_GOMP_TARGET, BUILT_IN_GOMP_TARGET_DATA,
>         BUILT_IN_GOMP_TARGET_END_DATA, BUILT_IN_GOMP_TARGET_UPDATE,
>         BUILT_IN_GOMP_TEAMS): New builtins.
>         * gimple-pretty-print.c (dump_gimple_omp_target): Print child_fn
>         for #pragma omp target.
>         * gimplify.c (omp_notice_variable): Diagnose if implicitly mapped
>         decl doesn't have mappable type.
>         (gimplify_scan_omp_clauses): For OMP_CLAUSE_MAP don't clear
>         notice_outer.
>         (gimplify_adjust_omp_clauses_1): Add OMP_CLAUSE_MAP even if the same
>         decl is already referenced in target data construct surrounding it.
>         (gimplify_adjust_omp_clauses): Likewise.  Handle
>         OMP_CLAUSE_THREAD_LIMIT.
>         (gimplify_omp_workshare): Fix up gimplification of target or target
>         data construct body.  For target data add GOMP_target_end_data
>         call in a try/finally cleanup.
>         * builtin-types.def (BT_FN_VOID_UINT_UINT,
>         BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
>         BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
> c/
>         * c-typeck.c (handle_omp_array_sections_1): Remove pointer_based_p
>         argument.  Do c_save_expr on low bound unconditionally.
>         (handle_omp_array_sections): Adjust caller.  Call c_fully_fold when
>         needed.  Add OMP_CLAUSE_MAP_POINTER map clause even for array based
>         array sections.
>         (c_finish_omp_clauses): Don't complain about non-mappable types
>         for OMP_CLAUSE_MAP_POINTER.  Check for duplicates even for map
>         clauses.  Handle OMP_CLAUSE_NUM_TEAMS.
>         * c-parser.c (c_parser_omp_target_data, c_parser_omp_target): Call
>         keep_next_level ().
> cp/
>         * parser.c (cp_parser_omp_target_data, cp_parser_omp_target): Call
>         keep_next_level (true).
>         * cp-objcp-common.h (LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
>         * semantics.c (handle_omp_array_sections_1): Remove pointer_based_p
>         argument.  Do cp_save_expr on low bound unconditionally.
>         (handle_omp_array_sections): Adjust caller.  Add OMP_CLAUSE_MAP_POINTER
>         map clause even for array based array sections.
>         (finish_omp_clauses): Don't complain about non-mappable types
>         for OMP_CLAUSE_MAP_POINTER.  Check for duplicates even for map
>         clauses.
> fortran/
>         * types.def (BT_FN_VOID_UINT_UINT, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
>         BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
> testsuite/
>         * c-c++-common/gomp/map-1.c (foo): Add dg-error for implicitly
>         mapped non-mappable var.
> libgomp/
>         * libgomp.map (GOMP_4.0): Add GOMP_target, GOMP_target_data,
>         GOMP_target_end_data, GOMP_target_update and GOMP_teams.
>         * Makefile.am (libgomp_la_SOURCES): Add target.c.
>         * Makefile.in: Regenerated.
>         * target.c: New file.
>         * libgomp_g.h (GOMP_target, GOMP_target_data,
>         GOMP_target_end_data, GOMP_target_update, GOMP_teams): New prototypes.
>         * testsuite/libgomp.c++/for-11.C (main): Uncomment
>         #pragma omp target teams directive.
>         * testsuite/libgomp.c/for-3.c (main): Likewise.
>         * testsuite/libgomp.c++/target-1.C: New test.
>         * testsuite/libgomp.c/target-1.c: New test.
>
> --- gcc/tree-cfg.c.jj   2013-08-27 20:53:36.000000000 +0200
> +++ gcc/tree-cfg.c      2013-09-02 13:55:36.591737508 +0200
> @@ -610,7 +610,6 @@ make_edges (void)
>             case GIMPLE_OMP_TASK:
>             case GIMPLE_OMP_FOR:
>             case GIMPLE_OMP_SINGLE:
> -           case GIMPLE_OMP_TARGET:
>             case GIMPLE_OMP_TEAMS:
>             case GIMPLE_OMP_MASTER:
>             case GIMPLE_OMP_ORDERED:
> @@ -620,6 +619,13 @@ make_edges (void)
>               fallthru = true;
>               break;
>
> +           case GIMPLE_OMP_TARGET:
> +             cur_region = new_omp_region (bb, code, cur_region);
> +             fallthru = true;
> +             if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE)
> +               cur_region = cur_region->outer;
> +             break;
> +
>             case GIMPLE_OMP_SECTIONS:
>               cur_region = new_omp_region (bb, code, cur_region);
>               fallthru = true;
> --- gcc/langhooks.c.jj  2013-08-27 20:53:03.000000000 +0200
> +++ gcc/langhooks.c     2013-09-04 19:30:55.551423852 +0200
> @@ -523,6 +523,15 @@ lhd_omp_firstprivatize_type_sizes (struc
>  {
>  }
>
> +/* Return true if TYPE is an OpenMP mappable type.  By default return true
> +   if type is complete.  */
> +
> +bool
> +lhd_omp_mappable_type (tree type)
> +{
> +  return COMPLETE_TYPE_P (type);
> +}
> +
>  /* Common function for add_builtin_function and
>     add_builtin_function_ext_scope.  */
>  static tree
> --- gcc/omp-low.c.jj    2013-08-27 22:44:31.000000000 +0200
> +++ gcc/omp-low.c       2013-09-04 19:58:30.320019227 +0200
> @@ -1453,6 +1453,9 @@ scan_sharing_clauses (tree clauses, omp_
>           break;
>
>         case OMP_CLAUSE_SHARED:
> +         /* Ignore shared directives in teams construct.  */
> +         if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
> +           break;
>           gcc_assert (is_taskreg_ctx (ctx));
>           decl = OMP_CLAUSE_DECL (c);
>           gcc_assert (!COMPLETE_TYPE_P (TREE_TYPE (decl))
> @@ -1533,6 +1536,9 @@ scan_sharing_clauses (tree clauses, omp_
>         case OMP_CLAUSE_FINAL:
>         case OMP_CLAUSE_IF:
>         case OMP_CLAUSE_NUM_THREADS:
> +       case OMP_CLAUSE_NUM_TEAMS:
> +       case OMP_CLAUSE_THREAD_LIMIT:
> +       case OMP_CLAUSE_DEVICE:
>         case OMP_CLAUSE_SCHEDULE:
>         case OMP_CLAUSE_DIST_SCHEDULE:
>         case OMP_CLAUSE_DEPEND:
> @@ -1540,6 +1546,66 @@ scan_sharing_clauses (tree clauses, omp_
>             scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
>           break;
>
> +       case OMP_CLAUSE_TO:
> +       case OMP_CLAUSE_FROM:
> +       case OMP_CLAUSE_MAP:
> +         if (ctx->outer)
> +           scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
> +         decl = OMP_CLAUSE_DECL (c);
> +         /* Global variables with "omp declare target" attribute
> +            don't need to be copied, the receiver side will use them
> +            directly.  */
> +         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +             && DECL_P (decl)
> +             && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
> +             && lookup_attribute ("omp declare target",
> +                                  DECL_ATTRIBUTES (decl)))
> +           break;
> +         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +             && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
> +           {
> +             /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
> +                #pragma omp target data, there is nothing to map for
> +                those.  */
> +             if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
> +                 && !POINTER_TYPE_P (TREE_TYPE (decl)))
> +               break;
> +           }
> +         if (DECL_P (decl))
> +           {
> +             install_var_field (decl, true, 3, ctx);
> +             if (gimple_omp_target_kind (ctx->stmt)
> +                 == GF_OMP_TARGET_KIND_REGION)
> +               install_var_local (decl, ctx);
> +           }
> +         else
> +           {
> +             tree base = get_base_address (decl);
> +             tree nc = OMP_CLAUSE_CHAIN (c);
> +             if (DECL_P (base)
> +                 && nc != NULL_TREE
> +                 && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
> +                 && OMP_CLAUSE_DECL (nc) == base
> +                 && OMP_CLAUSE_MAP_KIND (nc) == OMP_CLAUSE_MAP_POINTER
> +                 && integer_zerop (OMP_CLAUSE_SIZE (nc)))
> +               {
> +                 OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) = 1;
> +                 OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (nc) = 1;
> +               }
> +             else
> +               {
> +                 gcc_assert (!splay_tree_lookup (ctx->field_map,
> +                                                 (splay_tree_key) decl));
> +                 tree field
> +                   = build_decl (OMP_CLAUSE_LOCATION (c),
> +                                 FIELD_DECL, NULL_TREE, ptr_type_node);
> +                 insert_field_into_struct (ctx->record_type, field);
> +                 splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
> +                                    (splay_tree_value) field);
> +               }
> +           }
> +         break;
> +
>         case OMP_CLAUSE_NOWAIT:
>         case OMP_CLAUSE_ORDERED:
>         case OMP_CLAUSE_COLLAPSE:
> @@ -1590,16 +1656,46 @@ scan_sharing_clauses (tree clauses, omp_
>           break;
>
>         case OMP_CLAUSE_SHARED:
> +         /* Ignore shared directives in teams construct.  */
> +         if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
> +           break;
>           decl = OMP_CLAUSE_DECL (c);
>           if (! is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
>             fixup_remapped_decl (decl, ctx, false);
>           break;
>
> +       case OMP_CLAUSE_MAP:
> +         if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
> +           break;
> +         decl = OMP_CLAUSE_DECL (c);
> +         if (DECL_P (decl)
> +             && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
> +             && lookup_attribute ("omp declare target",
> +                                  DECL_ATTRIBUTES (decl)))
> +           break;
> +         if (DECL_P (decl))
> +           {
> +             if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
> +                 && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
> +                 && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
> +               {
> +                 tree new_decl = lookup_decl (decl, ctx);
> +                 TREE_TYPE (new_decl)
> +                   = remap_type (TREE_TYPE (decl), &ctx->cb);
> +               }
> +             else
> +               fixup_remapped_decl (decl, ctx, false);
> +           }
> +         break;
> +
>         case OMP_CLAUSE_COPYPRIVATE:
>         case OMP_CLAUSE_COPYIN:
>         case OMP_CLAUSE_DEFAULT:
>         case OMP_CLAUSE_IF:
>         case OMP_CLAUSE_NUM_THREADS:
> +       case OMP_CLAUSE_NUM_TEAMS:
> +       case OMP_CLAUSE_THREAD_LIMIT:
> +       case OMP_CLAUSE_DEVICE:
>         case OMP_CLAUSE_SCHEDULE:
>         case OMP_CLAUSE_DIST_SCHEDULE:
>         case OMP_CLAUSE_NOWAIT:
> @@ -1613,6 +1709,8 @@ scan_sharing_clauses (tree clauses, omp_
>         case OMP_CLAUSE_ALIGNED:
>         case OMP_CLAUSE_DEPEND:
>         case OMP_CLAUSE__LOOPTEMP_:
> +       case OMP_CLAUSE_TO:
> +       case OMP_CLAUSE_FROM:
>           break;
>
>         default:
> @@ -1677,6 +1775,26 @@ create_omp_child_function (omp_context *
>    DECL_EXTERNAL (decl) = 0;
>    DECL_CONTEXT (decl) = NULL_TREE;
>    DECL_INITIAL (decl) = make_node (BLOCK);
> +  bool target_p = false;
> +  if (lookup_attribute ("omp declare target",
> +                       DECL_ATTRIBUTES (current_function_decl)))
> +    target_p = true;
> +  else
> +    {
> +      omp_context *octx;
> +      for (octx = ctx; octx; octx = octx->outer)
> +       if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
> +           && gimple_omp_target_kind (octx->stmt)
> +              == GF_OMP_TARGET_KIND_REGION)
> +         {
> +           target_p = true;
> +           break;
> +         }
> +    }
> +  if (target_p)
> +    DECL_ATTRIBUTES (decl)
> +      = tree_cons (get_identifier ("omp declare target"),
> +                  NULL_TREE, DECL_ATTRIBUTES (decl));
>
>    t = build_decl (DECL_SOURCE_LOCATION (decl),
>                   RESULT_DECL, NULL_TREE, void_type_node);
> @@ -1975,6 +2093,53 @@ scan_omp_single (gimple stmt, omp_contex
>      layout_type (ctx->record_type);
>  }
>
> +/* Scan an OpenMP target{, data, update} directive.  */
> +
> +static void
> +scan_omp_target (gimple stmt, omp_context *outer_ctx)
> +{
> +  omp_context *ctx;
> +  tree name;
> +  int kind = gimple_omp_target_kind (stmt);
> +
> +  ctx = new_omp_context (stmt, outer_ctx);
> +  ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
> +  ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
> +  ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE);
> +  name = create_tmp_var_name (".omp_data_t");
> +  name = build_decl (gimple_location (stmt),
> +                    TYPE_DECL, name, ctx->record_type);
> +  DECL_ARTIFICIAL (name) = 1;
> +  DECL_NAMELESS (name) = 1;
> +  TYPE_NAME (ctx->record_type) = name;
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      create_omp_child_function (ctx, false);
> +      gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
> +    }
> +
> +  scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
> +  scan_omp (gimple_omp_body_ptr (stmt), ctx);
> +
> +  if (TYPE_FIELDS (ctx->record_type) == NULL)
> +    ctx->record_type = ctx->receiver_decl = NULL;
> +  else
> +    {
> +      layout_type (ctx->record_type);
> +      if (kind == GF_OMP_TARGET_KIND_REGION)
> +       fixup_child_record_type (ctx);
> +    }
> +}
> +
> +/* Scan an OpenMP teams directive.  */
> +
> +static void
> +scan_omp_teams (gimple stmt, omp_context *outer_ctx)
> +{
> +  omp_context *ctx = new_omp_context (stmt, outer_ctx);
> +  scan_sharing_clauses (gimple_omp_teams_clauses (stmt), ctx);
> +  scan_omp (gimple_omp_body_ptr (stmt), ctx);
> +}
>
>  /* Check OpenMP nesting restrictions.  */
>  static bool
> @@ -1992,7 +2157,7 @@ check_omp_nesting_restrictions (gimple s
>        else if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
>         {
>           if ((gimple_code (stmt) != GIMPLE_OMP_FOR
> -              || (gimple_omp_for_kind (ctx->stmt)
> +              || (gimple_omp_for_kind (stmt)
>                    != GF_OMP_FOR_KIND_DISTRIBUTE))
>               && gimple_code (stmt) != GIMPLE_OMP_PARALLEL)
>             {
> @@ -2345,6 +2510,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *g
>        scan_omp (gimple_omp_body_ptr (stmt), ctx);
>        break;
>
> +    case GIMPLE_OMP_TARGET:
> +      scan_omp_target (stmt, ctx);
> +      break;
> +
> +    case GIMPLE_OMP_TEAMS:
> +      scan_omp_teams (stmt, ctx);
> +      break;
> +
>      case GIMPLE_BIND:
>        {
>         tree var;
> @@ -2731,6 +2904,9 @@ lower_rec_input_clauses (tree clauses, g
>                 continue;
>               break;
>             case OMP_CLAUSE_SHARED:
> +             /* Ignore shared directives in teams construct.  */
> +             if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
> +               continue;
>               if (maybe_lookup_decl (OMP_CLAUSE_DECL (c), ctx) == NULL)
>                 {
>                   gcc_assert (is_global_var (OMP_CLAUSE_DECL (c)));
> @@ -2889,6 +3065,9 @@ lower_rec_input_clauses (tree clauses, g
>           switch (OMP_CLAUSE_CODE (c))
>             {
>             case OMP_CLAUSE_SHARED:
> +             /* Ignore shared directives in teams construct.  */
> +             if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
> +               continue;
>               /* Shared global vars are just accessed directly.  */
>               if (is_global_var (new_var))
>                 break;
> @@ -6728,7 +6907,8 @@ expand_omp_synch (struct omp_region *reg
>    gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_SINGLE
>               || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_MASTER
>               || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_ORDERED
> -             || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL);
> +             || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL
> +             || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_TEAMS);
>    gsi_remove (&si, true);
>    single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
>
> @@ -7317,6 +7497,318 @@ expand_omp_atomic (struct omp_region *re
>  }
>
>
> +/* Expand the OpenMP target{, data, update} directive starting at REGION.  */
> +
> +static void
> +expand_omp_target (struct omp_region *region)
> +{
> +  basic_block entry_bb, exit_bb, new_bb;
> +  struct function *child_cfun = NULL;
> +  tree child_fn = NULL_TREE, block, t;
> +  gimple_stmt_iterator gsi;
> +  gimple entry_stmt, stmt;
> +  edge e;
> +
> +  entry_stmt = last_stmt (region->entry);
> +  new_bb = region->entry;
> +  int kind = gimple_omp_target_kind (entry_stmt);
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      child_fn = gimple_omp_target_child_fn (entry_stmt);
> +      child_cfun = DECL_STRUCT_FUNCTION (child_fn);
> +    }
> +
> +  entry_bb = region->entry;
> +  exit_bb = region->exit;
> +
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      unsigned srcidx, dstidx, num;
> +
> +      /* If the target region needs data sent from the parent
> +        function, then the very first statement (except possible
> +        tree profile counter updates) of the parallel body
> +        is a copy assignment .OMP_DATA_I = &.OMP_DATA_O.  Since
> +        &.OMP_DATA_O is passed as an argument to the child function,
> +        we need to replace it with the argument as seen by the child
> +        function.
> +
> +        In most cases, this will end up being the identity assignment
> +        .OMP_DATA_I = .OMP_DATA_I.  However, if the parallel body had
> +        a function call that has been inlined, the original PARM_DECL
> +        .OMP_DATA_I may have been converted into a different local
> +        variable.  In which case, we need to keep the assignment.  */
> +      if (gimple_omp_target_data_arg (entry_stmt))
> +       {
> +         basic_block entry_succ_bb = single_succ (entry_bb);
> +         gimple_stmt_iterator gsi;
> +         tree arg;
> +         gimple tgtcopy_stmt = NULL;
> +         tree sender
> +           = TREE_VEC_ELT (gimple_omp_target_data_arg (entry_stmt), 0);
> +
> +         for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
> +           {
> +             gcc_assert (!gsi_end_p (gsi));
> +             stmt = gsi_stmt (gsi);
> +             if (gimple_code (stmt) != GIMPLE_ASSIGN)
> +               continue;
> +
> +             if (gimple_num_ops (stmt) == 2)
> +               {
> +                 tree arg = gimple_assign_rhs1 (stmt);
> +
> +                 /* We're ignoring the subcode because we're
> +                    effectively doing a STRIP_NOPS.  */
> +
> +                 if (TREE_CODE (arg) == ADDR_EXPR
> +                     && TREE_OPERAND (arg, 0) == sender)
> +                   {
> +                     tgtcopy_stmt = stmt;
> +                     break;
> +                   }
> +               }
> +           }
> +
> +         gcc_assert (tgtcopy_stmt != NULL);
> +         arg = DECL_ARGUMENTS (child_fn);
> +
> +         gcc_assert (gimple_assign_lhs (tgtcopy_stmt) == arg);
> +         gsi_remove (&gsi, true);
> +       }
> +
> +      /* Declare local variables needed in CHILD_CFUN.  */
> +      block = DECL_INITIAL (child_fn);
> +      BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
> +      /* The gimplifier could record temporaries in target block
> +        rather than in containing function's local_decls chain,
> +        which would mean cgraph missed finalizing them.  Do it now.  */
> +      for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
> +       if (TREE_CODE (t) == VAR_DECL
> +           && TREE_STATIC (t)
> +           && !DECL_EXTERNAL (t))
> +         varpool_finalize_decl (t);
> +      DECL_SAVED_TREE (child_fn) = NULL;
> +      /* We'll create a CFG for child_fn, so no gimple body is needed.  */
> +      gimple_set_body (child_fn, NULL);
> +      TREE_USED (block) = 1;
> +
> +      /* Reset DECL_CONTEXT on function arguments.  */
> +      for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
> +       DECL_CONTEXT (t) = child_fn;
> +
> +      /* Split ENTRY_BB at GIMPLE_OMP_TARGET,
> +        so that it can be moved to the child function.  */
> +      gsi = gsi_last_bb (entry_bb);
> +      stmt = gsi_stmt (gsi);
> +      gcc_assert (stmt && gimple_code (stmt) == GIMPLE_OMP_TARGET
> +                 && gimple_omp_target_kind (stmt)
> +                    == GF_OMP_TARGET_KIND_REGION);
> +      gsi_remove (&gsi, true);
> +      e = split_block (entry_bb, stmt);
> +      entry_bb = e->dest;
> +      single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
> +
> +      /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR.  */
> +      if (exit_bb)
> +       {
> +         gsi = gsi_last_bb (exit_bb);
> +         gcc_assert (!gsi_end_p (gsi)
> +                     && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
> +         stmt = gimple_build_return (NULL);
> +         gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
> +         gsi_remove (&gsi, true);
> +       }
> +
> +      /* Move the target region into CHILD_CFUN.  */
> +
> +      block = gimple_block (entry_stmt);
> +
> +      new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block);
> +      if (exit_bb)
> +       single_succ_edge (new_bb)->flags = EDGE_FALLTHRU;
> +      /* When the OMP expansion process cannot guarantee an up-to-date
> +        loop tree arrange for the child function to fixup loops.  */
> +      if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
> +       child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP;
> +
> +      /* Remove non-local VAR_DECLs from child_cfun->local_decls list.  */
> +      num = vec_safe_length (child_cfun->local_decls);
> +      for (srcidx = 0, dstidx = 0; srcidx < num; srcidx++)
> +       {
> +         t = (*child_cfun->local_decls)[srcidx];
> +         if (DECL_CONTEXT (t) == cfun->decl)
> +           continue;
> +         if (srcidx != dstidx)
> +           (*child_cfun->local_decls)[dstidx] = t;
> +         dstidx++;
> +       }
> +      if (dstidx != num)
> +       vec_safe_truncate (child_cfun->local_decls, dstidx);
> +
> +      /* Inform the callgraph about the new function.  */
> +      DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
> +      cgraph_add_new_function (child_fn, true);
> +
> +      /* Fix the callgraph edges for child_cfun.  Those for cfun will be
> +        fixed in a following pass.  */
> +      push_cfun (child_cfun);
> +      rebuild_cgraph_edges ();
> +
> +      /* Some EH regions might become dead, see PR34608.  If
> +        pass_cleanup_cfg isn't the first pass to happen with the
> +        new child, these dead EH edges might cause problems.
> +        Clean them up now.  */
> +      if (flag_exceptions)
> +       {
> +         basic_block bb;
> +         bool changed = false;
> +
> +         FOR_EACH_BB (bb)
> +           changed |= gimple_purge_dead_eh_edges (bb);
> +         if (changed)
> +           cleanup_tree_cfg ();
> +       }
> +      pop_cfun ();
> +    }
> +
> +  /* Emit a library call to launch the target region, or do data
> +     transfers.  */
> +  tree t1, t2, t3, t4, device, cond, c, clauses;
> +  enum built_in_function start_ix;
> +  location_t clause_loc;
> +
> +  clauses = gimple_omp_target_clauses (entry_stmt);
> +
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    start_ix = BUILT_IN_GOMP_TARGET;
> +  else if (kind == GF_OMP_TARGET_KIND_DATA)
> +    start_ix = BUILT_IN_GOMP_TARGET_DATA;
> +  else
> +    start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
> +
> +  /* By default, the value of DEVICE is -1 (let runtime library choose)
> +     and there is no conditional.  */
> +  cond = NULL_TREE;
> +  device = build_int_cst (integer_type_node, -1);
> +
> +  c = find_omp_clause (clauses, OMP_CLAUSE_IF);
> +  if (c)
> +    cond = OMP_CLAUSE_IF_EXPR (c);
> +
> +  c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
> +  if (c)
> +    {
> +      device = OMP_CLAUSE_DEVICE_ID (c);
> +      clause_loc = OMP_CLAUSE_LOCATION (c);
> +    }
> +  else
> +    clause_loc = gimple_location (entry_stmt);
> +
> +  /* Ensure 'device' is of the correct type.  */
> +  device = fold_convert_loc (clause_loc, integer_type_node, device);
> +
> +  /* If we found the clause 'if (cond)', build
> +     (cond ? device : -2).  */
> +  if (cond)
> +    {
> +      cond = gimple_boolify (cond);
> +
> +      basic_block cond_bb, then_bb, else_bb;
> +      edge e;
> +      tree tmp_var;
> +
> +      tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
> +      e = split_block (entry_bb, NULL);
> +      cond_bb = e->src;
> +      entry_bb = e->dest;
> +      remove_edge (e);
> +
> +      then_bb = create_empty_bb (cond_bb);
> +      else_bb = create_empty_bb (then_bb);
> +      set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
> +      set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
> +
> +      stmt = gimple_build_cond_empty (cond);
> +      gsi = gsi_start_bb (cond_bb);
> +      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
> +
> +      gsi = gsi_start_bb (then_bb);
> +      stmt = gimple_build_assign (tmp_var, device);
> +      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
> +
> +      gsi = gsi_start_bb (else_bb);
> +      stmt = gimple_build_assign (tmp_var,
> +                                 build_int_cst (integer_type_node, -2));
> +      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
> +
> +      make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
> +      make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
> +      if (current_loops)
> +       {
> +         add_bb_to_loop (then_bb, cond_bb->loop_father);
> +         add_bb_to_loop (else_bb, cond_bb->loop_father);
> +       }
> +      make_edge (then_bb, entry_bb, EDGE_FALLTHRU);
> +      make_edge (else_bb, entry_bb, EDGE_FALLTHRU);
> +
> +      device = tmp_var;
> +    }
> +
> +  gsi = gsi_last_bb (new_bb);
> +  t = gimple_omp_target_data_arg (entry_stmt);
> +  if (t == NULL)
> +    {
> +      t1 = size_zero_node;
> +      t2 = build_zero_cst (ptr_type_node);
> +      t3 = t2;
> +      t4 = t2;
> +    }
> +  else
> +    {
> +      t1 = TYPE_MAX_VALUE (TYPE_DOMAIN (TREE_TYPE (TREE_VEC_ELT (t, 1))));
> +      t1 = size_binop (PLUS_EXPR, t1, size_int (1));
> +      t2 = build_fold_addr_expr (TREE_VEC_ELT (t, 0));
> +      t3 = build_fold_addr_expr (TREE_VEC_ELT (t, 1));
> +      t4 = build_fold_addr_expr (TREE_VEC_ELT (t, 2));
> +    }
> +
> +  gimple g;
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      tree fnaddr = build_fold_addr_expr (child_fn);
> +      unsigned fnnamelen = IDENTIFIER_LENGTH (DECL_NAME (child_fn));
> +      tree fnname = build_string (fnnamelen,
> +                                 IDENTIFIER_POINTER (DECL_NAME (child_fn)));
> +      TREE_TYPE (fnname) = build_array_type_nelts (char_type_node,
> +                                                  fnnamelen);
> +      TREE_READONLY (fnname) = 1;
> +      TREE_STATIC (fnname) = 1;
> +      fnname = build_fold_addr_expr (fnname);
> +      g = gimple_build_call (builtin_decl_explicit (start_ix), 7,
> +                            device, fnaddr, fnname, t1, t2, t3, t4);
> +    }
> +  else
> +    g = gimple_build_call (builtin_decl_explicit (start_ix), 5,
> +                          device, t1, t2, t3, t4);
> +  gimple_set_location (g, gimple_location (entry_stmt));
> +  gsi_insert_before (&gsi, g, GSI_SAME_STMT);
> +  if (kind != GF_OMP_TARGET_KIND_REGION)
> +    {
> +      g = gsi_stmt (gsi);
> +      gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
> +      gsi_remove (&gsi, true);
> +    }
> +  if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
> +    {
> +      gsi = gsi_last_bb (region->exit);
> +      g = gsi_stmt (gsi);
> +      gcc_assert (g && gimple_code (g) == GIMPLE_OMP_RETURN);
> +      gsi_remove (&gsi, true);
> +    }
> +}
> +
> +
>  /* Expand the parallel region tree rooted at REGION.  Expansion
>     proceeds in depth-first order.  Innermost regions are expanded
>     first.  This way, parallel regions that require a new function to
> @@ -7374,6 +7866,7 @@ expand_omp (struct omp_region *region)
>         case GIMPLE_OMP_MASTER:
>         case GIMPLE_OMP_ORDERED:
>         case GIMPLE_OMP_CRITICAL:
> +       case GIMPLE_OMP_TEAMS:
>           expand_omp_synch (region);
>           break;
>
> @@ -7381,6 +7874,10 @@ expand_omp (struct omp_region *region)
>           expand_omp_atomic (region);
>           break;
>
> +       case GIMPLE_OMP_TARGET:
> +         expand_omp_target (region);
> +         break;
> +
>         default:
>           gcc_unreachable ();
>         }
> @@ -7445,6 +7942,9 @@ build_omp_regions_1 (basic_block bb, str
>              GIMPLE_OMP_SECTIONS, and we do nothing for it.  */
>           ;
>         }
> +      else if (code == GIMPLE_OMP_TARGET
> +              && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE)
> +       new_omp_region (bb, code, parent);
>        else
>         {
>           /* Otherwise, this directive becomes the parent for a new
> @@ -7825,7 +8325,7 @@ lower_omp_single (gimple_stmt_iterator *
>  {
>    tree block;
>    gimple t, bind, single_stmt = gsi_stmt (*gsi_p);
> -  gimple_seq bind_body, dlist;
> +  gimple_seq bind_body, bind_body_tail = NULL, dlist;
>    struct gimplify_ctx gctx;
>
>    push_gimplify_context (&gctx);
> @@ -7855,8 +8355,17 @@ lower_omp_single (gimple_stmt_iterator *
>    t = gimple_build_omp_return
>          (!!find_omp_clause (gimple_omp_single_clauses (single_stmt),
>                             OMP_CLAUSE_NOWAIT));
> -  gimple_seq_add_stmt (&bind_body, t);
> -  maybe_add_implicit_barrier_cancel (ctx, &bind_body);
> +  gimple_seq_add_stmt (&bind_body_tail, t);
> +  maybe_add_implicit_barrier_cancel (ctx, &bind_body_tail);
> +  if (ctx->record_type)
> +    {
> +      gimple_stmt_iterator gsi = gsi_start (bind_body_tail);
> +      tree clobber = build_constructor (ctx->record_type, NULL);
> +      TREE_THIS_VOLATILE (clobber) = 1;
> +      gsi_insert_after (&gsi, gimple_build_assign (ctx->sender_decl,
> +                                                  clobber), GSI_SAME_STMT);
> +    }
> +  gimple_seq_add_seq (&bind_body, bind_body_tail);
>    gimple_bind_set_body (bind, bind_body);
>
>    pop_gimplify_context (bind);
> @@ -8605,6 +9114,14 @@ lower_omp_taskreg (gimple_stmt_iterator
>    lower_send_clauses (clauses, &ilist, &olist, ctx);
>    lower_send_shared_vars (&ilist, &olist, ctx);
>
> +  if (ctx->record_type)
> +    {
> +      tree clobber = build_constructor (TREE_TYPE (ctx->sender_decl), NULL);
> +      TREE_THIS_VOLATILE (clobber) = 1;
> +      gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
> +                                                       clobber));
> +    }
> +
>    /* Once all the expansions are done, sequence all the different
>       fragments inside gimple_omp_body.  */
>
> @@ -8637,6 +9154,326 @@ lower_omp_taskreg (gimple_stmt_iterator
>    pop_gimplify_context (NULL);
>  }
>
> +/* Lower the OpenMP target directive in the current statement
> +   in GSI_P.  CTX holds context information for the directive.  */
> +
> +static void
> +lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> +{
> +  tree clauses;
> +  tree child_fn, t, c;
> +  gimple stmt = gsi_stmt (*gsi_p);
> +  gimple tgt_bind = NULL, bind;
> +  gimple_seq tgt_body = NULL, olist, ilist, new_body;
> +  struct gimplify_ctx gctx;
> +  location_t loc = gimple_location (stmt);
> +  int kind = gimple_omp_target_kind (stmt);
> +  unsigned int map_cnt = 0;
> +
> +  clauses = gimple_omp_target_clauses (stmt);
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
> +      tgt_body = gimple_bind_body (tgt_bind);
> +    }
> +  else if (kind == GF_OMP_TARGET_KIND_DATA)
> +    tgt_body = gimple_omp_body (stmt);
> +  child_fn = ctx->cb.dst_fn;
> +
> +  push_gimplify_context (&gctx);
> +
> +  for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
> +    switch (OMP_CLAUSE_CODE (c))
> +      {
> +       tree var, x;
> +
> +      default:
> +       break;
> +      case OMP_CLAUSE_MAP:
> +      case OMP_CLAUSE_TO:
> +      case OMP_CLAUSE_FROM:
> +       var = OMP_CLAUSE_DECL (c);
> +       if (!DECL_P (var))
> +         {
> +           if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
> +               || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
> +             map_cnt++;
> +           continue;
> +         }
> +       if (!lookup_sfield (var, ctx))
> +         continue;
> +
> +       if (kind == GF_OMP_TARGET_KIND_REGION)
> +         {
> +           x = build_receiver_ref (var, true, ctx);
> +           tree new_var = lookup_decl (var, ctx);
> +           SET_DECL_VALUE_EXPR (new_var, x);
> +           DECL_HAS_VALUE_EXPR_P (new_var) = 1;
> +         }
> +       map_cnt++;
> +      }
> +
> +  if (kind != GF_OMP_TARGET_KIND_UPDATE)
> +    lower_omp (&tgt_body, ctx);
> +
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      /* Declare all the variables created by mapping and the variables
> +        declared in the scope of the target body.  */
> +      record_vars_into (ctx->block_vars, child_fn);
> +      record_vars_into (gimple_bind_vars (tgt_bind), child_fn);
> +    }
> +
> +  olist = NULL;
> +  ilist = NULL;
> +  if (ctx->record_type)
> +    {
> +      ctx->sender_decl
> +       = create_tmp_var (ctx->record_type, ".omp_data_arr");
> +      DECL_NAMELESS (ctx->sender_decl) = 1;
> +      TREE_ADDRESSABLE (ctx->sender_decl) = 1;
> +      t = make_tree_vec (3);
> +      TREE_VEC_ELT (t, 0) = ctx->sender_decl;
> +      TREE_VEC_ELT (t, 1)
> +       = create_tmp_var (build_array_type_nelts (size_type_node, map_cnt),
> +                         ".omp_data_sizes");
> +      DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
> +      TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
> +      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
> +      TREE_VEC_ELT (t, 2)
> +       = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
> +                                                 map_cnt),
> +                         ".omp_data_kinds");
> +      DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
> +      TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
> +      TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
> +      gimple_omp_target_set_data_arg (stmt, t);
> +
> +      vec<constructor_elt, va_gc> *vsize;
> +      vec<constructor_elt, va_gc> *vkind;
> +      vec_alloc (vsize, map_cnt);
> +      vec_alloc (vkind, map_cnt);
> +      unsigned int map_idx = 0;
> +
> +      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
> +       switch (OMP_CLAUSE_CODE (c))
> +         {
> +           tree ovar, nc;
> +
> +         default:
> +           break;
> +         case OMP_CLAUSE_MAP:
> +         case OMP_CLAUSE_TO:
> +         case OMP_CLAUSE_FROM:
> +           nc = c;
> +           ovar = OMP_CLAUSE_DECL (c);
> +           if (!DECL_P (ovar))
> +             {
> +               if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +                   && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
> +                 {
> +                   gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c))
> +                                        == get_base_address (ovar));
> +                   nc = OMP_CLAUSE_CHAIN (c);
> +                   ovar = OMP_CLAUSE_DECL (nc);
> +                 }
> +               else
> +                 {
> +                   tree x = build_sender_ref (ovar, ctx);
> +                   tree v
> +                     = build_fold_addr_expr_with_type (ovar, ptr_type_node);
> +                   gimplify_assign (x, v, &ilist);
> +                   nc = NULL_TREE;
> +                 }
> +             }
> +           else if (!lookup_sfield (ovar, ctx))
> +             continue;
> +
> +           if (nc)
> +             {
> +               tree var = lookup_decl_in_outer_ctx (ovar, ctx);
> +               tree x = build_sender_ref (ovar, ctx);
> +               if (is_gimple_reg (var))
> +                 {
> +                   gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
> +                   tree avar = create_tmp_var (TREE_TYPE (var), NULL);
> +                   mark_addressable (avar);
> +                   if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
> +                       && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
> +                     gimplify_assign (avar, var, &ilist);
> +                   avar = build_fold_addr_expr (avar);
> +                   gimplify_assign (x, avar, &ilist);
> +                   if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
> +                        || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
> +                       && !TYPE_READONLY (TREE_TYPE (var)))
> +                     {
> +                       x = build_sender_ref (ovar, ctx);
> +                       x = build_simple_mem_ref (x);
> +                       gimplify_assign (var, x, &olist);
> +                     }
> +                 }
> +               else
> +                 {
> +                   var = build_fold_addr_expr (var);
> +                   gimplify_assign (x, var, &ilist);
> +                 }
> +             }
> +           tree s = OMP_CLAUSE_SIZE (c);
> +           if (s == NULL_TREE)
> +             s = TYPE_SIZE (TREE_TYPE (ovar));
> +           s = fold_convert (size_type_node, s);
> +           tree purpose = size_int (map_idx++);
> +           CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
> +           if (TREE_CODE (s) != INTEGER_CST)
> +             TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
> +
> +           unsigned char tkind = 0;
> +           switch (OMP_CLAUSE_CODE (c))
> +             {
> +             case OMP_CLAUSE_MAP:
> +               tkind = OMP_CLAUSE_MAP_KIND (c);
> +               break;
> +             case OMP_CLAUSE_TO:
> +               tkind = OMP_CLAUSE_MAP_TO;
> +               break;
> +             case OMP_CLAUSE_FROM:
> +               tkind = OMP_CLAUSE_MAP_FROM;
> +               break;
> +             default:
> +               gcc_unreachable ();
> +             }
> +           CONSTRUCTOR_APPEND_ELT (vkind, purpose,
> +                                   build_int_cst (unsigned_char_type_node,
> +                                                  tkind));
> +           if (nc && nc != c)
> +             c = nc;
> +         }
> +
> +      gcc_assert (map_idx == map_cnt);
> +
> +      DECL_INITIAL (TREE_VEC_ELT (t, 1))
> +       = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
> +      DECL_INITIAL (TREE_VEC_ELT (t, 2))
> +       = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
> +      if (!TREE_STATIC (TREE_VEC_ELT (t, 1)))
> +       {
> +         gimple_seq initlist = NULL;
> +         force_gimple_operand (build1 (DECL_EXPR, void_type_node,
> +                                       TREE_VEC_ELT (t, 1)),
> +                               &initlist, true, NULL_TREE);
> +         gimple_seq_add_seq (&ilist, initlist);
> +       }
> +
> +      tree clobber = build_constructor (ctx->record_type, NULL);
> +      TREE_THIS_VOLATILE (clobber) = 1;
> +      gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
> +                                                       clobber));
> +    }
> +
> +  /* Once all the expansions are done, sequence all the different
> +     fragments inside gimple_omp_body.  */
> +
> +  new_body = NULL;
> +
> +  if (ctx->record_type && kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
> +      /* fixup_child_record_type might have changed receiver_decl's type.  */
> +      t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
> +      gimple_seq_add_stmt (&new_body,
> +                          gimple_build_assign (ctx->receiver_decl, t));
> +    }
> +
> +  if (kind == GF_OMP_TARGET_KIND_REGION)
> +    {
> +      gimple_seq_add_seq (&new_body, tgt_body);
> +      new_body = maybe_catch_exception (new_body);
> +    }
> +  else if (kind == GF_OMP_TARGET_KIND_DATA)
> +    new_body = tgt_body;
> +  if (kind != GF_OMP_TARGET_KIND_UPDATE)
> +    {
> +      gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
> +      gimple_omp_set_body (stmt, new_body);
> +    }
> +
> +  bind = gimple_build_bind (NULL, NULL,
> +                           tgt_bind ? gimple_bind_block (tgt_bind)
> +                                    : NULL_TREE);
> +  gsi_replace (gsi_p, bind, true);
> +  gimple_bind_add_seq (bind, ilist);
> +  gimple_bind_add_stmt (bind, stmt);
> +  gimple_bind_add_seq (bind, olist);
> +
> +  pop_gimplify_context (NULL);
> +}
> +
> +/* Expand code for an OpenMP teams directive.  */
> +
> +static void
> +lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> +{
> +  gimple teams_stmt = gsi_stmt (*gsi_p);
> +  struct gimplify_ctx gctx;
> +  push_gimplify_context (&gctx);
> +
> +  tree block = make_node (BLOCK);
> +  gimple bind = gimple_build_bind (NULL, NULL, block);
> +  gsi_replace (gsi_p, bind, true);
> +  gimple_seq bind_body = NULL;
> +  gimple_seq dlist = NULL;
> +  gimple_seq olist = NULL;
> +
> +  tree num_teams = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
> +                                   OMP_CLAUSE_NUM_TEAMS);
> +  if (num_teams == NULL_TREE)
> +    num_teams = build_int_cst (unsigned_type_node, 0);
> +  else
> +    {
> +      num_teams = OMP_CLAUSE_NUM_TEAMS_EXPR (num_teams);
> +      num_teams = fold_convert (unsigned_type_node, num_teams);
> +      gimplify_expr (&num_teams, &bind_body, NULL, is_gimple_val, fb_rvalue);
> +    }
> +  tree thread_limit = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
> +                                      OMP_CLAUSE_THREAD_LIMIT);
> +  if (thread_limit == NULL_TREE)
> +    thread_limit = build_int_cst (unsigned_type_node, 0);
> +  else
> +    {
> +      thread_limit = OMP_CLAUSE_THREAD_LIMIT_EXPR (thread_limit);
> +      thread_limit = fold_convert (unsigned_type_node, thread_limit);
> +      gimplify_expr (&thread_limit, &bind_body, NULL, is_gimple_val,
> +                    fb_rvalue);
> +    }
> +
> +  lower_rec_input_clauses (gimple_omp_teams_clauses (teams_stmt),
> +                          &bind_body, &dlist, ctx, NULL);
> +  lower_omp (gimple_omp_body_ptr (teams_stmt), ctx);
> +  lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist, ctx);
> +  gimple_seq_add_stmt (&bind_body, teams_stmt);
> +
> +  location_t loc = gimple_location (teams_stmt);
> +  tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS);
> +  gimple call = gimple_build_call (decl, 2, num_teams, thread_limit);
> +  gimple_set_location (call, loc);
> +  gimple_seq_add_stmt (&bind_body, call);
> +
> +  gimple_seq_add_seq (&bind_body, gimple_omp_body (teams_stmt));
> +  gimple_omp_set_body (teams_stmt, NULL);
> +  gimple_seq_add_seq (&bind_body, olist);
> +  gimple_seq_add_seq (&bind_body, dlist);
> +  gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true));
> +  gimple_bind_set_body (bind, bind_body);
> +
> +  pop_gimplify_context (bind);
> +
> +  gimple_bind_append_vars (bind, ctx->block_vars);
> +  BLOCK_VARS (block) = ctx->block_vars;
> +  if (BLOCK_VARS (block))
> +    TREE_USED (block) = 1;
> +}
> +
> +
>  /* Callback for lower_omp_1.  Return non-NULL if *tp needs to be
>     regimplified.  If DATA is non-NULL, lower_omp_1 is outside
>     of OpenMP context, but with task_shared_vars set.  */
> @@ -8760,6 +9597,16 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p
>                         lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
>         gimple_regimplify_operands (stmt, gsi_p);
>        break;
> +    case GIMPLE_OMP_TARGET:
> +      ctx = maybe_lookup_ctx (stmt);
> +      gcc_assert (ctx);
> +      lower_omp_target (gsi_p, ctx);
> +      break;
> +    case GIMPLE_OMP_TEAMS:
> +      ctx = maybe_lookup_ctx (stmt);
> +      gcc_assert (ctx);
> +      lower_omp_teams (gsi_p, ctx);
> +      break;
>      case GIMPLE_CALL:
>        tree fndecl;
>        fndecl = gimple_call_fndecl (stmt);
> --- gcc/c/c-typeck.c.jj 2013-08-27 20:49:05.000000000 +0200
> +++ gcc/c/c-typeck.c    2013-09-04 17:21:17.295112407 +0200
> @@ -10773,8 +10773,7 @@ c_finish_omp_cancellation_point (locatio
>
>  static tree
>  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
> -                            bool &maybe_zero_len, unsigned int &first_non_one,
> -                            bool &pointer_based_p)
> +                            bool &maybe_zero_len, unsigned int &first_non_one)
>  {
>    tree ret, low_bound, length, type;
>    if (TREE_CODE (t) != TREE_LIST)
> @@ -10801,15 +10800,11 @@ handle_omp_array_sections_1 (tree c, tre
>                     omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
>           return error_mark_node;
>         }
> -      if (POINTER_TYPE_P (TREE_TYPE (t))
> -         && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> -       pointer_based_p = true;
>        return t;
>      }
>
>    ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
> -                                    maybe_zero_len, first_non_one,
> -                                    pointer_based_p);
> +                                    maybe_zero_len, first_non_one);
>    if (ret == error_mark_node || ret == NULL_TREE)
>      return ret;
>
> @@ -10989,16 +10984,12 @@ handle_omp_array_sections_1 (tree c, tre
>      }
>    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
>      types.safe_push (TREE_TYPE (ret));
> -  /* For pointer based array sections we will need to evaluate lb more
> -     than once.  */
> -  if (pointer_based_p)
> +  /* We will need to evaluate lb more than once.  */
> +  tree lb = c_save_expr (low_bound);
> +  if (lb != low_bound)
>      {
> -      tree lb = c_save_expr (low_bound);
> -      if (lb != low_bound)
> -       {
> -         TREE_PURPOSE (t) = lb;
> -         low_bound = lb;
> -       }
> +      TREE_PURPOSE (t) = lb;
> +      low_bound = lb;
>      }
>    ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
>    return ret;
> @@ -11010,12 +11001,10 @@ static bool
>  handle_omp_array_sections (tree c)
>  {
>    bool maybe_zero_len = false;
> -  bool pointer_based_p = false;
>    unsigned int first_non_one = 0;
>    vec<tree> types = vNULL;
>    tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
> -                                           maybe_zero_len, first_non_one,
> -                                           pointer_based_p);
> +                                           maybe_zero_len, first_non_one);
>    if (first == error_mark_node)
>      {
>        types.release ();
> @@ -11047,6 +11036,7 @@ handle_omp_array_sections (tree c)
>         }
>        if (tem)
>         first = build2 (COMPOUND_EXPR, TREE_TYPE (first), tem, first);
> +      first = c_fully_fold (first, false, NULL);
>        OMP_CLAUSE_DECL (c) = first;
>      }
>    else
> @@ -11155,27 +11145,31 @@ handle_omp_array_sections (tree c)
>        types.release ();
>        if (side_effects)
>         size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
> +      first = c_fully_fold (first, false, NULL);
>        OMP_CLAUSE_DECL (c) = first;
> +      if (size)
> +       size = c_fully_fold (size, false, NULL);
>        OMP_CLAUSE_SIZE (c) = size;
> -      if (pointer_based_p)
> -       {
> -         tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
> -         OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
> -         if (!c_mark_addressable (t))
> -           return false;
> -         OMP_CLAUSE_DECL (c2) = t;
> -         t = build_fold_addr_expr (first);
> -         t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> -                               build_pointer_type (char_type_node), t);
> -         t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
> -                              ptrdiff_type_node, t,
> -                              fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> -                                                TREE_TYPE (t),
> -                                                OMP_CLAUSE_DECL (c2)));
> -         OMP_CLAUSE_SIZE (c2) = t;
> -         OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
> -         OMP_CLAUSE_CHAIN (c) = c2;
> -       }
> +      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
> +       return false;
> +      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
> +      OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
> +      if (!c_mark_addressable (t))
> +       return false;
> +      OMP_CLAUSE_DECL (c2) = t;
> +      t = build_fold_addr_expr (first);
> +      t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
> +      tree ptr = OMP_CLAUSE_DECL (c2);
> +      if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
> +       ptr = build_fold_addr_expr (ptr);
> +      t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
> +                          ptrdiff_type_node, t,
> +                          fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> +                                            ptrdiff_type_node, ptr));
> +      t = c_fully_fold (t, false, NULL);
> +      OMP_CLAUSE_SIZE (c2) = t;
> +      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
> +      OMP_CLAUSE_CHAIN (c) = c2;
>      }
>    return false;
>  }
> @@ -11462,18 +11456,21 @@ c_finish_omp_clauses (tree clauses)
>             }
>           else if (!c_mark_addressable (t))
>             remove = true;
> -         else if (!COMPLETE_TYPE_P (TREE_TYPE (t)))
> +         else if (!COMPLETE_TYPE_P (TREE_TYPE (t))
> +                  && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +                       && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER))
>             {
>               error_at (OMP_CLAUSE_LOCATION (c),
>                         "%qD does not have a mappable type in %qs clause", t,
>                         omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
>               remove = true;
>             }
> -         else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> -           break;
>           else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
>             {
> -             error ("%qD appears more than once in motion clauses", t);
> +             if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
> +               error ("%qD appears more than once in motion clauses", t);
> +             else
> +               error ("%qD appears more than once in map clauses", t);
>               remove = true;
>             }
>           else
> @@ -11509,6 +11506,7 @@ c_finish_omp_clauses (tree clauses)
>
>         case OMP_CLAUSE_IF:
>         case OMP_CLAUSE_NUM_THREADS:
> +       case OMP_CLAUSE_NUM_TEAMS:
>         case OMP_CLAUSE_THREAD_LIMIT:
>         case OMP_CLAUSE_SCHEDULE:
>         case OMP_CLAUSE_ORDERED:
> --- gcc/c/c-parser.c.jj 2013-08-27 20:49:04.000000000 +0200
> +++ gcc/c/c-parser.c    2013-09-02 11:41:51.023319550 +0200
> @@ -12068,7 +12068,10 @@ c_parser_omp_target_data (location_t loc
>    OMP_TARGET_DATA_CLAUSES (stmt)
>      = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
>                                 "#pragma omp target data");
> -  OMP_TARGET_DATA_BODY (stmt) = c_parser_omp_structured_block (parser);
> +  keep_next_level ();
> +  tree block = c_begin_compound_stmt (true);
> +  add_stmt (c_parser_omp_structured_block (parser));
> +  OMP_TARGET_DATA_BODY (stmt) = c_end_compound_stmt (loc, block, true);
>
>    SET_EXPR_LOCATION (stmt, loc);
>    return add_stmt (stmt);
> @@ -12161,6 +12164,7 @@ c_parser_omp_target (c_parser *parser, e
>
>           c_parser_consume_token (parser);
>           strcpy (p_name, "#pragma omp target");
> +         keep_next_level ();
>           tree block = c_begin_compound_stmt (true);
>           tree ret = c_parser_omp_teams (loc, parser, p_name,
>                                          OMP_TARGET_CLAUSE_MASK, cclauses);
> @@ -12182,7 +12186,10 @@ c_parser_omp_target (c_parser *parser, e
>    OMP_TARGET_CLAUSES (stmt)
>      = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
>                                 "#pragma omp target");
> -  OMP_TARGET_BODY (stmt) = c_parser_omp_structured_block (parser);
> +  keep_next_level ();
> +  tree block = c_begin_compound_stmt (true);
> +  add_stmt (c_parser_omp_structured_block (parser));
> +  OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
>
>    SET_EXPR_LOCATION (stmt, loc);
>    add_stmt (stmt);
> --- gcc/tree.h.jj       2013-08-27 22:05:43.000000000 +0200
> +++ gcc/tree.h  2013-09-03 14:35:36.350942556 +0200
> @@ -629,6 +629,9 @@ struct GTY(()) tree_base {
>         OMP_CLAUSE_LINEAR_NO_COPYIN in
>            OMP_CLAUSE_LINEAR
>
> +       OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION in
> +          OMP_CLAUSE_MAP
> +
>         TRANSACTION_EXPR_RELAXED in
>            TRANSACTION_EXPR
>
> @@ -2050,15 +2053,21 @@ enum omp_clause_map_kind
>    OMP_CLAUSE_MAP_TO,
>    OMP_CLAUSE_MAP_FROM,
>    OMP_CLAUSE_MAP_TOFROM,
> -  /* This following is an internal only map kind, used for pointer based array
> -     sections.  OMP_CLAUSE_SIZE for these is not the pointer size, which is
> -     implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias.  */
> +  /* The following kind is an internal only map kind, used for pointer based
> +     array sections.  OMP_CLAUSE_SIZE for these is not the pointer size,
> +     which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias.  */
>    OMP_CLAUSE_MAP_POINTER
>  };
>
>  #define OMP_CLAUSE_MAP_KIND(NODE) \
>    (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
>
> +/* Nonzero if this map clause is for array (rather than pointer) based array
> +   section with zero bias.  Both the non-decl OMP_CLAUSE_MAP and
> +   correspoidng OMP_CLAUSE_MAP_POINTER clause are marked with this flag.  */
> +#define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \
> +  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag)
> +
>  enum omp_clause_proc_bind_kind
>  {
>    /* Numbers should match omp_proc_bind_t enum in omp.h.  */
> --- gcc/gimple.def.jj   2013-08-19 12:07:51.000000000 +0200
> +++ gcc/gimple.def      2013-09-02 11:55:57.892684945 +0200
> @@ -349,12 +349,18 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "g
>     CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
>  DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE)
>
> -/* GIMPLE_OMP_TARGET <BODY, CLAUSES> represents
> +/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
>     #pragma omp target {,data,update}
>     BODY is the sequence of statements inside the target construct
>     (NULL for target update).
> -   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
> -DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE)
> +   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
> +   CHILD_FN is set when outlining the body of the target region.
> +   All the statements in BODY are moved into this newly created
> +   function when converting OMP constructs into low-GIMPLE.
> +   DATA_ARG is a vec of 3 local variables in the parent function
> +   containing data to be mapped to CHILD_FN.  This is used to
> +   implement the MAP clauses.  */
> +DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL)
>
>  /* GIMPLE_OMP_TEAMS <BODY, CLAUSES> represents #pragma omp teams
>     BODY is the sequence of statements inside the single section.
> --- gcc/fortran/types.def.jj    2013-08-19 12:07:51.000000000 +0200
> +++ gcc/fortran/types.def       2013-09-04 20:19:40.880566458 +0200
> @@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
>  DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
>                      BT_CONST_VOLATILE_PTR)
>  DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
> +DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
>
>  DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
>
> @@ -167,6 +168,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_LONG_LON
>                      BT_PTR_LONG, BT_PTR_LONG)
>  DEF_FUNCTION_TYPE_5 (BT_FN_VOID_SIZE_VPTR_PTR_PTR_INT, BT_VOID, BT_SIZE,
>                      BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT)
> +DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
> +                    BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
>
>  DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
>                       BT_BOOL, BT_LONG, BT_LONG, BT_LONG, BT_LONG,
> @@ -203,6 +206,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
>                      BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
>                      BT_ULONGLONG, BT_ULONGLONG,
>                      BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
> +DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
> +                    BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
> +                    BT_PTR, BT_PTR, BT_PTR)
>
>  DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
>                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
> --- gcc/cp/parser.c.jj  2013-08-27 20:50:58.000000000 +0200
> +++ gcc/cp/parser.c     2013-09-04 17:37:20.511076432 +0200
> @@ -29451,6 +29451,7 @@ cp_parser_omp_target_data (cp_parser *pa
>    OMP_TARGET_DATA_CLAUSES (stmt)
>      = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
>                                  "#pragma omp target data", pragma_tok);
> +  keep_next_level (true);
>    OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser);
>
>    SET_EXPR_LOCATION (stmt, pragma_tok->location);
> @@ -29543,6 +29544,7 @@ cp_parser_omp_target (cp_parser *parser,
>
>           cp_lexer_consume_token (parser->lexer);
>           strcpy (p_name, "#pragma omp target");
> +         keep_next_level (true);
>           tree sb = begin_omp_structured_block ();
>           unsigned save = cp_parser_begin_omp_structured_block (parser);
>           tree ret = cp_parser_omp_teams (parser, pragma_tok, p_name,
> @@ -29566,6 +29568,7 @@ cp_parser_omp_target (cp_parser *parser,
>    OMP_TARGET_CLAUSES (stmt)
>      = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
>                                  "#pragma omp target", pragma_tok);
> +  keep_next_level (true);
>    OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
>
>    SET_EXPR_LOCATION (stmt, pragma_tok->location);
> --- gcc/cp/cp-objcp-common.h.jj 2013-03-20 10:07:19.000000000 +0100
> +++ gcc/cp/cp-objcp-common.h    2013-09-04 19:33:40.520583302 +0200
> @@ -145,6 +145,8 @@ extern void cp_common_init_ts (void);
>  #define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
>  #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
>  #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
> +#undef LANG_HOOKS_OMP_MAPPABLE_TYPE
> +#define LANG_HOOKS_OMP_MAPPABLE_TYPE cp_omp_mappable_type
>
>  #undef LANG_HOOKS_EH_USE_CXA_END_CLEANUP
>  #define LANG_HOOKS_EH_USE_CXA_END_CLEANUP true
> --- gcc/cp/semantics.c.jj       2013-08-27 20:50:56.000000000 +0200
> +++ gcc/cp/semantics.c  2013-09-04 17:21:48.507946784 +0200
> @@ -4101,8 +4101,7 @@ cxx_omp_create_clause_info (tree c, tree
>
>  static tree
>  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
> -                            bool &maybe_zero_len, unsigned int &first_non_one,
> -                            bool &pointer_based_p)
> +                            bool &maybe_zero_len, unsigned int &first_non_one)
>  {
>    tree ret, low_bound, length, type;
>    if (TREE_CODE (t) != TREE_LIST)
> @@ -4134,16 +4133,11 @@ handle_omp_array_sections_1 (tree c, tre
>           return error_mark_node;
>         }
>        t = convert_from_reference (t);
> -      if (!processing_template_decl
> -         && POINTER_TYPE_P (TREE_TYPE (t))
> -         && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> -       pointer_based_p = true;
>        return t;
>      }
>
>    ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
> -                                    maybe_zero_len, first_non_one,
> -                                    pointer_based_p);
> +                                    maybe_zero_len, first_non_one);
>    if (ret == error_mark_node || ret == NULL_TREE)
>      return ret;
>
> @@ -4326,16 +4320,12 @@ handle_omp_array_sections_1 (tree c, tre
>      }
>    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
>      types.safe_push (TREE_TYPE (ret));
> -  /* For pointer based array sections we will need to evaluate lb more
> -     than once.  */
> -  if (pointer_based_p)
> +  /* We will need to evaluate lb more than once.  */
> +  tree lb = cp_save_expr (low_bound);
> +  if (lb != low_bound)
>      {
> -      tree lb = cp_save_expr (low_bound);
> -      if (lb != low_bound)
> -       {
> -         TREE_PURPOSE (t) = lb;
> -         low_bound = lb;
> -       }
> +      TREE_PURPOSE (t) = lb;
> +      low_bound = lb;
>      }
>    ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, false);
>    return ret;
> @@ -4347,12 +4337,10 @@ static bool
>  handle_omp_array_sections (tree c)
>  {
>    bool maybe_zero_len = false;
> -  bool pointer_based_p = false;
>    unsigned int first_non_one = 0;
>    vec<tree> types = vNULL;
>    tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
> -                                           maybe_zero_len, first_non_one,
> -                                           pointer_based_p);
> +                                           maybe_zero_len, first_non_one);
>    if (first == error_mark_node)
>      {
>        types.release ();
> @@ -4506,26 +4494,27 @@ handle_omp_array_sections (tree c)
>             size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
>           OMP_CLAUSE_DECL (c) = first;
>           OMP_CLAUSE_SIZE (c) = size;
> -         if (pointer_based_p)
> -           {
> -             tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
> -                                         OMP_CLAUSE_MAP);
> -             OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
> -             if (!cxx_mark_addressable (t))
> -               return false;
> -             OMP_CLAUSE_DECL (c2) = t;
> -             t = build_fold_addr_expr (first);
> -             t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> -                                   build_pointer_type (char_type_node), t);
> -             t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
> -                                  ptrdiff_type_node, t,
> -                                  fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> -                                                    TREE_TYPE (t),
> -                                                    OMP_CLAUSE_DECL (c2)));
> -             OMP_CLAUSE_SIZE (c2) = t;
> -             OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
> -             OMP_CLAUSE_CHAIN (c) = c2;
> -           }
> +         if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
> +           return false;
> +         tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
> +                                     OMP_CLAUSE_MAP);
> +         OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
> +         if (!cxx_mark_addressable (t))
> +           return false;
> +         OMP_CLAUSE_DECL (c2) = t;
> +         t = build_fold_addr_expr (first);
> +         t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> +                               ptrdiff_type_node, t);
> +         tree ptr = OMP_CLAUSE_DECL (c2);
> +         if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
> +           ptr = build_fold_addr_expr (ptr);
> +         t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
> +                              ptrdiff_type_node, t,
> +                              fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> +                                                ptrdiff_type_node, ptr));
> +         OMP_CLAUSE_SIZE (c2) = t;
> +         OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
> +         OMP_CLAUSE_CHAIN (c) = c2;
>         }
>      }
>    return false;
> @@ -4972,21 +4961,24 @@ finish_omp_clauses (tree clauses)
>                    && TREE_CODE (TREE_TYPE (t)) != REFERENCE_TYPE
>                    && !cxx_mark_addressable (t))
>             remove = true;
> -         else if (!cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
> -                                          == REFERENCE_TYPE)
> -                                         ? TREE_TYPE (TREE_TYPE (t))
> -                                         : TREE_TYPE (t)))
> +         else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +                    && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
> +                  && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
> +                                             == REFERENCE_TYPE)
> +                                            ? TREE_TYPE (TREE_TYPE (t))
> +                                            : TREE_TYPE (t)))
>             {
>               error_at (OMP_CLAUSE_LOCATION (c),
>                         "%qD does not have a mappable type in %qs clause", t,
>                         omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
>               remove = true;
>             }
> -         else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> -           break;
>           else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
>             {
> -             error ("%qD appears more than once in motion clauses", t);
> +             if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
> +               error ("%qD appears more than once in motion clauses", t);
> +             else
> +               error ("%qD appears more than once in map clauses", t);
>               remove = true;
>             }
>           else
> --- gcc/langhooks.h.jj  2013-03-20 10:08:13.000000000 +0100
> +++ gcc/langhooks.h     2013-09-04 19:29:20.429909391 +0200
> @@ -111,6 +111,9 @@ struct lang_hooks_for_types
>       firstprivate variables.  */
>    void (*omp_firstprivatize_type_sizes) (struct gimplify_omp_ctx *, tree);
>
> +  /* Return true if TYPE is a mappable type.  */
> +  bool (*omp_mappable_type) (tree type);
> +
>    /* Return TRUE if TYPE1 and TYPE2 are identical for type hashing purposes.
>       Called only after doing all language independent checks.
>       At present, this function is only called when both TYPE1 and TYPE2 are
> --- gcc/langhooks-def.h.jj      2013-03-20 10:08:41.000000000 +0100
> +++ gcc/langhooks-def.h 2013-09-04 19:29:45.450781052 +0200
> @@ -77,6 +77,7 @@ extern tree lhd_omp_assignment (tree, tr
>  struct gimplify_omp_ctx;
>  extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
>                                                tree);
> +extern bool lhd_omp_mappable_type (tree);
>
>  #define LANG_HOOKS_NAME                        "GNU unknown"
>  #define LANG_HOOKS_IDENTIFIER_SIZE     sizeof (struct lang_identifier)
> @@ -166,6 +167,7 @@ extern tree lhd_make_node (enum tree_cod
>  #define LANG_HOOKS_TYPE_MAX_SIZE       lhd_return_null_const_tree
>  #define LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES \
>    lhd_omp_firstprivatize_type_sizes
> +#define LANG_HOOKS_OMP_MAPPABLE_TYPE   lhd_omp_mappable_type
>  #define LANG_HOOKS_TYPE_HASH_EQ                NULL
>  #define LANG_HOOKS_GET_ARRAY_DESCR_INFO        NULL
>  #define LANG_HOOKS_GET_SUBRANGE_BOUNDS NULL
> @@ -184,6 +186,7 @@ extern tree lhd_make_node (enum tree_cod
>    LANG_HOOKS_INCOMPLETE_TYPE_ERROR, \
>    LANG_HOOKS_TYPE_MAX_SIZE, \
>    LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES, \
> +  LANG_HOOKS_OMP_MAPPABLE_TYPE, \
>    LANG_HOOKS_TYPE_HASH_EQ, \
>    LANG_HOOKS_GET_ARRAY_DESCR_INFO, \
>    LANG_HOOKS_GET_SUBRANGE_BOUNDS, \
> --- gcc/gimple.h.jj     2013-08-27 21:55:11.000000000 +0200
> +++ gcc/gimple.h        2013-09-02 11:57:32.143164024 +0200
> @@ -4702,7 +4702,7 @@ static inline tree
>  gimple_omp_target_clauses (const_gimple gs)
>  {
>    GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> -  return gs->gimple_omp_single.clauses;
> +  return gs->gimple_omp_parallel.clauses;
>  }
>
>
> @@ -4712,7 +4712,7 @@ static inline tree *
>  gimple_omp_target_clauses_ptr (gimple gs)
>  {
>    GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> -  return &gs->gimple_omp_single.clauses;
> +  return &gs->gimple_omp_parallel.clauses;
>  }
>
>
> @@ -4722,7 +4722,7 @@ static inline void
>  gimple_omp_target_set_clauses (gimple gs, tree clauses)
>  {
>    GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> -  gs->gimple_omp_single.clauses = clauses;
> +  gs->gimple_omp_parallel.clauses = clauses;
>  }
>
>
> @@ -4747,6 +4747,67 @@ gimple_omp_target_set_kind (gimple g, in
>  }
>
>
> +/* Return the child function used to hold the body of OMP_TARGET GS.  */
> +
> +static inline tree
> +gimple_omp_target_child_fn (const_gimple gs)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  return gs->gimple_omp_parallel.child_fn;
> +}
> +
> +/* Return a pointer to the child function used to hold the body of
> +   OMP_TARGET GS.  */
> +
> +static inline tree *
> +gimple_omp_target_child_fn_ptr (gimple gs)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  return &gs->gimple_omp_parallel.child_fn;
> +}
> +
> +
> +/* Set CHILD_FN to be the child function for OMP_TARGET GS.  */
> +
> +static inline void
> +gimple_omp_target_set_child_fn (gimple gs, tree child_fn)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  gs->gimple_omp_parallel.child_fn = child_fn;
> +}
> +
> +
> +/* Return the artificial argument used to send variables and values
> +   from the parent to the children threads in OMP_TARGET GS.  */
> +
> +static inline tree
> +gimple_omp_target_data_arg (const_gimple gs)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  return gs->gimple_omp_parallel.data_arg;
> +}
> +
> +
> +/* Return a pointer to the data argument for OMP_TARGET GS.  */
> +
> +static inline tree *
> +gimple_omp_target_data_arg_ptr (gimple gs)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  return &gs->gimple_omp_parallel.data_arg;
> +}
> +
> +
> +/* Set DATA_ARG to be the data argument for OMP_TARGET GS.  */
> +
> +static inline void
> +gimple_omp_target_set_data_arg (gimple gs, tree data_arg)
> +{
> +  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
> +  gs->gimple_omp_parallel.data_arg = data_arg;
> +}
> +
> +
>  /* Return the clauses associated with OMP_TEAMS GS.  */
>
>  static inline tree
> --- gcc/omp-builtins.def.jj     2013-08-19 12:07:51.000000000 +0200
> +++ gcc/omp-builtins.def        2013-09-04 15:14:21.049713033 +0200
> @@ -225,3 +225,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C
>                   BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
>  DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
>                   BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
> +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target",
> +                 BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
> +                 ATTR_NOTHROW_LIST)
> +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data",
> +                 BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
> +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
> +                 BT_FN_VOID, ATTR_NOTHROW_LIST)
> +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
> +                 BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
> +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
> +                 BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
> --- gcc/gimple-pretty-print.c.jj        2013-08-27 21:53:53.000000000 +0200
> +++ gcc/gimple-pretty-print.c   2013-08-30 17:53:14.201347484 +0200
> @@ -1277,6 +1277,13 @@ dump_gimple_omp_target (pretty_printer *
>        pp_string (buffer, "#pragma omp target");
>        pp_string (buffer, kind);
>        dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
> +      if (gimple_omp_target_child_fn (gs))
> +       {
> +         pp_string (buffer, " [child fn: ");
> +         dump_generic_node (buffer, gimple_omp_target_child_fn (gs),
> +                            spc, flags, false);
> +         pp_right_bracket (buffer);
> +       }
>        if (!gimple_seq_empty_p (gimple_omp_body (gs)))
>         {
>           newline_and_indent (buffer, spc + 2);
> --- gcc/gimplify.c.jj   2013-08-27 22:03:35.000000000 +0200
> +++ gcc/gimplify.c      2013-09-04 19:36:07.442834807 +0200
> @@ -5976,7 +5976,16 @@ omp_notice_variable (struct gimplify_omp
>    if (ctx->region_type == ORT_TARGET)
>      {
>        if (n == NULL)
> -       omp_add_variable (ctx, decl, GOVD_MAP | flags);
> +       {
> +         if (!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);
> +           }
> +         else
> +           omp_add_variable (ctx, decl, GOVD_MAP | flags);
> +       }
>        else
>         n->value |= flags;
>        ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
> @@ -6283,7 +6292,6 @@ gimplify_scan_omp_clauses (tree *list_p,
>               break;
>             }
>           flags = GOVD_MAP | GOVD_EXPLICIT;
> -         notice_outer = false;
>           goto do_add;
>
>         case OMP_CLAUSE_DEPEND:
> @@ -6500,31 +6508,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
>    if (private_debug)
>      code = OMP_CLAUSE_PRIVATE;
>    else if (flags & GOVD_MAP)
> -    {
> -      /* If decl is already in the enclosing device data environment,
> -        the spec says that it should just be used and no init/assignment
> -        should be done.  If there was any privatization in between though,
> -        it means that original decl might be in the enclosing device data
> -        environment, but the privatized might not.  */
> -      struct gimplify_omp_ctx *ctx;
> -      for (ctx = gimplify_omp_ctxp->outer_context;
> -          ctx; ctx = ctx->outer_context)
> -       {
> -         n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
> -         if (n == NULL)
> -           continue;
> -         if (ctx->region_type == ORT_TARGET_DATA)
> -           {
> -             if ((n->value & GOVD_MAP) != 0)
> -               return 0;
> -           }
> -         else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
> -                               | GOVD_PRIVATE | GOVD_REDUCTION
> -                               | GOVD_LINEAR)) != 0)
> -           break;
> -       }
> -      code = OMP_CLAUSE_MAP;
> -    }
> +    code = OMP_CLAUSE_MAP;
>    else if (flags & GOVD_SHARED)
>      {
>        if (is_global_var (decl))
> @@ -6689,31 +6673,6 @@ gimplify_adjust_omp_clauses (tree *list_
>           n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
>           if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
>             remove = true;
> -         else
> -           {
> -             /* If decl is already in the enclosing device data environment,
> -                the spec says that it should just be used and no init/assignment
> -                should be done.  If there was any privatization in between though,
> -                it means that original decl might be in the enclosing device data
> -                environment, but the privatized might not.  */
> -             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 == NULL)
> -                   continue;
> -                 if (octx->region_type == ORT_TARGET_DATA)
> -                   {
> -                     if ((n->value & GOVD_MAP) != 0)
> -                       remove = true;
> -                   }
> -                 else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
> -                                       | GOVD_PRIVATE | GOVD_REDUCTION
> -                                       | GOVD_LINEAR)) != 0)
> -                   break;
> -               }
> -           }
>           break;
>
>         case OMP_CLAUSE_REDUCTION:
> @@ -6722,6 +6681,7 @@ gimplify_adjust_omp_clauses (tree *list_
>         case OMP_CLAUSE_IF:
>         case OMP_CLAUSE_NUM_THREADS:
>         case OMP_CLAUSE_NUM_TEAMS:
> +       case OMP_CLAUSE_THREAD_LIMIT:
>         case OMP_CLAUSE_DIST_SCHEDULE:
>         case OMP_CLAUSE_DEVICE:
>         case OMP_CLAUSE_SCHEDULE:
> @@ -7198,7 +7158,28 @@ gimplify_omp_workshare (tree *expr_p, gi
>        gcc_unreachable ();
>      }
>    gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
> -  gimplify_and_add (OMP_BODY (expr), &body);
> +  if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
> +    {
> +      struct gimplify_ctx gctx;
> +      push_gimplify_context (&gctx);
> +      gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
> +      if (gimple_code (g) == GIMPLE_BIND)
> +       pop_gimplify_context (g);
> +      else
> +       pop_gimplify_context (NULL);
> +      if (ort == ORT_TARGET_DATA)
> +       {
> +         gimple_seq cleanup = NULL;
> +         tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
> +         g = gimple_build_call (fn, 0);
> +         gimple_seq_add_stmt (&cleanup, g);
> +         g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
> +         body = NULL;
> +         gimple_seq_add_stmt (&body, g);
> +       }
> +    }
> +  else
> +    gimplify_and_add (OMP_BODY (expr), &body);
>    gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
>
>    switch (TREE_CODE (expr))
> --- gcc/testsuite/c-c++-common/gomp/map-1.c.jj  2013-07-06 19:18:46.161478820 +0200
> +++ gcc/testsuite/c-c++-common/gomp/map-1.c     2013-09-04 17:00:51.174322082 +0200
> @@ -33,7 +33,7 @@ foo (int g[3][10], int h[4][8], int i[2]
>    #pragma omp target map(to: o[2:5]) /* { dg-error "does not have pointer or array type" } */
>      ;
>    #pragma omp target map(to: a[:][:]) /* { dg-error "array type length expression is not optional" } */
> -    bar (&a[0][0]);
> +    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" } */
>    #pragma omp target map(tofrom: b[-1:]) /* { dg-error "negative low bound in array section" } */
>      bar (b);
>    #pragma omp target map(tofrom: c[:-3][:]) /* { dg-error "negative length in array section" } */
> --- gcc/builtin-types.def.jj    2013-08-19 12:07:51.000000000 +0200
> +++ gcc/builtin-types.def       2013-09-03 08:58:48.902658867 +0200
> @@ -345,6 +345,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
>  DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
>                      BT_CONST_VOLATILE_PTR)
>  DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
> +DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
>
>  DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
>
> @@ -472,6 +473,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_VPTR_PTR
>  DEF_FUNCTION_TYPE_5 (BT_FN_VOID_OMPFN_PTR_UINT_UINT_UINT,
>                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT,
>                      BT_UINT)
> +DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
> +                    BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
>
>  DEF_FUNCTION_TYPE_6 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VALIST_ARG,
>                      BT_INT, BT_STRING, BT_SIZE, BT_INT, BT_SIZE,
> @@ -511,6 +514,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
>                      BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
>                      BT_ULONGLONG, BT_ULONGLONG,
>                      BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
> +DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
> +                    BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
> +                    BT_PTR, BT_PTR, BT_PTR)
>
>  DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
>                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
> --- libgomp/libgomp.map.jj      2013-06-12 14:24:42.000000000 +0200
> +++ libgomp/libgomp.map 2013-09-04 15:17:42.694659418 +0200
> @@ -218,4 +218,9 @@ GOMP_4.0 {
>         GOMP_parallel;
>         GOMP_taskgroup_start;
>         GOMP_taskgroup_end;
> +       GOMP_target;
> +       GOMP_target_data;
> +       GOMP_target_end_data;
> +       GOMP_target_update;
> +       GOMP_teams;
>  } GOMP_3.0;
> --- libgomp/Makefile.am.jj      2013-03-20 10:02:06.000000000 +0100
> +++ libgomp/Makefile.am 2013-09-04 14:59:19.475033836 +0200
> @@ -60,7 +60,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_L
>  libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
>         iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
>         task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
> -       time.c fortran.c affinity.c
> +       time.c fortran.c affinity.c target.c
>
>  nodist_noinst_HEADERS = libgomp_f.h
>  nodist_libsubinclude_HEADERS = omp.h
> --- libgomp/Makefile.in.jj      2013-03-20 10:02:05.000000000 +0100
> +++ libgomp/Makefile.in 2013-09-04 14:59:52.243864494 +0200
> @@ -96,7 +96,7 @@ am_libgomp_la_OBJECTS = alloc.lo barrier
>         error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \
>         parallel.lo sections.lo single.lo task.lo team.lo work.lo \
>         lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \
> -       fortran.lo affinity.lo
> +       fortran.lo affinity.lo target.lo
>  libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
>  DEFAULT_INCLUDES = -I.@am__isrc@
>  depcomp = $(SHELL) $(top_srcdir)/../depcomp
> @@ -317,7 +317,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_L
>  libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
>         iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
>         task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
> -       time.c fortran.c affinity.c
> +       time.c fortran.c affinity.c target.c
>
>  nodist_noinst_HEADERS = libgomp_f.h
>  nodist_libsubinclude_HEADERS = omp.h
> @@ -474,6 +474,7 @@ distclean-compile:
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sections.Plo@am__quote@
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@
> +@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@
>  @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/time.Plo@am__quote@
> --- libgomp/libgomp_g.h.jj      2013-04-05 18:15:57.000000000 +0200
> +++ libgomp/libgomp_g.h 2013-09-04 15:25:55.982080005 +0200
> @@ -199,4 +199,14 @@ extern bool GOMP_single_start (void);
>  extern void *GOMP_single_copy_start (void);
>  extern void GOMP_single_copy_end (void *);
>
> +/* target.c */
> +
> +extern void GOMP_target (int, void (*) (void *), const char *,
> +                        size_t, void **, size_t *, unsigned char *);
> +extern void GOMP_target_data (int, size_t, void **, size_t *, unsigned char *);
> +extern void GOMP_target_end_data (void);
> +extern void GOMP_target_update (int, size_t, void **, size_t *,
> +                               unsigned char *);
> +extern void GOMP_teams (unsigned int, unsigned int);
> +
>  #endif /* LIBGOMP_G_H */
> --- libgomp/testsuite/libgomp.c++/for-11.C.jj   2013-06-21 08:38:19.000000000 +0200
> +++ libgomp/testsuite/libgomp.c++/for-11.C      2013-09-04 15:28:21.542379101 +0200
> @@ -75,10 +75,7 @@ int
>  main ()
>  {
>    int err = 0;
> -// FIXME: distribute construct must be closely nested
> -// in teams region, but we don't handle target expansions
> -// yet.  Enable when it works.
> -// #pragma omp target teams reduction(|:err)
> +  #pragma omp target teams reduction(|:err)
>      {
>        err |= test_d_normal ();
>        err |= test_d_ds128_normal ();
> --- libgomp/testsuite/libgomp.c++/target-1.C.jj 2013-09-04 18:55:02.294387946 +0200
> +++ libgomp/testsuite/libgomp.c++/target-1.C    2013-09-04 18:54:56.430417616 +0200
> @@ -0,0 +1 @@
> +#include "../libgomp.c/target-1.c"
> --- libgomp/testsuite/libgomp.c/for-3.c.jj      2013-07-06 19:44:29.000000000 +0200
> +++ libgomp/testsuite/libgomp.c/for-3.c 2013-09-04 15:28:53.986223342 +0200
> @@ -77,10 +77,7 @@ int
>  main ()
>  {
>    int err = 0;
> -/* FIXME: distribute construct must be closely nested
> -   in teams region, but we don't handle target expansions
> -   yet.  Enable when it works.  */
> -/* #pragma omp target teams reduction(|:err) */
> +  #pragma omp target teams reduction(|:err)
>      {
>        err |= test_d_normal ();
>        err |= test_d_ds128_normal ();
> --- libgomp/testsuite/libgomp.c/target-1.c.jj   2013-09-04 18:54:18.520609855 +0200
> +++ libgomp/testsuite/libgomp.c/target-1.c      2013-09-04 18:54:10.000000000 +0200
> @@ -0,0 +1,90 @@
> +extern
> +#ifdef __cplusplus
> +"C"
> +#endif
> +void abort (void);
> +
> +void
> +fn1 (double *x, double *y, int z)
> +{
> +  int i;
> +  for (i = 0; i < z; i++)
> +    {
> +      x[i] = i & 31;
> +      y[i] = (i & 63) - 30;
> +    }
> +}
> +
> +#pragma omp declare target
> +int tgtv = 6;
> +int
> +tgt (void)
> +{
> +  #pragma omp atomic update
> +    tgtv++;
> +  return 0;
> +}
> +#pragma omp end declare target
> +
> +double
> +fn2 (int x, int y, int z)
> +{
> +  double b[1024], c[1024], s = 0;
> +  int i, j;
> +  fn1 (b, c, x);
> +  #pragma omp target data map(to: b)
> +  {
> +    #pragma omp target map(tofrom: c)
> +      #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
> +       #pragma omp distribute dist_schedule(static, 4) collapse(1)
> +         for (j=0; j < x; j += y)
> +           #pragma omp parallel for reduction(+:s)
> +             for (i = j; i < j + y; i++)
> +               tgt (), s += b[i] * c[i];
> +    #pragma omp target update from(b, tgtv)
> +  }
> +  return s;
> +}
> +
> +double
> +fn3 (int x)
> +{
> +  double b[1024], c[1024], s = 0;
> +  int i;
> +  fn1 (b, c, x);
> +  #pragma omp target map(to: b, c)
> +    #pragma omp parallel for reduction(+:s)
> +      for (i = 0; i < x; i++)
> +       tgt (), s += b[i] * c[i];
> +  return s;
> +}
> +
> +double
> +fn4 (int x, double *p)
> +{
> +  double b[1024], c[1024], d[1024], s = 0;
> +  int i;
> +  fn1 (b, c, x);
> +  fn1 (d + x, p + x, x);
> +  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
> +    #pragma omp parallel for reduction(+:s)
> +      for (i = 0; i < x; i++)
> +       s += b[i] * c[i] + d[x + i] + p[x + i];
> +  return s;
> +}
> +
> +int
> +main ()
> +{
> +  double a = fn2 (128, 4, 6);
> +  int b = tgtv;
> +  double c = fn3 (61);
> +  #pragma omp target update from(tgtv)
> +  int d = tgtv;
> +  double e[1024];
> +  double f = fn4 (64, e);
> +  if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
> +      || f != 8032.0)
> +    abort ();
> +  return 0;
> +}
> --- libgomp/target.c.jj 2013-09-04 14:58:30.488287438 +0200
> +++ libgomp/target.c    2013-09-04 15:18:31.363404929 +0200
> @@ -0,0 +1,83 @@
> +/* Copyright (C) 2013 Free Software Foundation, Inc.
> +   Contributed by Jakub Jelinek <jakub@redhat.com>.
> +
> +   This file is part of the GNU OpenMP Library (libgomp).
> +
> +   Libgomp is free software; you can redistribute it and/or modify it
> +   under the terms of the GNU General Public License as published by
> +   the Free Software Foundation; either version 3, or (at your option)
> +   any later version.
> +
> +   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
> +   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
> +   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
> +   more details.
> +
> +   Under Section 7 of GPL version 3, you are granted additional
> +   permissions described in the GCC Runtime Library Exception, version
> +   3.1, as published by the Free Software Foundation.
> +
> +   You should have received a copy of the GNU General Public License and
> +   a copy of the GCC Runtime Library Exception along with this program;
> +   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
> +   <http://www.gnu.org/licenses/>.  */
> +
> +/* This file handles the maintainence of threads in response to team
> +   creation and termination.  */
> +
> +#include "libgomp.h"
> +#include <stdlib.h>
> +#include <string.h>
> +
> +static int
> +resolve_device (int device)
> +{
> +  return -1;
> +}
> +
> +/* Called when encountering a target directive.  If DEVICE
> +   is -1, it means use device-var ICV.  If it is -2 (or any other value
> +   larger than last available hw device, use host fallback.
> +   FN is address of host code, FNNAME corresponding name to lookup
> +   in the target code.  HOSTADDRS, SIZES and KINDS are arrays
> +   with MAPNUM entries, with addresses of the host objects,
> +   sizes of the host objects (resp. for pointer kind pointer bias
> +   and assumed sizeof (void *) size) and kinds.  */
> +
> +void
> +GOMP_target (int device, void (*fn) (void *), const char *fnname,
> +            size_t mapnum, void **hostaddrs, size_t *sizes,
> +            unsigned char *kinds)
> +{
> +  if (resolve_device (device) == -1)
> +    {
> +      fn (hostaddrs);
> +      return;
> +    }
> +}
> +
> +void
> +GOMP_target_data (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
> +                 unsigned char *kinds)
> +{
> +  if (resolve_device (device) == -1)
> +    return;
> +}
> +
> +void
> +GOMP_target_end_data (void)
> +{
> +}
> +
> +void
> +GOMP_target_update (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
> +                   unsigned char *kinds)
> +{
> +  if (resolve_device (device) == -1)
> +    return;
> +}
> +
> +void
> +GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
> +{
> +}
>
>         Jakub