[llvm] [NFC][OpenMP] Add small class-member use_device_ptr/addr unit tests. (PR #164039)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Fri Oct 17 18:35:50 PDT 2025
https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/164039
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.
>From c8a68ec9185d918511b4dbcbaf9e3797b384ed28 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 17 Oct 2025 18:25:33 -0700
Subject: [PATCH] [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 <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..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 <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..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 <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..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 <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..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 <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..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 <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..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 <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..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 <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();
+}
More information about the llvm-commits
mailing list