[llvm] [NFC][OpenMP][Offload] Add tests for `use_device_ptr(fb_preserve/nullify)`. (PR #170948)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Fri Dec 5 15:21:21 PST 2025
https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/170948
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.
>From 67163666c3f84ffa89634df5ff730ac40b13ba11 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 5 Dec 2025 15:10:27 -0800
Subject: [PATCH] [NFC][OpenMP][Offload] Add tests for
`use_device_ptr(fb_preserve/nullify)`.
The fallback modifiers are 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.
---
...a_use_device_ptr_class_member_fallback.cpp | 34 ++++++++++++++++++
...vice_ptr_class_member_fallback_nullify.cpp | 30 ++++++++++++++++
...ice_ptr_class_member_fallback_preserve.cpp | 30 ++++++++++++++++
...e_device_ptr_class_member_ref_fallback.cpp | 35 +++++++++++++++++++
..._ptr_class_member_ref_fallback_nullify.cpp | 31 ++++++++++++++++
...ptr_class_member_ref_fallback_preserve.cpp | 31 ++++++++++++++++
...rget_data_use_device_ptr_var_fallback.cpp} | 6 +++-
...ta_use_device_ptr_var_fallback_nullify.cpp | 23 ++++++++++++
...a_use_device_ptr_var_fallback_preserve.cpp | 23 ++++++++++++
...t_data_use_device_ptr_var_ref_fallback.cpp | 26 ++++++++++++++
...se_device_ptr_var_ref_fallback_nullify.cpp | 25 +++++++++++++
...e_device_ptr_var_ref_fallback_preserve.cpp | 24 +++++++++++++
12 files changed, 317 insertions(+), 1 deletion(-)
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp
rename offload/test/mapping/use_device_ptr/{target_data_use_device_ptr_var_fallback.c => target_data_use_device_ptr_var_fallback.cpp} (82%)
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp
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(); }
More information about the llvm-commits
mailing list