diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c new file mode 100644 index 0000000000000..4b67a3bc2aa7f --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c @@ -0,0 +1,24 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// Test that when a use_device_addr lookup fails, the +// list-item retains its original address by default. +// +// This is necessary because we must assume that the +// list-item is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include +int h[10]; +int *ph = &h[0]; + +void f1() { + printf("%p\n", &h[2]); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_addr(h[2]) + printf("%p\n", &h[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +#pragma omp target data use_device_addr(ph[2]) + printf("%p\n", &ph[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c new file mode 100644 index 0000000000000..4495a46b6d204 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c @@ -0,0 +1,21 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// Test that when a use_device_addr lookup fails, the +// list-item retains its original address by default. +// +// This is necessary because we must assume that the +// list-item is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include +int x; + +void f1() { + printf("%p\n", &x); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_addr(x) + printf("%p\n", &x); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } 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.c new file mode 100644 index 0000000000000..e8fa3b69e9296 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c @@ -0,0 +1,32 @@ +// RUN: %libomptarget-compilexx-run-and-check-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. +// +// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31: +// If a list item that appears in a use_device_ptr clause ... does not point to +// a mapped object, it must contain a valid device address for the target +// device, and the list item references are instead converted to references to a +// local device pointer that refers to this device address. +// +// Note: OpenMP 6.1 will have a way to change the +// fallback behavior: preserve or 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(xp) + printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); }