[llvm] [NFC][OpenMP] Add small class-member use_device_ptr/addr unit tests. (PR #164039)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-offload

Author: Abhinav Gaba (abhinavgaba)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/164039.diff


8 Files Affected:

- (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member.cpp (+34) 
- (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref.cpp (+34) 
- (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_ref_with_map.cpp (+49) 
- (added) offload/test/mapping/use_device_addr/target_data_use_device_addr_class_member_with_map.cpp (+43) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member.cpp (+34) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref.cpp (+34) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_with_map.cpp (+36) 
- (added) offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_with_map.cpp (+36) 


``````````diff
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();
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/164039