From c8a68ec9185d918511b4dbcbaf9e3797b384ed28 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba Date: Fri, 17 Oct 2025 18:25:33 -0700 Subject: [PATCH 1/3] [NFC][OpenMP] Add small class-member use_device_ptr/addr unit tests. Two of the tests are currently asserting, and two are emitting unexpected results. The asserting tests will be fixed using the ATTACH-style codegen from #153683. The other two involve use_device_addr on byrefs, and need more follow-up codegen changes, that have been noted in a FIXME comment. --- ...rget_data_use_device_addr_class_member.cpp | 34 +++++++++++++ ..._data_use_device_addr_class_member_ref.cpp | 34 +++++++++++++ ..._device_addr_class_member_ref_with_map.cpp | 49 +++++++++++++++++++ ..._use_device_addr_class_member_with_map.cpp | 43 ++++++++++++++++ ...arget_data_use_device_ptr_class_member.cpp | 34 +++++++++++++ ...t_data_use_device_ptr_class_member_ref.cpp | 34 +++++++++++++ ...e_device_ptr_class_member_ref_with_map.cpp | 36 ++++++++++++++ ...a_use_device_ptr_class_member_with_map.cpp | 36 ++++++++++++++ 8 files changed, 300 insertions(+) create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member.cpp create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref.cpp create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref_with_map.cpp create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_with_map.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_with_map.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_with_map.cpp 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..429853048c583 --- /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..944ff50dd3962 --- /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..bc63f5c4224ad --- /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..cfa3860213f3d --- /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..36ac5e5782296 --- /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..d208b34c1da74 --- /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..28569954e13b0 --- /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..c096942be5b0d --- /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(); +} From efe0e95a3ba75914d7dbc21ffc8b7d473ca2af99 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba Date: Fri, 17 Oct 2025 18:42:14 -0700 Subject: [PATCH 2/3] Formatting fixes. --- .../target_data_use_device_addr_class_member.cpp | 4 ++-- .../target_data_use_device_addr_class_member_ref.cpp | 4 ++-- .../target_data_use_device_addr_class_member_ref_with_map.cpp | 2 +- .../target_data_use_device_addr_class_member_with_map.cpp | 4 ++-- .../target_data_use_device_ptr_class_member.cpp | 4 ++-- .../target_data_use_device_ptr_class_member_ref.cpp | 4 ++-- .../target_data_use_device_ptr_class_member_ref_with_map.cpp | 4 ++-- .../target_data_use_device_ptr_class_member_with_map.cpp | 4 ++-- 8 files changed, 15 insertions(+), 15 deletions(-) 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 index 429853048c583..6fef34f665b66 100644 --- 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 @@ -23,9 +23,9 @@ struct ST { #pragma omp target data use_device_addr(c) { printf("%d\n", &c == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { 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 index 944ff50dd3962..8ca02ddd0425c 100644 --- 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 @@ -23,9 +23,9 @@ struct ST { #pragma omp target data use_device_addr(d) { printf("%d\n", &d == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { 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 index bc63f5c4224ad..8f58d4cb49fc3 100644 --- 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 @@ -40,7 +40,7 @@ struct ST { (uintptr_t)&d == (uintptr_t)mapped_ptr - offset); } } - }; + } }; int main() { 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 index cfa3860213f3d..f71631ec21526 100644 --- 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 @@ -32,9 +32,9 @@ struct ST { // CHECK: 0 1 printf("%d %d\n", &c == mapped_ptr, (uintptr_t)&c == (uintptr_t)mapped_ptr - offset); - }; + } } - }; + } }; int main() { 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 index 36ac5e5782296..b0253cdbe20d9 100644 --- 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 @@ -23,9 +23,9 @@ struct ST { #pragma omp target data use_device_ptr(a) { printf("%d\n", a == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { 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 index d208b34c1da74..4de34487c2b04 100644 --- 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 @@ -23,9 +23,9 @@ struct ST { #pragma omp target data use_device_ptr(b) { printf("%d\n", b == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { 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 index 28569954e13b0..27fda743b989e 100644 --- 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 @@ -25,9 +25,9 @@ struct ST { #pragma omp target data map(b[0], m) use_device_ptr(b) { printf("%d\n", b == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { 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 index c096942be5b0d..38a369659d13d 100644 --- 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 @@ -25,9 +25,9 @@ struct ST { #pragma omp target data map(a[0], m) use_device_ptr(a) { printf("%d\n", a == mapped_ptr); // CHECK: 1 - }; + } } - }; + } }; int main() { From a6a37a650f83d6141644e1d0bc850a66e333bc8e Mon Sep 17 00:00:00 2001 From: Abhinav Gaba Date: Fri, 17 Oct 2025 18:49:17 -0700 Subject: [PATCH 3/3] More formatting fixes --- .../target_data_use_device_addr_class_member_ref_with_map.cpp | 2 +- .../target_data_use_device_addr_class_member_with_map.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) 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 index 8f58d4cb49fc3..5e8769eb3079d 100644 --- 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 @@ -20,7 +20,7 @@ struct ST { #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 + 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: 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 index f71631ec21526..f5db4ecc66175 100644 --- 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 @@ -23,8 +23,8 @@ struct ST { 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: + // 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.