--- gcc/gimplify.c.jj	2015-07-31 16:55:01.482411392 +0200
+++ gcc/gimplify.c	2015-07-31 16:57:22.307320290 +0200
@@ -6202,6 +6202,7 @@ gimplify_scan_omp_clauses (tree *list_p,
 {
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
+  hash_map<tree, tree> *struct_map_to_clause = NULL;
 
   ctx = new_omp_context (region_type);
   outer_ctx = ctx->outer_context;
@@ -6442,6 +6443,11 @@ gimplify_scan_omp_clauses (tree *list_p,
 	    }
 	  if (!DECL_P (decl))
 	    {
+	      if (TREE_CODE (decl) == COMPONENT_REF)
+		{
+		  while (TREE_CODE (decl) == COMPONENT_REF)
+		    decl = TREE_OPERAND (decl, 0);
+		}
 	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
 				 NULL, is_gimple_lvalue, fb_lvalue)
 		  == GS_ERROR)
@@ -6449,6 +6455,128 @@ gimplify_scan_omp_clauses (tree *list_p,
 		  remove = true;
 		  break;
 		}
+	      if (DECL_P (decl))
+		{
+		  if (error_operand_p (decl))
+		    {
+		      remove = true;
+		      break;
+		    }
+
+		  if (TYPE_SIZE_UNIT (TREE_TYPE (decl)) == NULL
+		      || (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (decl)))
+			  != INTEGER_CST))
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"mapping field %qE of variable length "
+				"structure", OMP_CLAUSE_DECL (c));
+		      remove = true;
+		      break;
+		    }
+
+		  tree offset;
+		  HOST_WIDE_INT bitsize, bitpos;
+		  machine_mode mode;
+		  int unsignedp, volatilep = 0;
+		  tree base
+		    = get_inner_reference (OMP_CLAUSE_DECL (c), &bitsize,
+					   &bitpos, &offset, &mode, &unsignedp,
+					   &volatilep, false);
+		  gcc_assert (base == decl
+			      && (offset == NULL_TREE
+				  || TREE_CODE (offset) == INTEGER_CST));
+
+		  splay_tree_node n
+		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+		  if (n == NULL || (n->value & GOVD_MAP) == 0)
+		    {
+		      *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
+		      OMP_CLAUSE_DECL (*list_p) = decl;
+		      OMP_CLAUSE_SIZE (*list_p) = size_int (1);
+		      OMP_CLAUSE_CHAIN (*list_p) = c;
+		      if (struct_map_to_clause == NULL)
+			struct_map_to_clause = new hash_map<tree, tree>;
+		      struct_map_to_clause->put (decl, *list_p);
+		      list_p = &OMP_CLAUSE_CHAIN (*list_p);
+		      flags = GOVD_MAP | GOVD_EXPLICIT;
+		      if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+			flags |= GOVD_SEEN;
+		      goto do_add_decl;
+		    }
+		  else
+		    {
+		      tree *osc = struct_map_to_clause->get (decl), *sc;
+		      if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+			n->value |= GOVD_SEEN;
+		      offset_int o1, o2;
+		      if (offset)
+			o1 = wi::to_offset (offset);
+		      else
+			o1 = 0;
+		      if (bitpos)
+			o1 = o1 + bitpos / BITS_PER_UNIT;
+		      for (sc = &OMP_CLAUSE_CHAIN (*osc); *sc != c;
+			   sc = &OMP_CLAUSE_CHAIN (*sc))
+			if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF)
+			  break;
+			else
+			  {
+			    tree offset2;
+			    HOST_WIDE_INT bitsize2, bitpos2;
+			    base = get_inner_reference (OMP_CLAUSE_DECL (*sc),
+							&bitsize2, &bitpos2,
+							&offset2, &mode,
+							&unsignedp, &volatilep,
+							false);
+			    if (base != decl)
+			      break;
+			    gcc_assert (offset == NULL_TREE
+					|| TREE_CODE (offset) == INTEGER_CST);
+			    tree d1 = OMP_CLAUSE_DECL (*sc);
+			    tree d2 = OMP_CLAUSE_DECL (c);
+			    while (TREE_CODE (d1) == COMPONENT_REF)
+			      if (TREE_CODE (d2) == COMPONENT_REF
+				  && TREE_OPERAND (d1, 1)
+				     == TREE_OPERAND (d2, 1))
+				{
+				  d1 = TREE_OPERAND (d1, 0);
+				  d2 = TREE_OPERAND (d2, 0);
+				}
+			      else
+				break;
+			    if (d1 == d2)
+			      {
+				error_at (OMP_CLAUSE_LOCATION (c),
+					  "%qE appears more than once in map "
+					  "clauses", OMP_CLAUSE_DECL (c));
+				remove = true;
+				break;
+			      }
+			    if (offset2)
+			      o2 = wi::to_offset (offset2);
+			    else
+			      o2 = 0;
+			    if (bitpos2)
+			      o2 = o2 + bitpos2 / BITS_PER_UNIT;
+			    if (wi::ltu_p (o1, o2)
+				|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
+			      break;
+			  }
+		      if (!remove)
+			OMP_CLAUSE_SIZE (*osc)
+			  = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+					size_one_node);
+		      if (!remove && *sc != c)
+			{
+			  *list_p = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = *sc;
+			  *sc = c;
+			  continue;
+			}
+		    }
+		}
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -6790,6 +6918,8 @@ gimplify_scan_omp_clauses (tree *list_p,
     }
 
   gimplify_omp_ctxp = ctx;
