[llvm-branch-commits] [llvm] [NFC][OpenMP][Offload] Add tests for `use_device_ptr(fb_preserve/nullify)`. (3/4) (PR #173930)
Abhinav Gaba via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Dec 29 15:21:32 PST 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/173930
>From df4ef8fa806260ac5122e238a818cc36201af6b1 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 1/3] [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} (69%)
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 69%
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 33a363495e24a..5be209a8c5b0d 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(); }
>From 1d3d08d654a0fcf059fa4abd469e9b2812e5dd45 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Dec 2025 15:03:49 -0800
Subject: [PATCH 2/3] Update test checks to incorporate Parsing/Sema changes.
---
.../target_data_use_device_ptr_class_member_fallback.cpp | 2 --
...et_data_use_device_ptr_class_member_fallback_nullify.cpp | 6 +++---
...t_data_use_device_ptr_class_member_fallback_preserve.cpp | 2 --
...target_data_use_device_ptr_class_member_ref_fallback.cpp | 2 --
...ata_use_device_ptr_class_member_ref_fallback_nullify.cpp | 6 +++---
...ta_use_device_ptr_class_member_ref_fallback_preserve.cpp | 2 --
.../target_data_use_device_ptr_var_fallback_nullify.cpp | 6 +++---
.../target_data_use_device_ptr_var_fallback_preserve.cpp | 2 --
.../target_data_use_device_ptr_var_ref_fallback_nullify.cpp | 5 ++---
...target_data_use_device_ptr_var_ref_fallback_preserve.cpp | 2 --
10 files changed, 11 insertions(+), 24 deletions(-)
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
index 2dd33732a7d2f..5c232d5db02e0 100644
--- 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
@@ -12,8 +12,6 @@
// pointee is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int x = 0;
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
index 61f3367f537ee..e85f51736f4f7 100644
--- 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
@@ -7,8 +7,6 @@
// 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;
@@ -18,8 +16,10 @@ struct ST {
void f1() {
printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]]
+ // FIXME: Update this with codegen changes for fb_nullify
#pragma omp target data use_device_ptr(fb_nullify : a)
- printf("%p\n", a); // OFFLOAD-NEXT: (nil)
+ printf("%p\n", a); // EXPECTED-OFFLOAD-NEXT: (nil)
+ // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
};
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
index b7af1f39cc3bf..51944c561dc2d 100644
--- 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
@@ -8,8 +8,6 @@
// privatized pointer retains its original value
// because of fb_preserve.
-// XFAIL: *
-
#include <stdio.h>
int x = 0;
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
index 45f89d0ee92cc..59a8facdf2f5c 100644
--- 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
@@ -12,8 +12,6 @@
// pointee is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int x = 0;
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
index 39f39a974577c..00e6372a0f588 100644
--- 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
@@ -7,8 +7,6 @@
// 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;
@@ -19,8 +17,10 @@ struct ST {
void f2() {
printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]]
+ // FIXME: Update this with codegen changes for fb_nullify
#pragma omp target data use_device_ptr(fb_nullify : b)
- printf("%p\n", b); // OFFLOAD-NEXT: (nil)
+ printf("%p\n", b); // EXPECTED-OFFLOAD-NEXT: (nil)
+ // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
};
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
index ad861ab12001e..beeb7526e1625 100644
--- 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
@@ -8,8 +8,6 @@
// privatized pointer retains its original value
// because of fb_preserve.
-// XFAIL: *
-
#include <stdio.h>
int x = 0;
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
index cb11f52645dd8..2d4cd11463801 100644
--- 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
@@ -7,16 +7,16 @@
// 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:]]
+ // FIXME: Update this with codegen changes for fb_nullify
#pragma omp target data use_device_ptr(fb_nullify : xp)
- printf("%p\n", xp); // OFFLOAD-NEXT: (nil)
+ printf("%p\n", xp); // EXPECTED-OFFLOAD-NEXT: (nil)
+ // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
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
index 31ce803fc1ed0..197704f14f86a 100644
--- 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
@@ -8,8 +8,6 @@
// privatized pointer retains its original value
// because of fb_preserve.
-// XFAIL: *
-
#include <stdio.h>
int x;
int *xp = &x;
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
index 230ffda4fad9a..7fa76dd69e7c0 100644
--- 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
@@ -7,8 +7,6 @@
// 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;
@@ -18,7 +16,8 @@ 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)
+ printf("%p\n", xpr); // EXPECTED-OFFLOAD-NEXT: (nil)
+ // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
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
index 443739814ed0d..e7f8bd48ec4fe 100644
--- 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
@@ -8,8 +8,6 @@
// privatized pointer retains its original value
// because of fb_preserve.
-// XFAIL: *
-
#include <stdio.h>
int x;
int *xp = &x;
>From 85241dcaef55f29227ad911257198c173aa0e7f8 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Dec 2025 15:21:04 -0800
Subject: [PATCH 3/3] Clang-format fixes.
---
...target_data_use_device_ptr_class_member_fallback_nullify.cpp | 2 +-
...et_data_use_device_ptr_class_member_ref_fallback_nullify.cpp | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
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
index e85f51736f4f7..9745276294078 100644
--- 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
@@ -16,7 +16,7 @@ struct ST {
void f1() {
printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]]
- // FIXME: Update this with codegen changes for fb_nullify
+ // FIXME: Update this with codegen changes for fb_nullify
#pragma omp target data use_device_ptr(fb_nullify : a)
printf("%p\n", a); // EXPECTED-OFFLOAD-NEXT: (nil)
// OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
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
index 00e6372a0f588..76610a95af512 100644
--- 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
@@ -17,7 +17,7 @@ struct ST {
void f2() {
printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]]
- // FIXME: Update this with codegen changes for fb_nullify
+ // FIXME: Update this with codegen changes for fb_nullify
#pragma omp target data use_device_ptr(fb_nullify : b)
printf("%p\n", b); // EXPECTED-OFFLOAD-NEXT: (nil)
// OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
More information about the llvm-branch-commits
mailing list