diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10005,8 +10005,10 @@
   llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd(
       MapType,
       MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
-                                   MappableExprsHandler::OMP_MAP_FROM |
-                                   MappableExprsHandler::OMP_MAP_MEMBER_OF)));
+                                   MappableExprsHandler::OMP_MAP_FROM)));
+  MapTypeArg = MapperCGF.Builder.CreateOr(
+      MapTypeArg,
+      MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_IMPLICIT));
 
   // Call the runtime API __tgt_push_mapper_component to fill up the runtime
   // data structure.
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -118,8 +118,11 @@
 // CK0: [[INIT]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK0: br label %[[LHEAD:[^,]+]]
 
 // CK0: [[LHEAD]]
@@ -228,8 +231,11 @@
 // CK0: [[EVALDEL]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK0: br label %[[DONE]]
 // CK0: [[DONE]]
 // CK0: ret void
@@ -672,8 +678,11 @@
 
 // CK1: [[INITEVALDEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK1: br label %[[LHEAD:[^,]+]]
 
 // CK1: [[LHEAD]]
@@ -718,8 +727,11 @@
 // CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 // CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK1: br label %[[DONE]]
 // CK1: [[DONE]]
 // CK1: ret void
@@ -793,8 +805,11 @@
 
 // CK2: [[INITEVALDEL]]
 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK2: br label %[[LHEAD:[^,]+]]
 
 // CK2: [[LHEAD]]
@@ -841,8 +856,11 @@
 // CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK2: [[EVALDEL]]
 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK2: br label %[[DONE]]
 // CK2: [[DONE]]
 // CK2: ret void
@@ -998,8 +1016,11 @@
 // CK4: [[INITEVALDEL]]
 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 // CK4: br label %[[LHEAD:[^,]+]]
 
 // CK4: [[LHEAD]]
@@ -1108,8 +1129,11 @@
 // CK4: [[EVALDEL]]
 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
+
+// Remove movement mappings and mark as implicit
+// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 // CK4: br label %[[DONE]]
 // CK4: [[DONE]]
 // CK4: ret void
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -472,7 +472,8 @@
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
-    bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF);
+    bool UpdateRef =
+        !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && i == 0);
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
       // Base is address of pointer.
@@ -665,8 +666,7 @@
     bool IsLast, IsHostPtr;
     bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
     bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