+  if (struct_map_to_clause)
+    delete struct_map_to_clause;
 }
 
 struct gimplify_adjust_omp_clauses_data
--- gcc/omp-low.c.jj	2015-07-31 16:55:01.272414510 +0200
+++ gcc/omp-low.c	2015-07-31 16:57:22.317320141 +0200
@@ -12954,6 +12954,7 @@ lower_omp_target (gimple_stmt_iterator *
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	  case GOMP_MAP_STRUCT:
 	    break;
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
@@ -13303,6 +13304,7 @@ lower_omp_target (gimple_stmt_iterator *
 		    case GOMP_MAP_ALWAYS_TO:
 		    case GOMP_MAP_ALWAYS_FROM:
 		    case GOMP_MAP_ALWAYS_TOFROM:
+		    case GOMP_MAP_RELEASE:
 		      tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
 		      break;
 		    default:
--- gcc/tree-pretty-print.c.jj	2015-07-31 16:55:01.484411362 +0200
+++ gcc/tree-pretty-print.c	2015-07-31 16:57:22.320320097 +0200
@@ -643,6 +643,9 @@ dump_omp_clause (pretty_printer *pp, tre
 	case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  pp_string (pp, "firstprivate");
 	  break;
+	case GOMP_MAP_STRUCT:
+	  pp_string (pp, "struct");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
--- gcc/c/c-parser.c.jj	2015-07-31 16:55:01.481411407 +0200
+++ gcc/c/c-parser.c	2015-07-31 16:57:22.313320201 +0200
@@ -10190,10 +10190,25 @@ c_parser_omp_variable_list (c_parser *pa
 		  t = error_mark_node;
 		  break;
 		}
-	      /* FALL THROUGH.  */
+	      /* FALLTHROUGH  */
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
+	      while (c_parser_next_token_is (parser, CPP_DOT))
+		{
+		  location_t op_loc = c_parser_peek_token (parser)->location;
+		  c_parser_consume_token (parser);
+		  if (!c_parser_next_token_is (parser, CPP_NAME))
+		    {
+		      c_parser_error (parser, "expected identifier");
+		      t = error_mark_node;
+		      break;
+		    }
+		  tree ident = c_parser_peek_token (parser)->value;
+		  c_parser_consume_token (parser);
+		  t = build_component_ref (op_loc, t, ident);
+		}
+	      /* FALLTHROUGH  */
 	    case OMP_CLAUSE_DEPEND:
 	    case OMP_CLAUSE_REDUCTION:
 	      while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
--- gcc/c/c-typeck.c.jj	2015-07-31 16:55:01.482411392 +0200
+++ gcc/c/c-typeck.c	2015-07-31 16:58:09.246623290 +0200
@@ -12040,6 +12040,7 @@ handle_omp_array_sections (tree c, bool
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_RELEASE:
 	    OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 	    break;
 	  default:
@@ -12117,7 +12118,7 @@ tree
 c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
-  bitmap_head aligned_head, map_head;
+  bitmap_head aligned_head, map_head, map_field_head;
   tree c, t, type, *pc;
   tree simdlen = NULL_TREE, safelen = NULL_TREE;
   bool branch_seen = false;
@@ -12130,6 +12131,7 @@ c_finish_omp_clauses (tree clauses, bool
   bitmap_initialize (&lastprivate_head, &bitmap_default_obstack);
   bitmap_initialize (&aligned_head, &bitmap_default_obstack);
   bitmap_initialize (&map_head, &bitmap_default_obstack);
+  bitmap_initialize (&map_field_head, &bitmap_default_obstack);
 
   for (pc = &clauses, c = clauses; c ; c = *pc)
     {
@@ -12574,8 +12576,49 @@ c_finish_omp_clauses (tree clauses, bool
 	      break;
 	    }
 	  if (t == error_mark_node)
-	    remove = true;
-	  else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+	    {
+	      remove = true;
+	      break;
+	    }
+	  if (TREE_CODE (t) == COMPONENT_REF
+	      && is_omp
+	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+	    {
+	      if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "bit-field %qE in %qs clause",
+			    t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  remove = true;
+		}
+	      else if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%qE does not have a mappable type in %qs clause",
+			    t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  remove = true;
+		}
+	      while (TREE_CODE (t) == COMPONENT_REF)
+		{
+		  if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
+		      == UNION_TYPE)
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"%qE is a member of a union", t);
+		      remove = true;
+		      break;
+		    }
+		  t = TREE_OPERAND (t, 0);
+		}
+	      if (remove)
+		break;
+	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+		{
+		  if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+		    break;
+		}
+	    }
+	  if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qE is not a variable in %qs clause", t,
@@ -12597,6 +12640,7 @@ c_finish_omp_clauses (tree clauses, bool
 			     == GOMP_MAP_FIRSTPRIVATE_POINTER)
 			 || (OMP_CLAUSE_MAP_KIND (c)
 			     == GOMP_MAP_FORCE_DEVICEPTR)))
