[NFC][OpenMP] Add use_device_ptr/addr tests for when the lookup fails.

[abhinavgaba]

As per OpenMP 5.1, the pointers are expected to retain their original values when a lookup fails and there is no device pointer to translate to.

[llvmbot]

@llvm/pr-subscribers-offload

Author: Abhinav Gaba (abhinavgaba)

Changes

As per OpenMP 5.1, the pointers are expected to retain their original values when a lookup fails and there is no device pointer to translate to.

---

Full diff: https://github.com/llvm/llvm-project/pull/169428.diff

3 Files Affected:

• (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c (+24)
• (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c (+21)
• (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c (+32)

diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
new file mode 100644
index 0000000000000..4b67a3bc2aa7f
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// Test that when a use_device_addr lookup fails, the
+// list-item retains its original address by default.
+//
+// This is necessary because we must assume that the
+// list-item is device-accessible, even if it was not
+// previously mapped.
+
+// XFAIL: *
+
+#include <stdio.h>
+int h[10];
+int *ph = &h[0];
+
+void f1() {
+  printf("%p\n", &h[2]); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_addr(h[2])
+  printf("%p\n", &h[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+#pragma omp target data use_device_addr(ph[2])
+  printf("%p\n", &ph[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
new file mode 100644
index 0000000000000..4495a46b6d204
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
@@ -0,0 +1,21 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// Test that when a use_device_addr lookup fails, the
+// list-item retains its original address by default.
+//
+// This is necessary because we must assume that the
+// list-item is device-accessible, even if it was not
+// previously mapped.
+
+// XFAIL: *
+
+#include <stdio.h>
+int x;
+
+void f1() {
+  printf("%p\n", &x); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_addr(x)
+  printf("%p\n", &x); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
new file mode 100644
index 0000000000000..e8fa3b69e9296
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value by
+// default.
+//
+// This is necessary because we must assume that the
+// pointee is device-accessible, even if it was not
+// previously mapped.
+//
+// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
+// If a list item that appears in a use_device_ptr clause ... does not point to
+// a mapped object, it must contain a valid device address for the target
+// device, and the list item references are instead converted to references to a
+// local device pointer that refers to this device address.
+//
+// Note: OpenMP 6.1 will have a way to change the
+// fallback behavior: preserve or nullify.
+
+// XFAIL: *
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+
+void f1() {
+  printf("%p\n", xp); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(xp)
+  printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }