1 files changed, 61 insertions, 0 deletions
diff --git a/offload/test/offloading/force-usm.cpp b/offload/test/offloading/force-usm.cpp
new file mode 100644
index 000000000000..a043ba47f54a
--- /dev/null
+++ b/offload/test/offloading/force-usm.cpp
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic
+// RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=NO-USM
+//
+// RUN: %libomptarget-compilexxx-generic-force-usm
+// RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=32 \
+// RUN:       %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=FORCE-USM
+//
+// REQUIRES: unified_shared_memory
+//
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// clang-format on
+
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+
+int GI;
+#pragma omp declare target
+int *pGI;
+#pragma omp end declare target
+
+int main(void) {
+
+  GI = 0;
+  // Implicit mappings
+  int alpha = 1;
+  int beta[3] = {2, 5, 8};
+
+  // Require map clauses for non-USM execution
+  pGI = (int *)malloc(sizeof(int));
+  *pGI = 42;
+
+#pragma omp target map(pGI[ : 1], GI)
+  {
+    GI = 1 * alpha;
+    *pGI = 2 * beta[1];
+  }
+
+  assert(GI == 1);
+  assert(*pGI == 10);
+
+  printf("SUCCESS\n");
+
+  return 0;
+}
+
+// clang-format off
+// NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
+// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12
+// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
+// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI
+// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
+// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=12
+// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
+// NO-USM-NEXT: SUCCESS
+
+// FORCE-USM: SUCCESS
+//
+// clang-format on