+		   && t == OMP_CLAUSE_DECL (c)
 		   && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
@@ -12613,7 +12657,12 @@ c_finish_omp_clauses (tree clauses, bool
 	      remove = true;
 	    }
 	  else
-	    bitmap_set_bit (&map_head, DECL_UID (t));
+	    {
+	      bitmap_set_bit (&map_head, DECL_UID (t));
+	      if (t != OMP_CLAUSE_DECL (c)
+		  && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+		bitmap_set_bit (&map_field_head, DECL_UID (t));
+	    }
 	  break;
 
 	case OMP_CLAUSE_TO_DECLARE:
--- gcc/cp/semantics.c.jj	2015-07-31 16:55:01.485411348 +0200
+++ gcc/cp/semantics.c	2015-07-31 16:57:22.303320349 +0200
@@ -4836,6 +4836,7 @@ handle_omp_array_sections (tree c, bool
 	      case GOMP_MAP_ALWAYS_TO:
 	      case GOMP_MAP_ALWAYS_FROM:
 	      case GOMP_MAP_ALWAYS_TOFROM:
+	      case GOMP_MAP_RELEASE:
 		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 		break;
 	      default:
--- include/gomp-constants.h.jj	2015-07-31 16:55:01.604409581 +0200
+++ include/gomp-constants.h	2015-07-31 16:55:38.711858574 +0200
@@ -102,6 +102,14 @@ enum gomp_map_kind
     /* If not already present, allocate.  And unconditionally copy to and from
        device.  */
     GOMP_MAP_ALWAYS_TOFROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
+    /* Map a sparse struct; the address is the base of the structure, alignment
+       it's required alignment, and size is the number of adjacent entries
+       that belong to the struct.  The adjacent entries should be sorted by
+       increasing address, so it is easy to determine lowest needed address
+       (address of the first adjacent entry) and highest needed address
+       (address of the last adjacent entry plus its size).  */
+    GOMP_MAP_STRUCT =			(GOMP_MAP_FLAG_ALWAYS
+					 | GOMP_MAP_FLAG_SPECIAL | 0),
     /* OpenMP 4.1 alias for forced deallocation.  */
     GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,
     /* Decrement usage count and deallocate if zero.  */
--- libgomp/target.c.jj	2015-07-31 16:55:01.981403983 +0200
+++ libgomp/target.c	2015-07-31 16:55:38.710858589 +0200
@@ -245,6 +245,66 @@ gomp_map_pointer (struct target_mem_desc
 			  sizeof (void *));
 }
 
+static void
+gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n,
+			  size_t first, size_t i, void **hostaddrs,
+			  size_t *sizes, void *kinds)
+{
+  struct gomp_device_descr *devicep = tgt->device_descr;
+  struct splay_tree_s *mem_map = &devicep->mem_map;
+  struct splay_tree_key_s cur_node;
+  int kind;
+  const bool short_mapkind = true;
+  const int typemask = short_mapkind ? 0xff : 0x7;
+
+  cur_node.host_start = (uintptr_t) hostaddrs[i];
+  cur_node.host_end = cur_node.host_start + sizes[i];
+  splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
+  kind = get_kind (short_mapkind, kinds, i);
+  if (n2
+      && n2->tgt == n->tgt
+      && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
+    {
+      gomp_map_vars_existing (devicep, n2, &cur_node,
+			      &tgt->list[i], kind & typemask);
+      return;
+    }
+  if (sizes[i] == 0)
+    {
+      if (cur_node.host_start > (uintptr_t) hostaddrs[first - 1])
+	{
+	  cur_node.host_start--;
+	  n2 = splay_tree_lookup (mem_map, &cur_node);
+	  cur_node.host_start++;
+	  if (n2
+	      && n2->tgt == n->tgt
+	      && n2->host_start - n->host_start
+		 == n2->tgt_offset - n->tgt_offset)
+	    {
+	      gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
+				      kind & typemask);
+	      return;
+	    }
+	}
+      cur_node.host_end++;
+      n2 = splay_tree_lookup (mem_map, &cur_node);
+      cur_node.host_end--;
+      if (n2
+	  && n2->tgt == n->tgt
+	  && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
+	{
+	  gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
+				  kind & typemask);
+	  return;
+	}
+    }
+  gomp_mutex_unlock (&devicep->lock);
+  gomp_fatal ("Trying to map into device [%p..%p) structure element when "
+	      "other mapped elements from the same structure weren't mapped "
+	      "together with it", (void *) cur_node.host_start,
+	      (void *) cur_node.host_end);
+}
+
 attribute_hidden struct target_mem_desc *
 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -304,6 +364,37 @@ gomp_map_vars (struct gomp_device_descr
 	  tgt->list[i].offset = ~(uintptr_t) 0;
 	  continue;
 	}
