[NFC][OpenMP] Add small class-member use_device_ptr/addr unit tests. #164039
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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 llvm#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.
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
|
@llvm/pr-subscribers-offload Author: Abhinav Gaba (abhinavgaba) ChangesTwo 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 Full diff: https://github.com/llvm/llvm-project/pull/164039.diff 8 Files Affected:
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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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();
+}
|
|
The new tests are failing for me on Gentoo/arm64: Full build log: llvm-runtimes:offload-22.0.0.9999:20251021-082951.log.gz |
abhinavgaba
added a commit
to abhinavgaba/llvm-project
that referenced
this pull request
Oct 21, 2025
The failure was reported here: llvm#164039 (comment) The test is checking for the "bad" behavior so as to keep track of it. It could just be updated to check the "good" behavior and marked as XFAIL, but then it would not be as informative. The update for now is to fix the pointer arithmetic. If that isn't sufficient, we can fall back to doing the above.
|
I've been testing LLVM 21.x release in the meantime. Need 30 minutes or so to finish rebuilding. |
abhinavgaba
added a commit
that referenced
this pull request
Oct 22, 2025
The failure was reported here: #164039 (comment) The test was checking for the "bad" behavior so as to keep track of it, but there seem to be some issues with the pointer arithmetic specific to aarch64. The update for now is to not check for the "bad" behavior fully. We may need to debug further if similar issues are encountered eventually once the codegen has been fixed.
llvm-sync bot
pushed a commit
to arm/arm-toolchain
that referenced
this pull request
Oct 22, 2025
…#164456) The failure was reported here: llvm/llvm-project#164039 (comment) The test was checking for the "bad" behavior so as to keep track of it, but there seem to be some issues with the pointer arithmetic specific to aarch64. The update for now is to not check for the "bad" behavior fully. We may need to debug further if similar issues are encountered eventually once the codegen has been fixed.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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_addron byrefs, and need more follow-up codegen changes, that have been noted in a FIXME comment.