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 <[email protected]> Date: Fri, 5 Dec 2025 15:10:27 -0800 Subject: [PATCH 1/4] [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 <[email protected]> Date: Mon, 29 Dec 2025 15:03:49 -0800 Subject: [PATCH 2/4] 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 <[email protected]> Date: Mon, 29 Dec 2025 15:21:04 -0800 Subject: [PATCH 3/4] 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]] >From fdfa6ed54c24f5118e56ccaf6280727eb21f1b0f Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 30 Dec 2025 13:37:37 -0800 Subject: [PATCH 4/4] Clang-format changed its mind. --- ...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 9745276294078..3094446f8b44d 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 76610a95af512..39a987b08a505 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]] _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
