diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member.cpp new file mode 100644 index 0000000000000..6fef34f665b66 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f7() { +#pragma omp target data map(to : c) + { + void *mapped_ptr = omp_get_mapped_ptr(&c, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data use_device_addr(c) + { + printf("%d\n", &c == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f7(); +} diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref.cpp new file mode 100644 index 0000000000000..8ca02ddd0425c --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f8() { +#pragma omp target enter data map(to : d) + { + void *mapped_ptr = omp_get_mapped_ptr(&d, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data use_device_addr(d) + { + printf("%d\n", &d == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f8(); +} diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref_with_map.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref_with_map.cpp new file mode 100644 index 0000000000000..5e8769eb3079d --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref_with_map.cpp @@ -0,0 +1,49 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f6() { + uintptr_t offset = (uintptr_t)&d - n; +#pragma omp target data map(to : m, d) + { + void *mapped_ptr = omp_get_mapped_ptr(&d, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data map(m, d) use_device_addr(d) + { + // FIXME: Clang is mapping class member references using: + // &this[0], &ref_ptee(this[0].d), 4, PTR_AND_OBJ + // but a load from `this[0]` cannot be used to compute the offset + // in the runtime, because for example in this case, it would mean + // that the base address of the pointee is a load from `n`, i.e. 111. + // clang should be emitting the following instead: + // &ref_ptr(this[0].d), &ref_ptee(this[0].d), 4, PTR_AND_OBJ + // And eventually, the following that's compatible with the + // ref/attach modifiers: + // &ref_ptee(this[0].[d])), &ref_ptee(this[0].d), TO | FROM + // &ref_ptr(this[0].d), &ref_ptee(this[0].d), 4, ATTACH + // EXPECTED: 1 0 + // CHECK: 0 1 + printf("%d %d\n", &d == mapped_ptr, + (uintptr_t)&d == (uintptr_t)mapped_ptr - offset); + } + } + } +}; + +int main() { + ST s; + s.f6(); +} diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_with_map.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_with_map.cpp new file mode 100644 index 0000000000000..f5db4ecc66175 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_with_map.cpp @@ -0,0 +1,43 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f5() { + uintptr_t offset = (uintptr_t)&c - (uintptr_t)this; +#pragma omp target data map(to : m, c) + { + void *mapped_ptr = omp_get_mapped_ptr(&c, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data map(m, c) use_device_addr(c) + { + // FIXME: RT is currently doing the translation for "&this[0]" instead + // of &this->c, for a map like: + // this, &this->c, ..., RETURN_PARAM + // We either need to fix RT, or emit a separate entry for such + // use_device_addr, even if there is a matching map entry already. + // EXPECTED: 1 0 + // CHECK: 0 1 + printf("%d %d\n", &c == mapped_ptr, + (uintptr_t)&c == (uintptr_t)mapped_ptr - offset); + } + } + } +}; + +int main() { + ST s; + s.f5(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member.cpp new file mode 100644 index 0000000000000..b0253cdbe20d9 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f3() { +#pragma omp target data map(to : a[0]) + { + void *mapped_ptr = omp_get_mapped_ptr(a, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data use_device_ptr(a) + { + printf("%d\n", a == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f3(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref.cpp new file mode 100644 index 0000000000000..4de34487c2b04 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f4() { +#pragma omp target data map(to : b[0]) + { + void *mapped_ptr = omp_get_mapped_ptr(b, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data use_device_ptr(b) + { + printf("%d\n", b == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f4(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_with_map.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_with_map.cpp new file mode 100644 index 0000000000000..27fda743b989e --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_with_map.cpp @@ -0,0 +1,36 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f2() { +#pragma omp target data map(to : b[0]) + { + void *mapped_ptr = omp_get_mapped_ptr(b, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data map(b[0], m) use_device_ptr(b) + { + printf("%d\n", b == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f2(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_with_map.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_with_map.cpp new file mode 100644 index 0000000000000..38a369659d13d --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_with_map.cpp @@ -0,0 +1,36 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include +#include + +int x = 0; +int *y = &x; +int z = 0; + +struct ST { + int n = 111; + int *a = &x; + int *&b = y; + int c = 0; + int &d = z; + int m = 0; + + void f1() { +#pragma omp target data map(to : a[0]) + { + void *mapped_ptr = omp_get_mapped_ptr(a, omp_get_default_device()); + printf("%d\n", mapped_ptr != NULL); // CHECK: 1 +#pragma omp target data map(a[0], m) use_device_ptr(a) + { + printf("%d\n", a == mapped_ptr); // CHECK: 1 + } + } + } +}; + +int main() { + ST s; + s.f1(); +}