+      else if ((kind & typemask) == GOMP_MAP_STRUCT)
+	{
+	  size_t first = i + 1;
+	  size_t last = i + sizes[i];
+	  cur_node.host_start = (uintptr_t) hostaddrs[i];
+	  cur_node.host_end = (uintptr_t) hostaddrs[last]
+			      + sizes[last];
+	  tgt->list[i].key = NULL;
+	  tgt->list[i].offset = ~(uintptr_t) 2;
+	  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+	  if (n == NULL)
+	    {
+	      size_t align = (size_t) 1 << (kind >> rshift);
+	      if (tgt_align < align)
+		tgt_align = align;
+	      tgt_size -= (uintptr_t) hostaddrs[first]
+			  - (uintptr_t) hostaddrs[i];
+	      tgt_size = (tgt_size + align - 1) & ~(align - 1);
+	      tgt_size += cur_node.host_end - (uintptr_t) hostaddrs[i];
+	      not_found_cnt += last - i;
+	      for (i = first; i <= last; i++)
+		tgt->list[i].key = NULL;
+	      i--;
+	      continue;
+	    }
+	  for (i = first; i <= last; i++)
+	    gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
+				      sizes, kinds);
+	  i--;
+	  continue;
+	}
       cur_node.host_start = (uintptr_t) hostaddrs[i];
       if (!GOMP_MAP_POINTER_P (kind & typemask))
 	cur_node.host_end = cur_node.host_start + sizes[i];
@@ -406,7 +497,8 @@ gomp_map_vars (struct gomp_device_descr
       if (not_found_cnt)
 	tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
       splay_tree_node array = tgt->array;
-      size_t j;
+      size_t j, field_tgt_offset = 0, field_tgt_clear = ~(size_t) 0;
+      uintptr_t field_tgt_base = 0;
 
       for (i = 0; i < mapnum; i++)
 	if (tgt->list[i].key == NULL)
@@ -414,24 +506,53 @@ gomp_map_vars (struct gomp_device_descr
 	    int kind = get_kind (short_mapkind, kinds, i);
 	    if (hostaddrs[i] == NULL)
 	      continue;
-	    if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE)
+	    switch (kind & typemask)
 	      {
-		size_t align = (size_t) 1 << (kind >> rshift);
+		size_t align, len, first, last;
+		splay_tree_key n;
+	      case GOMP_MAP_FIRSTPRIVATE:
+		align = (size_t) 1 << (kind >> rshift);
 		tgt_size = (tgt_size + align - 1) & ~(align - 1);
 		tgt->list[i].offset = tgt_size;
-		size_t len = sizes[i];
+		len = sizes[i];
 		devicep->host2dev_func (devicep->target_id,
 					(void *) (tgt->tgt_start + tgt_size),
 					(void *) hostaddrs[i], len);
 		tgt_size += len;
 		continue;
-	      }
-	    switch (kind & typemask)
-	      {
 	      case GOMP_MAP_FIRSTPRIVATE_INT:
 	      case GOMP_MAP_USE_DEVICE_PTR:
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
 		continue;
+	      case GOMP_MAP_STRUCT:
+		first = i + 1;
+		last = i + sizes[i];
+		cur_node.host_start = (uintptr_t) hostaddrs[i];
+		cur_node.host_end = (uintptr_t) hostaddrs[last]
+				    + sizes[last];
+		if (tgt->list[first].key != NULL)
+		  continue;
+		n = splay_tree_lookup (mem_map, &cur_node);
+		if (n == NULL)
+		  {
+		    size_t align = (size_t) 1 << (kind >> rshift);
+		    tgt_size -= (uintptr_t) hostaddrs[first]
+				- (uintptr_t) hostaddrs[i];
+		    tgt_size = (tgt_size + align - 1) & ~(align - 1);
+		    tgt_size += (uintptr_t) hostaddrs[first]
+				- (uintptr_t) hostaddrs[i];
+		    field_tgt_base = (uintptr_t) hostaddrs[first];
+		    field_tgt_offset = tgt_size;
+		    field_tgt_clear = last;
+		    tgt_size += cur_node.host_end
+				- (uintptr_t) hostaddrs[first];
+		    continue;
+		  }
+		for (i = first; i <= last; i++)
+		  gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
+					    sizes, kinds);
+		i--;
+		continue;
 	      default:
 		break;
 	      }
