diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 068c0c69e58..6d30a9ed97d 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7923,6 +7923,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			      STRIP_NOPS (t);
 			      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
 				t = TREE_OPERAND (t, 0);
+			      if (REFERENCE_REF_P (t))
+				t = TREE_OPERAND (t, 0);
 			    }
 			}
 		      while (TREE_CODE (t) == COMPONENT_REF);
@@ -8021,6 +8023,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    {
 	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
 	      indir_component_ref_p = true;
+	      if (REFERENCE_REF_P (t))
+		t = TREE_OPERAND (t, 0);
 	      STRIP_NOPS (t);
 	      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
 		t = TREE_OPERAND (t, 0);
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 56f0aaaf979..8f07da8a991 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9802,7 +9802,10 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
   /* FIXME: If we're not mapping the base pointer in some other clause on this
      directive, I think we want to create ALLOC/RELEASE here -- i.e. not
      early-exit.  */
-  if (openmp && attach_detach)
+  if (openmp
+      && attach_detach
+      && !(TREE_CODE (TREE_TYPE (ocd)) == REFERENCE_TYPE
+	   && TREE_CODE (TREE_TYPE (TREE_TYPE (ocd))) != POINTER_TYPE))
     return NULL;
 
   if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
@@ -9851,9 +9854,32 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
 
       tree noind = strip_indirections (base);
 
-      if (!openmp
+      if (openmp
+	  && TREE_CODE (TREE_TYPE (noind)) == REFERENCE_TYPE
 	  && (region_type & ORT_TARGET)
 	  && TREE_CODE (noind) == COMPONENT_REF)
+	{
+	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_TO);
+	  OMP_CLAUSE_DECL (c2) = unshare_expr (base);
+	  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+
+	  tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
+	  OMP_CLAUSE_DECL (c3) = unshare_expr (noind);
+	  OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+	  OMP_CLAUSE_CHAIN (c2) = c3;
+	  OMP_CLAUSE_CHAIN (c3) = NULL_TREE;
+
+	  *inner = c2;
+	  return NULL;
+	}
+      else if (!openmp
+	       && (region_type & ORT_TARGET)
+	       && TREE_CODE (noind) == COMPONENT_REF)
 	{
 	  /* The base for this component access is a struct component access
 	     itself.  Insert a node to be processed on the next iteration of
@@ -9867,13 +9893,30 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
 	  OMP_CLAUSE_DECL (c2) = unshare_expr (noind);
 	  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+	  OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
 	  *inner = c2;
 	  return NULL;
 	}
 
-      tree sdecl = strip_components_and_deref (base);
+      tree sdecl = base;
+      /* There are too many places we need to do things like this.  */
+      if (TREE_CODE (sdecl) == INDIRECT_REF
+	  || TREE_CODE (sdecl) == MEM_REF)
+	{
+	  sdecl = TREE_OPERAND (sdecl, 0);
+	  if (TREE_CODE (sdecl) == INDIRECT_REF
+	      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (sdecl, 0)))
+		  == REFERENCE_TYPE))
+	    sdecl = TREE_OPERAND (sdecl, 0);
+	}
 
-      if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET))
+      while (TREE_CODE (sdecl) == COMPONENT_REF
+	     || TREE_CODE (sdecl) == POINTER_PLUS_EXPR)
+	sdecl = TREE_OPERAND (sdecl, 0);
+
+      if (DECL_P (sdecl)
+	  && POINTER_TYPE_P (TREE_TYPE (sdecl))
+	  && (region_type & ORT_TARGET))
 	{
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
 				      OMP_CLAUSE_MAP);
