Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -501,7 +501,8 @@
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast,
-                                              UpdateRef, IsHostPtr);
+                                              UpdateRef, IsHostPtr,
+                                              /*MustContain=*/true);
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));
Index: openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c
@@ -0,0 +1,140 @@
+// --------------------------------------------------
+// Check extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=BEFORE
+// XUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// XUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=AFTER
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DEXTENDS=AFTER
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// END.
+
+#include <stdio.h>
+
+#define BEFORE 0
+#define AFTER  1
+
+#define SIZE 100
+
+#if EXTENDS == BEFORE
+# define SMALL_BEG (SIZE-2)
+# define SMALL_END SIZE
+# define LARGE_BEG 0
+# define LARGE_END SIZE
+#elif EXTENDS == AFTER
+# define SMALL_BEG 0
+# define SMALL_END 2
+# define LARGE_BEG 0
+# define LARGE_END SIZE
+#else
+# error EXTENDS undefined
+#endif
+
+#define SMALL SMALL_BEG:(SMALL_END-SMALL_BEG)
+#define LARGE LARGE_BEG:(LARGE_END-LARGE_BEG)
+
+void check_not_present() {
+  int arr[SIZE];
+
+  for (int i = 0; i < SIZE; ++i)
+    arr[i] = 99;
+
+  // CHECK-LABEL: checking not present
+  fprintf(stderr, "checking not present\n");
+
+  // arr[LARGE] isn't (fully) present at the end of the target data region, so
+  // the device-to-host transfer should not be performed, or it might fail.
+#pragma omp target data map(tofrom: arr[LARGE])
+  {
+#pragma omp target exit data map(delete: arr[LARGE])
+#pragma omp target enter data map(alloc: arr[SMALL])
+#pragma omp target map(alloc: arr[SMALL])
+    for (int i = SMALL_BEG; i < SMALL_END; ++i)
+      arr[i] = 88;
+  }
+
+  // CHECK-NOT: Libomptarget
+  // CHECK-NOT: error
+  for (int i = 0; i < SIZE; ++i) {
+    if (arr[i] != 99)
+      fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+  }
+}
+
+void check_is_present() {
+  int arr[SIZE];
+
+  for (int i = 0; i < SIZE; ++i)
+    arr[i] = 99;
+
+  // CHECK-LABEL: checking is present
+  fprintf(stderr, "checking is present\n");
+
+  // arr[SMALL] is (fully) present at the end of the target data region, and the
+  // device-to-host transfer should be performed only for it even though more
+  // of the array is then present.
+#pragma omp target data map(tofrom: arr[SMALL])
+  {
+#pragma omp target exit data map(delete: arr[SMALL])
+#pragma omp target enter data map(alloc: arr[LARGE])
+#pragma omp target map(alloc: arr[LARGE])
+    for (int i = LARGE_BEG; i < LARGE_END; ++i)
+      arr[i] = 88;
+  }
+
+  // CHECK-NOT: Libomptarget
+  // CHECK-NOT: error
+  for (int i = 0; i < SIZE; ++i) {
+    if (SMALL_BEG <= i && i < SMALL_END) {
+      if (arr[i] != 88)
+        fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+    } else if (arr[i] != 99) {
+      fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+    }
+  }
+}
+
+int main() {
+  check_not_present();
+  check_is_present();
+  return 0;
+}