-                     (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ &&
-                      (!FromMapper || I != ArgNum - 1));
+                     (!FromMapper || I != ArgNum - 1);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
@@ -719,8 +719,7 @@
     // exists.
     if (((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
          !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) ||
-        (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && FromMapper &&
-         I == ArgNum - 1)) {
+        (FromMapper && I == ArgNum - 1)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
@@ -0,0 +1,70 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+// XFAIL: clang
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D sa[2];
+  double x[2], y[2];
+  double x1[2], y1[2];
+  y[1] = x[1] = 20;
+
+  sa[0].e = 111;
+  sa[0].f.a = 222;
+  sa[0].f.c.a = 777;
+  sa[0].f.b = &x[0];
+  sa[0].f.c.b = &x1[0];
+  sa[0].h = N;
+
+  sa[1].e = 111;
+  sa[1].f.a = 222;
+  sa[1].f.c.a = 777;
+  sa[1].f.b = &y[0];
+  sa[1].f.c.b = &y1[0];
+  sa[1].h = N;
+
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&y[0]);
+#pragma omp target map(tofrom : sa) firstprivate(p)
+  {
+    printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a,
+           sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sa[1].e = 333;
+    sa[1].f.a = 444;
+    sa[1].f.c.a = 555;
+    sa[1].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp
@@ -0,0 +1,60 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D sa[10];
+  sa[1].e = 111;
+  sa[1].f.a = 222;
+  sa[1].f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  sa[1].f.b = &x[0];
+  sa[1].f.c.b = &x1[0];
+  sa[1].h = N;
+
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+#pragma omp target map(tofrom : sa[1]) firstprivate(p)
+  {
+    printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a,
+           sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sa[1].e = 333;
+    sa[1].f.a = 444;
+    sa[1].f.c.a = 555;
+    sa[1].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
+         sa[1].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
@@ -0,0 +1,102 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 2
+
+class MyObjectA {
+public:
+  MyObjectA() {
+    data1 = 1;
+    data2 = 2;
+  }
+  void show() {
+    printf("\t\tObject A Contents:\n");
+    printf("\t\t\tdata1 = %d  data2 = %d\n", data1, data2);
+  }
+  void foo() {
+    data1 += 10;
+    data2 += 20;
+  }
+  int data1;
+  int data2;
+};
+
+class MyObjectB {
+public:
+  MyObjectB() {
+    arr = new MyObjectA[N];
+    len = N;
+  }
+  void show() {
+    printf("\tObject B Contents:\n");
+    for (int i = 0; i < len; i++)
+      arr[i].show();
+  }
+  void foo() {
+    for (int i = 0; i < len; i++)
+      arr[i].foo();
+  }
+  MyObjectA *arr;
+  int len;
+};
+#pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[:obj.len])
+
+class MyObjectC {
+public:
+  MyObjectC() {
+    arr = new MyObjectB[N];
+    len = N;
+  }
+  void show() {
+    printf("Object C Contents:\n");
+    for (int i = 0; i < len; i++)
+      arr[i].show();
+  }
+  void foo() {
+    for (int i = 0; i < len; i++)
+      arr[i].foo();
+  }
+  MyObjectB *arr;
+  int len;
+};
+#pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[:obj.len])
+
+int main(void) {
+  MyObjectC *outer = new MyObjectC[N];
+
+  printf("Original data hierarchy:\n");
+  for (int i = 0; i < N; i++)
+    outer[i].show();
+
+  printf("Sending data to device...\n");
+#pragma omp target enter data map(to : outer[:N])
+
+  printf("Calling foo()...\n");
+#pragma omp target teams distribute parallel for
+  for (int i = 0; i < N; i++)
+    outer[i].foo();
+
+  printf("foo() complete!\n");
+
+  printf("Sending data back to host...\n");
+#pragma omp target exit data map(from : outer[:N])
+
+  printf("Modified Data Hierarchy:\n");
+  for (int i = 0; i < N; i++)
+    outer[i].show();
+
+  printf("Testing for correctness...\n");
+  printf("outer[1].arr[1].arr[1].data2 = %d.\n", outer[1].arr[1].arr[1].data2);
+  // CHECK: outer[1].arr[1].arr[1].data2 = 22.
+  assert(outer[1].arr[1].arr[1].data2 == 22);
+
+  return 0;
+}
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp
@@ -0,0 +1,62 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D s;
+  s.e = 111;
+  s.f.a = 222;
+  s.f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  s.f.b = &x[0];
+  s.f.c.b = &x1[0];
+  s.h = N;
+
+  D *sp = &s;
+
+  printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1],
+         sp[0].f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+#pragma omp target map(tofrom : sp[0]) firstprivate(p)
+  {
+    printf("%d %d %d\n", sp[0].f.a, sp[0].f.c.a,
+           sp[0].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    sp[0].e = 333;
+    sp[0].f.a = 444;
+    sp[0].f.c.a = 555;
+    sp[0].f.b[1] = 40;
+  }
+  printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1],
+         sp[0].f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp
@@ -0,0 +1,62 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  int a;
+  double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+  int a;
+  double *b;
+  C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+  int e;
+  C f;
+  int h;
+} D;
+
+int main() {
+  constexpr int N = 10;
+  D s;
+  s.e = 111;
+  s.f.a = 222;
+  s.f.c.a = 777;
+  double x[2];
+  double x1[2];
+  x[1] = 20;
+  s.f.b = &x[0];
+  s.f.c.b = &x1[0];
+  s.h = N;
+
+  printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1],
+         s.f.b == &x[0] ? 1 : 0);
+  // CHECK: 111 222 777 20.00000 1
+
+  __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+
+#pragma omp target map(tofrom : s) firstprivate(p)
+  {
+    printf("%d %d %d\n", s.f.a, s.f.c.a,
+           s.f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+    // CHECK: 222 777 0
+    s.e = 333;
+    s.f.a = 444;
+    s.f.c.a = 555;
+    s.f.b[1] = 40;
+  }
+
+  printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1],
+         s.f.b == &x[0] ? 1 : 0);
+  // CHECK: 333 222 777 40.00000 1
+}