@@ -10212,11 +10255,10 @@ omp_build_struct_sibling_lists (enum tree_code code,
 	      else
 		*tail = inner;
 
-	      OMP_CLAUSE_CHAIN (inner) = NULL_TREE;
-
 	      omp_mapping_group newgrp;
 	      newgrp.grp_start = new_next ? new_next : tail;
-	      newgrp.grp_end = inner;
+	      newgrp.grp_end = (OMP_CLAUSE_CHAIN (inner)
+				? OMP_CLAUSE_CHAIN (inner) : inner);
 	      newgrp.mark = UNVISITED;
 	      newgrp.sibling = NULL;
 	      newgrp.next = NULL;
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
index cabeb7c2b7a..39a48a40920 100644
--- a/libgomp/testsuite/libgomp.c++/baseptrs-3.C
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
@@ -1,9 +1,58 @@
-/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */
-
 #include <cstdlib>
 #include <cstring>
 #include <cassert>
 
+struct sa0
+{
+  int *ptr;
+};
+
+struct sb0
+{
+  int arr[10];
+};
+
+struct sc0
+{
+  sa0 a;
+  sb0 b;
+  sc0 (sa0 &my_a, sb0 &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo0 ()
+{
+  sa0 my_a;
+  sb0 my_b;
+
+  my_a.ptr = (int *) malloc (sizeof (int) * 10);
+  sc0 my_c(my_a, my_b);
+
+  memset (my_c.a.ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.a.ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.a.ptr[i] == i);
+
+  memset (my_c.b.arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.b.arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.b.arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.b.arr[i] == i);
+
+  free (my_a.ptr);
+}
+
 struct sa
 {
   int *ptr;
@@ -90,6 +139,49 @@ bar ()
   free (my_a.ptr);
 }
 
+struct scp0
+{
+  sa *a;
+  sb *b;
+  scp0 (sa *my_a, sb *my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop0 ()
+{
+  sa *my_a = new sa;
+  sb *my_b = new sb;
+
+  my_a->ptr = new int[10];
+  scp0 *my_c = new scp0(my_a, my_b);
+
+  memset (my_c->a->ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->a->ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->a->ptr[i] == i);
+
+  memset (my_c->b->arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->b->arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->b->arr[i] == i);
+
+  delete[] my_a->ptr;
+  delete my_a;
+  delete my_b;
+}
+
 struct scp
 {
   sa *&a;
@@ -108,7 +200,7 @@ foop ()
 
   memset (my_c->a->ptr, 0, sizeof (int) * 10);
 
-  #pragma omp target map (my_c->a->ptr, my_c->a->ptr[:10])
+  #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
   {
     for (int i = 0; i < 10; i++)
       my_c->a->ptr[i] = i;
@@ -119,8 +211,7 @@ foop ()
 
   memset (my_c->b->arr, 0, sizeof (int) * 10);
 
-/* FIXME: This currently ICEs.  */
-/*  #pragma omp target map (my_c->b->arr[:10]) */
+  #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
   {
     for (int i = 0; i < 10; i++)
       my_c->b->arr[i] = i;
@@ -146,7 +237,8 @@ barp ()
 
   memset (my_cref->a->ptr, 0, sizeof (int) * 10);
 
-  #pragma omp target map (my_cref->a->ptr, my_cref->a->ptr[:10])
+  #pragma omp target map (my_cref->a, my_cref->a[:1], my_cref->a->ptr, \
+			  my_cref->a->ptr[:10])
   {
     for (int i = 0; i < 10; i++)
       my_cref->a->ptr[i] = i;
@@ -157,8 +249,7 @@ barp ()
 
   memset (my_cref->b->arr, 0, sizeof (int) * 10);
 
-/* FIXME: This currently ICEs.  */
-/*  #pragma omp target map (my_cref->b->arr[:10]) */
+  #pragma omp target map (my_cref->b, my_cref->b[:1], my_cref->b->arr[:10])
   {
     for (int i = 0; i < 10; i++)
       my_cref->b->arr[i] = i;
@@ -174,8 +265,10 @@ barp ()
 
 int main (int argc, char *argv[])
 {
+  foo0 ();
   foo ();
   bar ();
+  foop0 ();
   foop ();
   barp ();
   return 0;