@@ -449,10 +570,20 @@ gomp_map_vars (struct gomp_device_descr
 	      {
 		size_t align = (size_t) 1 << (kind >> rshift);
 		tgt->list[i].key = k;
-		tgt_size = (tgt_size + align - 1) & ~(align - 1);
 		k->tgt = tgt;
-		k->tgt_offset = tgt_size;
-		tgt_size += k->host_end - k->host_start;
+		if (field_tgt_clear != ~(size_t) 0)
+		  {
+		    k->tgt_offset = k->host_start - field_tgt_base
+				    + field_tgt_offset;
+		    if (i == field_tgt_clear)
+		      field_tgt_clear = ~(size_t) 0;
+		  }
+		else
+		  {
+		    tgt_size = (tgt_size + align - 1) & ~(align - 1);
+		    k->tgt_offset = tgt_size;
+		    tgt_size += k->host_end - k->host_start;
+		  }
 		tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
 		tgt->list[i].always_copy_from
 		  = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
@@ -571,6 +702,12 @@ gomp_map_vars (struct gomp_device_descr
 		cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
 	      else if (tgt->list[i].offset == ~(uintptr_t) 1)
 		cur_node.tgt_offset = 0;
+	      else if (tgt->list[i].offset == ~(uintptr_t) 2)
+		cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start
+				      + tgt->list[i + 1].key->tgt_offset
+				      + tgt->list[i + 1].offset
+				      + (uintptr_t) hostaddrs[i]
+				      - (uintptr_t) hostaddrs[i + 1];
 	      else
 		cur_node.tgt_offset = tgt->tgt_start
 				      + tgt->list[i].offset;
--- libgomp/testsuite/libgomp.c/target-21.c.jj	2015-07-31 17:00:30.415527080 +0200
+++ libgomp/testsuite/libgomp.c/target-21.c	2015-07-31 17:32:56.098638516 +0200
@@ -0,0 +1,55 @@
+extern void abort (void);
+union U { int x; long long y; };
+struct T { int a; union U b; int c; };
+struct S { int s; int u; struct T v; union U w; };
+
+int
+main ()
+{
+  struct S s;
+  s.s = 0;
+  s.u = 1;
+  s.v.a = 2;
+  s.v.b.y = 3LL;
+  s.v.c = 19;
+  s.w.x = 4;
+  int err = 0;
+  #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err)
+  {
+    err = 0;
+    if (s.u != 1 || s.v.b.y != 3LL)
+      err = 1;
+    s.w.x = 6;
+  }
+  if (err || s.w.x != 6)
+    abort ();
+  s.u++;
+  s.v.a++;
+  s.v.b.y++;
+  s.w.x++;
+  #pragma omp target data map (tofrom: s)
+  #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+  {
+    err = 0;
+    if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7)
+      err = 1;
+    s.w.x = 8;
+  }
+  if (err || s.w.x != 8)
+    abort ();
+  s.u++;
+  s.v.a++;
+  s.v.b.y++;
+  s.w.x++;
+  #pragma omp target data map (from: s.w) map (to: s.v.b, s.u)
+  #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+  {
+    err = 0;
+    if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9)
+      err = 1;
+    s.w.x = 11;
+  }
+  if (err || s.w.x != 11)
+    abort ();
+  return 0;
+}
