diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
new file mode 100644
index 0000000000000..3b1a8192bf2cf
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+  int *paa[10][10];
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    int *original_ph3 = &ph[3];
+    int **original_paa02 = &paa[0][2];
+
+#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
+    int *mapped_ptr_ph3 =
+        (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+    int **mapped_ptr_paa02 =
+        (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+
+    // CHECK-COUNT-4: 1
+    printf("%d\n", mapped_ptr_ph3 != nullptr);
+    printf("%d\n", mapped_ptr_paa02 != nullptr);
+    printf("%d\n", original_ph3 != mapped_ptr_ph3);
+    printf("%d\n", original_paa02 != mapped_ptr_paa02);
+
+// (A) use_device_addr operand within mapped address range.
+// CHECK: A: 1
+#pragma omp target data use_device_addr(ph[3 : 4])
+    printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (B) use_device_addr operand in extended address range, but not
+// mapped address range.
+// CHECK: B: 1
+#pragma omp target data use_device_addr(ph[2])
+    printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (C) use_device_addr/map: same base-array, different first-location.
+// CHECK: C: 1
+#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
+    printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (D) use_device_addr/map: different base-array/pointers.
+// CHECK: D: 1
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
+    printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (E) use_device_addr operand within mapped range of previous map.
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa[0])
+    printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+// (F) use_device_addr/map: different operands, same base-array.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+    printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+// (G) use_device_addr/map: different base-array/pointers.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+    printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
new file mode 100644
index 0000000000000..b9ebde431e7bf
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -0,0 +1,143 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+  int *paa[10][10];
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    int *original_ph3 = &ph[3];
+    int **original_paa02 = &paa[0][2];
+
+// (A) No corresponding map, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+    }
+
+// (B) use_device_addr/map: different operands, same base-pointer.
+// use_device_addr operand within mapped address range.
+// CHECK: B: 1 1 1
+#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1])
+    {
+      int *mapped_ptr_ph4 =
+          (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr,
+             mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+    }
+
+// (C) use_device_addr/map: different base-pointers.
+// No corresponding storage, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+    }
+
+// (D) use_device_addr/map: one of two maps with matching base-pointer.
+// use_device_addr operand within mapped address range of second map,
+// lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (E) No corresponding map, lookup should fail
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == (int **)nullptr + 2);
+    }
+
+// (F) use_device_addr/map: different operands, same base-array.
+// use_device_addr within mapped address range. Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == mapped_ptr_paa02);
+    }
+
+// (G) use_device_addr/map: different operands, same base-array.
+// use_device_addr extends beyond existing mapping. Not spec compliant.
+// But the lookup succeeds because we use the base-address for translation.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr(
+          original_paa02 + 2, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr,
+             mapped_ptr_paa04 != original_paa02 + 2,
+             &paa[0][4] == mapped_ptr_paa04);
+    }
+
+    int *original_paa020 = &paa[0][2][0];
+    int **original_paa0 = (int **)&paa[0];
+
+// (H) use_device_addr/map: different base-pointers.
+// No corresponding storage for use_device_addr opnd, lookup should fail.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa020 =
+          (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+      int **mapped_ptr_paa0 =
+          (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
+             mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+    }
+
+// (I) use_device_addr/map: one map with different, one with same base-ptr.
+// Lookup should succeed.
+// CHECK: I: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == mapped_ptr_paa02);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
new file mode 100644
index 0000000000000..e9a1124bc4612
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
@@ -0,0 +1,98 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+  int *(&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    int *original_ph3 = &ph[3];
+    int **original_paa02 = &paa[0][2];
+
+#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
+    int *mapped_ptr_ph3 =
+        (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+    int **mapped_ptr_paa02 =
+        (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+
+    // CHECK-COUNT-4: 1
+    printf("%d\n", mapped_ptr_ph3 != nullptr);
+    printf("%d\n", mapped_ptr_paa02 != nullptr);
+    printf("%d\n", original_ph3 != mapped_ptr_ph3);
+    printf("%d\n", original_paa02 != mapped_ptr_paa02);
+
+// (A) use_device_addr operand within mapped address range.
+// EXPECTED: A: 1
+// CHECK:    A: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[3 : 4])
+    printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (B) use_device_addr operand in extended address range, but not
+// mapped address range.
+// EXPECTED: B: 1
+// CHECK:    B: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[2])
+    printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (C) use_device_addr/map: same base-array, different first-location.
+// EXPECTED: C: 1
+// CHECK:    C: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
+    printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (D) use_device_addr/map: different base-array/pointers.
+// EXPECTED: D: 1
+// CHECK:    D: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
+    printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (E) use_device_addr operand within mapped range of previous map.
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa[0])
+    printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+// (F) use_device_addr/map: different operands, same base-array.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+    printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+// (G) use_device_addr/map: different base-array/pointers.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+    printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
new file mode 100644
index 0000000000000..0090cdb095366
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -0,0 +1,158 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+  int *(&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    int *original_ph3 = &ph[3];
+    int **original_paa02 = &paa[0][2];
+
+// (A) No corresponding map, lookup should fail.
+// EXPECTED: A: 1 1 1
+// CHECK:    A: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+    }
+
+// (B) use_device_addr/map: different operands, same base-pointer.
+// use_device_addr operand within mapped address range.
+// EXPECTED: B: 1 1 1
+// CHECK:    B: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1])
+    {
+      int *mapped_ptr_ph4 =
+          (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr,
+             mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+    }
+
+// (C) use_device_addr/map: different base-pointers.
+// No corresponding storage, lookup should fail.
+// EXPECTED: C: 1 1 1
+// CHECK:    C: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+    }
+
+// (D) use_device_addr/map: one of two maps with matching base-pointer.
+// use_device_addr operand within mapped address range of second map,
+// lookup should succeed.
+// EXPECTED: D: 1 1 1
+// CHECK:    D: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4])
+    {
+      int *mapped_ptr_ph3 =
+          (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (E) No corresponding map, lookup should fail
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == (int **)nullptr + 2);
+    }
+
+// (F) use_device_addr/map: different operands, same base-array.
+// use_device_addr within mapped address range. Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == mapped_ptr_paa02);
+    }
+
+// (G) use_device_addr/map: different operands, same base-array.
+// use_device_addr extends beyond existing mapping. Not spec compliant.
+// But the lookup succeeds because we use the base-address for translation.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr(
+          original_paa02 + 2, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr,
+             mapped_ptr_paa04 != original_paa02 + 2,
+             &paa[0][4] == mapped_ptr_paa04);
+    }
+
+    int *original_paa020 = &paa[0][2][0];
+    int **original_paa0 = (int **)&paa[0];
+
+// (H) use_device_addr/map: different base-pointers.
+// No corresponding storage for use_device_addr opnd, lookup should fail.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+    {
+      int **mapped_ptr_paa020 =
+          (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+      int **mapped_ptr_paa0 =
+          (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
+             mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+    }
+
+// (I) use_device_addr/map: one map with different, one with same base-ptr.
+// Lookup should succeed.
+// CHECK: I: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+    {
+      int **mapped_ptr_paa02 =
+          (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+      printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+             mapped_ptr_paa02 != original_paa02,
+             &paa[0][2] == mapped_ptr_paa02);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
new file mode 100644
index 0000000000000..883297f7e90cd
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
@@ -0,0 +1,93 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on a variable (not a section).
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+  int *paa[10][10];
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    void *original_addr_g = &g;
+    void *original_addr_h = &h;
+    void *original_addr_ph = &ph;
+    void *original_addr_paa = &paa;
+
+#pragma omp target enter data map(to : g, h, ph, paa)
+    void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
+    void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
+    void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
+    void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device());
+
+    // CHECK-COUNT-8: 1
+    printf("%d\n", mapped_ptr_g != nullptr);
+    printf("%d\n", mapped_ptr_h != nullptr);
+    printf("%d\n", mapped_ptr_ph != nullptr);
+    printf("%d\n", mapped_ptr_paa != nullptr);
+    printf("%d\n", original_addr_g != mapped_ptr_g);
+    printf("%d\n", original_addr_h != mapped_ptr_h);
+    printf("%d\n", original_addr_ph != mapped_ptr_ph);
+    printf("%d\n", original_addr_paa != mapped_ptr_paa);
+
+// (A)
+// CHECK: A: 1
+#pragma omp target data use_device_addr(g)
+    printf("A: %d\n", mapped_ptr_g == &g);
+
+// (B)
+// CHECK: B: 1
+#pragma omp target data use_device_addr(h)
+    printf("B: %d\n", mapped_ptr_h == &h);
+
+// (C)
+// CHECK: C: 1
+#pragma omp target data use_device_addr(ph)
+    printf("C: %d\n", mapped_ptr_ph == &ph);
+
+// (D) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &ph, not &ph[0/1].
+// CHECK: D: 1
+#pragma omp target data map(ph[1 : 2]) use_device_addr(ph)
+    printf("D: %d\n", mapped_ptr_ph == &ph);
+
+// (E)
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa)
+    printf("E: %d\n", mapped_ptr_paa == &paa);
+
+// (F) use_device_addr/map with same base-array, paa.
+// Address translation should happen for &paa.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][2]) use_device_addr(paa)
+    printf("F: %d\n", mapped_ptr_paa == &paa);
+
+// (G) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &paa.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+    printf("G: %d\n", mapped_ptr_paa == &paa);
+
+#pragma omp target exit data map(release : g, h, ph, paa)
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
new file mode 100644
index 0000000000000..79c6f69edba8e
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
@@ -0,0 +1,159 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on a variable (not a section).
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+  int *paa[10][10];
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    void *original_addr_g = &g;
+    void *original_addr_h = &h;
+    void *original_addr_ph = &ph;
+    void *original_addr_paa = &paa;
+
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(g)
+    {
+      void *mapped_ptr_g =
+          omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
+             mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+    }
+
+// (B) Lookup should succeed.
+// CHECK: B: 1 1 1
+#pragma omp target data map(g) use_device_addr(g)
+    {
+      void *mapped_ptr_g =
+          omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_g != nullptr,
+             mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+    }
+
+// (C) No corresponding item, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data use_device_addr(h)
+    {
+      void *mapped_ptr_h =
+          omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
+             mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+    }
+
+// (D) Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(h) use_device_addr(h)
+    {
+      void *mapped_ptr_h =
+          omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_h != nullptr,
+             mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+    }
+
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
+             mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+    }
+
+// (F) Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_ph != nullptr,
+             mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+    }
+
+// (G) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: G: 1 1 1
+#pragma omp target data map(ph[0 : 1]) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
+             mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+    }
+
+// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_ph != nullptr,
+             mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+    }
+
+// (I) No corresponding item, lookup should fail.
+// CHECK: I: 1 1 1
+#pragma omp target data use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
+             mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+    }
+
+// (J) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: J: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
+             mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+    }
+
+// (K) Lookup should succeed.
+// CHECK: K: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("K: %d %d %d\n", mapped_ptr_paa != nullptr,
+             mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+    }
+
+// (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: L: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("L: %d %d %d\n", mapped_ptr_paa != nullptr,
+             mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
new file mode 100644
index 0000000000000..f018c65f36ec5
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
@@ -0,0 +1,100 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on a reference variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+  int *(&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    void *original_addr_g = &g;
+    void *original_addr_h = &h;
+    void *original_addr_ph = &ph;
+    void *original_addr_paa = &paa;
+
+#pragma omp target enter data map(to : g, h, ph, paa)
+    void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
+    void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
+    void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
+    void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device());
+
+    // CHECK-COUNT-8: 1
+    printf("%d\n", mapped_ptr_g != nullptr);
+    printf("%d\n", mapped_ptr_h != nullptr);
+    printf("%d\n", mapped_ptr_ph != nullptr);
+    printf("%d\n", mapped_ptr_paa != nullptr);
+    printf("%d\n", original_addr_g != mapped_ptr_g);
+    printf("%d\n", original_addr_h != mapped_ptr_h);
+    printf("%d\n", original_addr_ph != mapped_ptr_ph);
+    printf("%d\n", original_addr_paa != mapped_ptr_paa);
+
+// (A)
+// CHECK: A: 1
+#pragma omp target data use_device_addr(g)
+    printf("A: %d\n", mapped_ptr_g == &g);
+
+// (B)
+// CHECK: B: 1
+#pragma omp target data use_device_addr(h)
+    printf("B: %d\n", mapped_ptr_h == &h);
+
+// (C)
+// CHECK: C: 1
+#pragma omp target data use_device_addr(ph)
+    printf("C: %d\n", mapped_ptr_ph == &ph);
+
+// (D) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &ph, not &ph[0/1].
+// CHECK: D: 1
+#pragma omp target data map(ph[1 : 2]) use_device_addr(ph)
+    printf("D: %d\n", mapped_ptr_ph == &ph);
+
+// (E)
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa)
+    printf("E: %d\n", mapped_ptr_paa == &paa);
+
+// (F) use_device_addr/map with same base-array, paa.
+// Address translation should happen for &paa.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][2]) use_device_addr(paa)
+    printf("F: %d\n", mapped_ptr_paa == &paa);
+
+// (G) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &paa.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+    printf("G: %d\n", mapped_ptr_paa == &paa);
+
+#pragma omp target exit data map(release : g, h, ph, paa)
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
new file mode 100644
index 0000000000000..9360db4195041
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
@@ -0,0 +1,166 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on a reference variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+  int *(&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa[0][2] = &g;
+
+    void *original_addr_g = &g;
+    void *original_addr_h = &h;
+    void *original_addr_ph = &ph;
+    void *original_addr_paa = &paa;
+
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(g)
+    {
+      void *mapped_ptr_g =
+          omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
+             mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+    }
+
+// (B) Lookup should succeed.
+// CHECK: B: 1 1 1
+#pragma omp target data map(g) use_device_addr(g)
+    {
+      void *mapped_ptr_g =
+          omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_g != nullptr,
+             mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+    }
+
+// (C) No corresponding item, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data use_device_addr(h)
+    {
+      void *mapped_ptr_h =
+          omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
+             mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+    }
+
+// (D) Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(h) use_device_addr(h)
+    {
+      void *mapped_ptr_h =
+          omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_h != nullptr,
+             mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+    }
+
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
+             mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+    }
+
+// (F) Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_ph != nullptr,
+             mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+    }
+
+// (G) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: G: 1 1 1
+#pragma omp target data map(ph[0 : 1]) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
+             mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+    }
+
+// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph)
+    {
+      void *mapped_ptr_ph =
+          omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_ph != nullptr,
+             mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+    }
+
+// (I) No corresponding item, lookup should fail.
+// CHECK: I: 1 1 1
+#pragma omp target data use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
+             mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+    }
+
+// (J) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: J: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
+             mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+    }
+
+// (K) Lookup should succeed.
+// CHECK: K: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("K: %d %d %d\n", mapped_ptr_paa != nullptr,
+             mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+    }
+
+// (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: L: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+    {
+      void *mapped_ptr_paa =
+          omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+      printf("L: %d %d %d\n", mapped_ptr_paa != nullptr,
+             mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c
similarity index 94%
rename from offload/test/mapping/target_use_device_addr.c
rename to offload/test/mapping/use_device_addr/target_use_device_addr.c
index 5c2bb8a48f6e6..4a9dbe252f761 100644
--- a/offload/test/mapping/target_use_device_addr.c
+++ b/offload/test/mapping/use_device_addr/target_use_device_addr.c
@@ -12,7 +12,9 @@ int main() {
   printf("%d, %p\n", xp[1], &xp[1]);
 #pragma omp target data use_device_addr(xp[1 : 3]) map(tofrom : x)
 #pragma omp target is_device_ptr(xp)
-  { xp[1] = 222; }
+  {
+    xp[1] = 222;
+  }
   // CHECK: 222
   printf("%d, %p\n", xp[1], &xp[1]);
 }
diff --git a/offload/test/mapping/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
similarity index 93%
rename from offload/test/mapping/target_wrong_use_device_addr.c
rename to offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
index 7a5babd692530..28ec6857fa1a8 100644
--- a/offload/test/mapping/target_wrong_use_device_addr.c
+++ b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
@@ -14,7 +14,7 @@ int main() {
   // CHECK: host addr=0x[[#%x,HOST_ADDR:]]
   fprintf(stderr, "host addr=%p\n", x);
 
-#pragma omp target data map(to : x [0:10])
+#pragma omp target data map(to : x[0 : 10])
   {
 // CHECK: omptarget device 0 info: variable x does not have a valid device
 // counterpart
@@ -27,4 +27,3 @@ int main() {
 
   return 0;
 }
-
diff --git a/offload/test/mapping/array_section_use_device_ptr.c b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c
similarity index 97%
rename from offload/test/mapping/array_section_use_device_ptr.c
rename to offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c
index 86e2875c35c4a..4cfcce28c1127 100644
--- a/offload/test/mapping/array_section_use_device_ptr.c
+++ b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c
@@ -20,7 +20,9 @@ int main() {
 
   float *A_dev = NULL;
 #pragma omp target data use_device_ptr(A)
-  { A_dev = A; }
+  {
+    A_dev = A;
+  }
 #pragma omp target exit data map(delete : A[FROM : LENGTH])
 
   // CHECK: Success
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
new file mode 100644
index 0000000000000..a7745de53298e
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
@@ -0,0 +1,100 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_ptr on a variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int h[10];
+int *ph = &h[0];
+
+struct S {
+  int (*paa)[10][10] = &aa;
+
+  void f1(int i) {
+    paa--;
+    void *original_ph3 = &ph[3];
+    void *original_paa102 = &paa[1][0][2];
+
+#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5])
+    void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+    void *mapped_ptr_paa102 =
+        omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+
+    // CHECK-COUNT-4: 1
+    printf("%d\n", mapped_ptr_ph3 != nullptr);
+    printf("%d\n", mapped_ptr_paa102 != nullptr);
+    printf("%d\n", original_ph3 != mapped_ptr_ph3);
+    printf("%d\n", original_paa102 != mapped_ptr_paa102);
+
+// (A) Mapped data is within extended address range. Lookup should succeed.
+// CHECK: A: 1
+#pragma omp target data use_device_ptr(ph)
+    printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (B) use_device_ptr/map on pointer, and pointee already exists.
+// Lookup should succeed.
+// CHECK: B: 1
+#pragma omp target data map(ph) use_device_ptr(ph)
+    printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: C: 1
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
+    printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: D: 1
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
+    printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (E) Mapped data is within extended address range. Lookup should succeed.
+// Lookup should succeed.
+// CHECK: E: 1
+#pragma omp target data use_device_ptr(paa)
+    printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (F) use_device_ptr/map on pointer, and pointee already exists.
+// &paa[0] should be in extended address-range of the existing paa[1][...]
+// Lookup should succeed.
+// FIXME: However, it currently does not. Might need an RT fix.
+// EXPECTED: F: 1
+// CHECK:    F: 0
+#pragma omp target data map(paa) use_device_ptr(paa)
+    printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+    printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+    printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5])
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
new file mode 100644
index 0000000000000..fe3cdb56e4baa
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
@@ -0,0 +1,125 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_ptr on a variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int h[10];
+int *ph = &h[0];
+
+struct S {
+  int (*paa)[10][10] = &aa;
+
+  void f1(int i) {
+    paa--;
+    void *original_addr_ph3 = &ph[3];
+    void *original_addr_paa102 = &paa[1][0][2];
+
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+    }
+
+// (B) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: B: 1 1 1
+#pragma omp target data map(ph) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+    }
+
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: C: 1 1 1
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+             mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+    }
+
+// (F) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+             mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+    }
+
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+             mapped_ptr_paa102 != original_addr_paa102,
+             &paa[1][0][2] == mapped_ptr_paa102);
+    }
+
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+             mapped_ptr_paa102 != original_addr_paa102,
+             &paa[1][0][2] == mapped_ptr_paa102);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
new file mode 100644
index 0000000000000..66e65de4195a4
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
@@ -0,0 +1,111 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_ptr on a reference variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int (*paa_ptee)[10][10] = &aa;
+
+int h[10];
+int *ph_ptee = &h[0];
+int *&ph = ph_ptee;
+
+struct S {
+  int (*&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa--;
+    void *original_ph3 = &ph[3];
+    void *original_paa102 = &paa[1][0][2];
+
+#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5])
+    void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+    void *mapped_ptr_paa102 =
+        omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+
+    // CHECK-COUNT-4: 1
+    printf("%d\n", mapped_ptr_ph3 != nullptr);
+    printf("%d\n", mapped_ptr_paa102 != nullptr);
+    printf("%d\n", original_ph3 != mapped_ptr_ph3);
+    printf("%d\n", original_paa102 != mapped_ptr_paa102);
+
+// (A) Mapped data is within extended address range. Lookup should succeed.
+// EXPECTED: A: 1
+// CHECK:    A: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_ptr(ph)
+    printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (B) use_device_ptr/map on pointer, and pointee already exists.
+// Lookup should succeed.
+// EXPECTED: B: 1
+// CHECK:    B: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_ptr(ph)
+    printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: C: 1
+// CHECK:    C: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
+    printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: D: 1
+// CHECK:    D: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
+    printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+// (E) Mapped data is within extended address range. Lookup should succeed.
+// Lookup should succeed.
+// CHECK: E: 1
+#pragma omp target data use_device_ptr(paa)
+    printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (F) use_device_ptr/map on pointer, and pointee already exists.
+// &paa[0] should be in extended address-range of the existing paa[1][...]
+// Lookup should succeed.
+// FIXME: However, it currently does not. Might need an RT fix.
+// EXPECTED: F: 1
+// CHECK:    F: 0
+#pragma omp target data map(paa) use_device_ptr(paa)
+    printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+    printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+    printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5])
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
new file mode 100644
index 0000000000000..419ab3eb33d4d
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
@@ -0,0 +1,136 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_ptr on a reference variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int (*paa_ptee)[10][10] = &aa;
+
+int h[10];
+int *ph_ptee = &h[0];
+int *&ph = ph_ptee;
+
+struct S {
+  int (*&paa)[10][10] = paa_ptee;
+
+  void f1(int i) {
+    paa--;
+    void *original_addr_ph3 = &ph[3];
+    void *original_addr_paa102 = &paa[1][0][2];
+
+// (A) No corresponding item, lookup should fail.
+// EXPECTED: A: 1 1 1
+// CHECK:    A: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+    }
+
+// (B) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// EXPECTED: B: 1 1 1
+// CHECK:    B: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+    }
+
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: C: 1 1 1
+// CHECK:    C: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: D: 1 1 1
+// CHECK:    D: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
+    {
+      void *mapped_ptr_ph3 =
+          omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+      printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+             mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+    }
+
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+             mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+    }
+
+// (F) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+             mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+    }
+
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+             mapped_ptr_paa102 != original_addr_paa102,
+             &paa[1][0][2] == mapped_ptr_paa102);
+    }
+
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+    {
+      void *mapped_ptr_paa102 =
+          omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+      printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+             mapped_ptr_paa102 != original_addr_paa102,
+             &paa[1][0][2] == mapped_ptr_paa102);
+    }
+  }
+};
+
+S s1;
+int main() { s1.f1(1); }