[llvm] [NFC][OpenMP] Add several use_device_ptr/addr tests. (PR #154939)

[Abhinav Gaba via llvm-commits]

================
@@ -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)
----------------
abhinavgaba wrote:

@dreachem, please check if the expectation that this lookup should succeed, is valid. The pointee is mapped on the outer target_enter_data.

https://github.com/llvm/llvm-project/pull/154939