[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:49:31 PDT 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/164039
>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 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 <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();
+}
>From efe0e95a3ba75914d7dbc21ffc8b7d473ca2af99 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
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 <abhinav.gaba at intel.com>
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.
More information about the llvm-commits
mailing list