[llvm] [NFC][OpenMP][Offload] Add tests for `use_device_ptr(fb_preserve/nullify)`. (PR #170948)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-offload

Author: Abhinav Gaba (abhinavgaba)

<details>
<summary>Changes</summary>

The `fallback` modifier for `use_device_ptr` is currently part of OpenMP 6.1. The tests mostly fail for now. The associated libomptarget and clang parsing/sema changes are in #<!-- -->169438, #<!-- -->169603 and #<!-- -->170578, with clang codegen to follow.

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


12 Files Affected:

- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp (+34) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp (+30) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp (+30) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp (+35) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp (+31) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp (+31) 
- (renamed) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp (+5-1) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp (+23) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp (+23) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp (+26) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp (+25) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp (+24) 


``````````diff
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp
new file mode 100644
index 0000000000000..2dd33732a7d2f
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp
@@ -0,0 +1,34 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-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.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(a)
+    printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp
new file mode 100644
index 0000000000000..61f3367f537ee
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp
@@ -0,0 +1,30 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:          0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_nullify : a)
+    printf("%p\n", a); // OFFLOAD-NEXT:   (nil)
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp
new file mode 100644
index 0000000000000..b7af1f39cc3bf
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp
@@ -0,0 +1,30 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : a)
+    printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp
new file mode 100644
index 0000000000000..45f89d0ee92cc
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp
@@ -0,0 +1,35 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-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.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(b)
+    printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp
new file mode 100644
index 0000000000000..39f39a974577c
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp
@@ -0,0 +1,31 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:          0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_nullify : b)
+    printf("%p\n", b); // OFFLOAD-NEXT:   (nil)
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp
new file mode 100644
index 0000000000000..ad861ab12001e
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp
@@ -0,0 +1,31 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+// XFAIL: *
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : b)
+    printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
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.cpp
similarity index 82%
rename from offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
rename to offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp
index e8fa3b69e9296..54faea65aec08 100644
--- 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.cpp
@@ -1,4 +1,8 @@
-// RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
 
 // Test that when a use_device_ptr lookup fails, the
 // privatized pointer retains its original value by
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp
new file mode 100644
index 0000000000000..cb11f52645dd8
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp
@@ -0,0 +1,23 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_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(fb_nullify : xp)
+  printf("%p\n", xp); // OFFLOAD-NEXT:   (nil)
+                      // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp
new file mode 100644
index 0000000000000..31ce803fc1ed0
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp
@@ -0,0 +1,23 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+// 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(fb_preserve : xp)
+  printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp
new file mode 100644
index 0000000000000..1060ed9cdbc70
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp
@@ -0,0 +1,26 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-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.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(xpr)
+  printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp
new file mode 100644
index 0000000000000..230ffda4fad9a
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp
@@ -0,0 +1,25 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+// XFAIL: *
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:          0x[[#%x,ADDR:]]
+  // FIXME: We won't get "nil" until we start privatizing xpr.
+#pragma omp target data use_device_ptr(fb_nullify : xpr)
+  printf("%p\n", xpr); // OFFLOAD-NEXT:   (nil)
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp
new file mode 100644
index 0000000000000..443739814ed0d
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+// XFAIL: *
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : xpr)
+  printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }

``````````

</details>


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