From 67163666c3f84ffa89634df5ff730ac40b13ba11 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba 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 + +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 + +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 + +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 + +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 + +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 + +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 +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 +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 +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 +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 +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(); }