[llvm] [NFC][OpenMP] Add several use_device_ptr/addr tests. (PR #154939)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 22 06:13:23 PDT 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/154939
>From b86bd0addcf626eac8a4b8922916926207e6ec01 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 22 Aug 2025 04:22:53 -0700
Subject: [PATCH 1/3] [NFC][OpenMP] Add various combinations of
use_device_ptr/addr tests.
Most of the non-reference tests should start passing once we start using
ATTACH map-type based codegen.
The reference tests have a different issue wherein the clause operand is
not being privatized, and the target-data region is accessing the
original. That needs to be fixed separately.
---
...t_data_use_device_addr_arrsec_existing.cpp | 85 +++++++++++
...ta_use_device_addr_arrsec_not_existing.cpp | 121 +++++++++++++++
...ta_use_device_addr_arrsec_ref_existing.cpp | 98 ++++++++++++
...se_device_addr_arrsec_ref_not_existing.cpp | 136 +++++++++++++++++
...rget_data_use_device_addr_var_existing.cpp | 95 ++++++++++++
..._data_use_device_addr_var_not_existing.cpp | 137 +++++++++++++++++
..._data_use_device_addr_var_ref_existing.cpp | 102 +++++++++++++
...a_use_device_addr_var_ref_not_existing.cpp | 144 ++++++++++++++++++
.../target_use_device_addr.c | 0
.../target_wrong_use_device_addr.c | 0
.../array_section_use_device_ptr.c | 0
.../target_data_use_device_ptr_existing.cpp | 102 +++++++++++++
...arget_data_use_device_ptr_not_existing.cpp | 109 +++++++++++++
...arget_data_use_device_ptr_ref_existing.cpp | 113 ++++++++++++++
...t_data_use_device_ptr_ref_not_existing.cpp | 120 +++++++++++++++
15 files changed, 1362 insertions(+)
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
create mode 100644 offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
rename offload/test/mapping/{ => use_device_addr}/target_use_device_addr.c (100%)
rename offload/test/mapping/{ => use_device_addr}/target_wrong_use_device_addr.c (100%)
rename offload/test/mapping/{ => use_device_ptr}/array_section_use_device_ptr.c (100%)
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
new file mode 100644
index 0000000000000..eeb18ef57ca12
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+ int *paa[10][10];
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ int *original_ph3 = &ph[3];
+ int **original_paa02 = &paa[0][2];
+
+ #pragma omp target enter data map(to:ph[3:4], paa[0][2:5])
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+
+ // CHECK-COUNT-4: 1
+ printf("%d\n", mapped_ptr_ph3 != nullptr);
+ printf("%d\n", mapped_ptr_paa02 != nullptr);
+ printf("%d\n", original_ph3 != mapped_ptr_ph3);
+ printf("%d\n", original_paa02 != mapped_ptr_paa02);
+
+ // (A) use_device_addr operand within mapped address range.
+ // CHECK: A: 1
+ #pragma omp target data use_device_addr(ph[3:4])
+ printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (B) use_device_addr operand in extended address range, but not
+ // mapped address range.
+ // CHECK: B: 1
+ #pragma omp target data use_device_addr(ph[2])
+ printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (C) use_device_addr/map: same base-array, different first-location.
+ // CHECK: C: 1
+ #pragma omp target data map(ph[3:2]) use_device_addr(ph[4:1])
+ printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (D) use_device_addr/map: different base-array/pointers.
+ // CHECK: D: 1
+ #pragma omp target data map(ph) use_device_addr(ph[3:4])
+ printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (E) use_device_addr operand within mapped range of previous map.
+ // CHECK: E: 1
+ #pragma omp target data use_device_addr(paa[0])
+ printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ // (F) use_device_addr/map: different operands, same base-array.
+ // CHECK: F: 1
+ #pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+ printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ // (G) use_device_addr/map: different base-array/pointers.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+ printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ #pragma omp target exit data map(release:ph[3:4], paa[0][2:5])
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
new file mode 100644
index 0000000000000..11543dffcce6e
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -0,0 +1,121 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+ int *paa[10][10];
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ int *original_ph3 = &ph[3];
+ int **original_paa02 = &paa[0][2];
+
+ // (A) No corresponding map, lookup should fail.
+ // CHECK: A: 1 1 1
+ #pragma omp target data use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ }
+
+ // (B) use_device_addr/map: different operands, same base-pointer.
+ // use_device_addr operand within mapped address range.
+ // CHECK: B: 1 1 1
+ #pragma omp target data map(ph[2:3]) use_device_addr(ph[3:1])
+ {
+ int *mapped_ptr_ph4 = (int*) omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+ }
+
+ // (C) use_device_addr/map: different base-pointers.
+ // No corresponding storage, lookup should fail.
+ // CHECK: C: 1 1 1
+ #pragma omp target data map(ph) use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ }
+
+ // (D) use_device_addr/map: one of two maps with matching base-pointer.
+ // use_device_addr operand within mapped address range of second map,
+ // lookup should succeed.
+ // CHECK: D: 1 1 1
+ #pragma omp target data map(ph) map(ph[2:5]) use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (E) No corresponding map, lookup should fail
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == (int**) nullptr + 2);
+ }
+
+ // (F) use_device_addr/map: different operands, same base-array.
+ // use_device_addr within mapped address range. Lookup should succeed.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(paa) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ }
+
+ // (G) use_device_addr/map: different operands, same base-array.
+ // use_device_addr extends beyond existing mapping. Not spec compliant.
+ // But the lookup succeeds because we use the base-address for translation.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa04 = (int**) omp_get_mapped_ptr(original_paa02 + 2, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, mapped_ptr_paa04 != original_paa02 + 2, &paa[0][4] == mapped_ptr_paa04);
+ }
+
+ int *original_paa020 = &paa[0][2][0];
+ int **original_paa0 = (int**) &paa[0];
+ // (H) use_device_addr/map: different base-pointers.
+ // No corresponding storage for use_device_addr opnd, lookup should fail.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa020 = (int**) omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+ int **mapped_ptr_paa0 = (int**) omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ }
+
+ // (I) use_device_addr/map: one map with different, one with same base-ptr.
+ // Lookup should succeed.
+ // CHECK: I: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
new file mode 100644
index 0000000000000..502bcdf2ad2b8
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
@@ -0,0 +1,98 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+ int *(&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ int *original_ph3 = &ph[3];
+ int **original_paa02 = &paa[0][2];
+
+ #pragma omp target enter data map(to:ph[3:4], paa[0][2:5])
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+
+ // CHECK-COUNT-4: 1
+ printf("%d\n", mapped_ptr_ph3 != nullptr);
+ printf("%d\n", mapped_ptr_paa02 != nullptr);
+ printf("%d\n", original_ph3 != mapped_ptr_ph3);
+ printf("%d\n", original_paa02 != mapped_ptr_paa02);
+
+ // (A) use_device_addr operand within mapped address range.
+ // EXPECTED: A: 1
+ // CHECK: A: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data use_device_addr(ph[3:4])
+ printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (B) use_device_addr operand in extended address range, but not
+ // mapped address range.
+ // EXPECTED: B: 1
+ // CHECK: B: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data use_device_addr(ph[2])
+ printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (C) use_device_addr/map: same base-array, different first-location.
+ // EXPECTED: C: 1
+ // CHECK: C: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph[3:2]) use_device_addr(ph[4:1])
+ printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (D) use_device_addr/map: different base-array/pointers.
+ // EXPECTED: D: 1
+ // CHECK: D: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) use_device_addr(ph[3:4])
+ printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (E) use_device_addr operand within mapped range of previous map.
+ // CHECK: E: 1
+ #pragma omp target data use_device_addr(paa[0])
+ printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ // (F) use_device_addr/map: different operands, same base-array.
+ // CHECK: F: 1
+ #pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+ printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ // (G) use_device_addr/map: different base-array/pointers.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+ printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
+
+ #pragma omp target exit data map(release:ph[3:4], paa[0][2:5])
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
new file mode 100644
index 0000000000000..18436dbee79a6
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -0,0 +1,136 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+ int *(&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ int *original_ph3 = &ph[3];
+ int **original_paa02 = &paa[0][2];
+
+ // (A) No corresponding map, lookup should fail.
+ // EXPECTED: A: 1 1 1
+ // CHECK: A: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ }
+
+ // (B) use_device_addr/map: different operands, same base-pointer.
+ // use_device_addr operand within mapped address range.
+ // EXPECTED: B: 1 1 1
+ // CHECK: B: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph[2:3]) use_device_addr(ph[3:1])
+ {
+ int *mapped_ptr_ph4 = (int*) omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+ }
+
+ // (C) use_device_addr/map: different base-pointers.
+ // No corresponding storage, lookup should fail.
+ // EXPECTED: C: 1 1 1
+ // CHECK: C: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ }
+
+ // (D) use_device_addr/map: one of two maps with matching base-pointer.
+ // use_device_addr operand within mapped address range of second map,
+ // lookup should succeed.
+ // EXPECTED: D: 1 1 1
+ // CHECK: D: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) map(ph[2:5]) use_device_addr(ph[3:4])
+ {
+ int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (E) No corresponding map, lookup should fail
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == (int**) nullptr + 2);
+ }
+
+ // (F) use_device_addr/map: different operands, same base-array.
+ // use_device_addr within mapped address range. Lookup should succeed.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(paa) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ }
+
+ // (G) use_device_addr/map: different operands, same base-array.
+ // use_device_addr extends beyond existing mapping. Not spec compliant.
+ // But the lookup succeeds because we use the base-address for translation.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa04 = (int**) omp_get_mapped_ptr(original_paa02 + 2, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, mapped_ptr_paa04 != original_paa02 + 2, &paa[0][4] == mapped_ptr_paa04);
+ }
+
+ int *original_paa020 = &paa[0][2][0];
+ int **original_paa0 = (int**) &paa[0];
+ // (H) use_device_addr/map: different base-pointers.
+ // No corresponding storage for use_device_addr opnd, lookup should fail.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+ {
+ int **mapped_ptr_paa020 = (int**) omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+ int **mapped_ptr_paa0 = (int**) omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ }
+
+ // (I) use_device_addr/map: one map with different, one with same base-ptr.
+ // Lookup should succeed.
+ // CHECK: I: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+ {
+ int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
new file mode 100644
index 0000000000000..ae61142827652
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on a variable (not a section).
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+ int *paa[10][10];
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ void *original_addr_g = &g;
+ void *original_addr_h = &h;
+ void *original_addr_ph = &ph;
+ void *original_addr_paa = &paa;
+
+ #pragma omp target enter data map(to:g, h, ph, paa)
+ void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
+ void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
+ void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
+ void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device());
+
+ // CHECK-COUNT-8: 1
+ printf("%d\n", mapped_ptr_g != nullptr);
+ printf("%d\n", mapped_ptr_h != nullptr);
+ printf("%d\n", mapped_ptr_ph != nullptr);
+ printf("%d\n", mapped_ptr_paa != nullptr);
+ printf("%d\n", original_addr_g != mapped_ptr_g);
+ printf("%d\n", original_addr_h != mapped_ptr_h);
+ printf("%d\n", original_addr_ph != mapped_ptr_ph);
+ printf("%d\n", original_addr_paa != mapped_ptr_paa);
+
+ // (A)
+ // CHECK: A: 1
+ #pragma omp target data use_device_addr(g)
+ printf("A: %d\n", mapped_ptr_g == &g);
+
+ // (B)
+ // CHECK: B: 1
+ #pragma omp target data use_device_addr(h)
+ printf("B: %d\n", mapped_ptr_h == &h);
+
+ // (C)
+ // CHECK: C: 1
+ #pragma omp target data use_device_addr(ph)
+ printf("C: %d\n", mapped_ptr_ph == &ph);
+
+ // (D) use_device_addr/map with different base-array/pointer.
+ // Address translation should happen for &ph, not &ph[0/1].
+ // CHECK: D: 1
+ #pragma omp target data map(ph[1:2]) use_device_addr(ph)
+ printf("D: %d\n", mapped_ptr_ph == &ph);
+
+ // (E)
+ // CHECK: E: 1
+ #pragma omp target data use_device_addr(paa)
+ printf("E: %d\n", mapped_ptr_paa == &paa);
+
+ // (F) use_device_addr/map with same base-array, paa.
+ // Address translation should happen for &paa.
+ // CHECK: F: 1
+ #pragma omp target data map(paa[0][2]) use_device_addr(paa)
+ printf("F: %d\n", mapped_ptr_paa == &paa);
+
+ // (G) use_device_addr/map with different base-array/pointer.
+ // Address translation should happen for &paa.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+ printf("G: %d\n", mapped_ptr_paa == &paa);
+
+ #pragma omp target exit data map(release:g, h, ph, paa)
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
new file mode 100644
index 0000000000000..5fadd36eb36b0
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
@@ -0,0 +1,137 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on a variable (not a section).
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g, h[10];
+int *ph = &h[0];
+
+struct S {
+ int *paa[10][10];
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ void *original_addr_g = &g;
+ void *original_addr_h = &h;
+ void *original_addr_ph = &ph;
+ void *original_addr_paa = &paa;
+
+ // (A) No corresponding item, lookup should fail.
+ // CHECK: A: 1 1 1
+ #pragma omp target data use_device_addr(g)
+ {
+ void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_g == nullptr, mapped_ptr_g != original_addr_g, (void*) &g == nullptr);
+ }
+
+ // (B) Lookup should succeed.
+ // CHECK: B: 1 1 1
+ #pragma omp target data map(g) use_device_addr(g)
+ {
+ void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_g != nullptr, mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+ }
+
+ // (C) No corresponding item, lookup should fail.
+ // CHECK: C: 1 1 1
+ #pragma omp target data use_device_addr(h)
+ {
+ void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_h == nullptr, mapped_ptr_h != original_addr_h, (void*) &h == nullptr);
+ }
+
+ // (D) Lookup should succeed.
+ // CHECK: D: 1 1 1
+ #pragma omp target data map(h) use_device_addr(h)
+ {
+ void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_h != nullptr, mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+ }
+
+ // (E) No corresponding item, lookup should fail.
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ }
+
+ // (F) Lookup should succeed.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(ph) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ }
+
+ // (G) Maps pointee only, but use_device_addr operand is pointer.
+ // Lookup should fail.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(ph[0:1]) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ }
+
+ // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(ph[0:1]) map(ph) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ }
+
+ // (I) No corresponding item, lookup should fail.
+ // CHECK: I: 1 1 1
+ #pragma omp target data use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ }
+
+ // (J) Maps pointee only, but use_device_addr operand is pointer.
+ // Lookup should fail.
+ // CHECK: J: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ }
+
+ // (K) Lookup should succeed.
+ // CHECK: K: 1 1 1
+ #pragma omp target data map(paa) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ }
+
+ // (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+ // CHECK: L: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
new file mode 100644
index 0000000000000..aad1afb265885
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
@@ -0,0 +1,102 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on a reference variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+ int *(&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ void *original_addr_g = &g;
+ void *original_addr_h = &h;
+ void *original_addr_ph = &ph;
+ void *original_addr_paa = &paa;
+
+ #pragma omp target enter data map(to:g, h, ph, paa)
+ void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
+ void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
+ void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
+ void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device());
+
+ // CHECK-COUNT-8: 1
+ printf("%d\n", mapped_ptr_g != nullptr);
+ printf("%d\n", mapped_ptr_h != nullptr);
+ printf("%d\n", mapped_ptr_ph != nullptr);
+ printf("%d\n", mapped_ptr_paa != nullptr);
+ printf("%d\n", original_addr_g != mapped_ptr_g);
+ printf("%d\n", original_addr_h != mapped_ptr_h);
+ printf("%d\n", original_addr_ph != mapped_ptr_ph);
+ printf("%d\n", original_addr_paa != mapped_ptr_paa);
+
+ // (A)
+ // CHECK: A: 1
+ #pragma omp target data use_device_addr(g)
+ printf("A: %d\n", mapped_ptr_g == &g);
+
+ // (B)
+ // CHECK: B: 1
+ #pragma omp target data use_device_addr(h)
+ printf("B: %d\n", mapped_ptr_h == &h);
+
+ // (C)
+ // CHECK: C: 1
+ #pragma omp target data use_device_addr(ph)
+ printf("C: %d\n", mapped_ptr_ph == &ph);
+
+ // (D) use_device_addr/map with different base-array/pointer.
+ // Address translation should happen for &ph, not &ph[0/1].
+ // CHECK: D: 1
+ #pragma omp target data map(ph[1:2]) use_device_addr(ph)
+ printf("D: %d\n", mapped_ptr_ph == &ph);
+
+ // (E)
+ // CHECK: E: 1
+ #pragma omp target data use_device_addr(paa)
+ printf("E: %d\n", mapped_ptr_paa == &paa);
+
+ // (F) use_device_addr/map with same base-array, paa.
+ // Address translation should happen for &paa.
+ // CHECK: F: 1
+ #pragma omp target data map(paa[0][2]) use_device_addr(paa)
+ printf("F: %d\n", mapped_ptr_paa == &paa);
+
+ // (G) use_device_addr/map with different base-array/pointer.
+ // Address translation should happen for &paa.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+ printf("G: %d\n", mapped_ptr_paa == &paa);
+
+ #pragma omp target exit data map(release:g, h, ph, paa)
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
new file mode 100644
index 0000000000000..6fcdd220d4f37
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
@@ -0,0 +1,144 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_addr on a reference variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int g_ptee;
+int &g = g_ptee;
+
+int h_ptee[10];
+int (&h)[10] = h_ptee;
+
+int *ph_ptee = &h_ptee[0];
+int *&ph = ph_ptee;
+int *paa_ptee[10][10];
+
+struct S {
+ int *(&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa[0][2] = &g;
+
+ void *original_addr_g = &g;
+ void *original_addr_h = &h;
+ void *original_addr_ph = &ph;
+ void *original_addr_paa = &paa;
+
+ // (A) No corresponding item, lookup should fail.
+ // CHECK: A: 1 1 1
+ #pragma omp target data use_device_addr(g)
+ {
+ void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_g == nullptr, mapped_ptr_g != original_addr_g, (void*) &g == nullptr);
+ }
+
+ // (B) Lookup should succeed.
+ // CHECK: B: 1 1 1
+ #pragma omp target data map(g) use_device_addr(g)
+ {
+ void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_g != nullptr, mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+ }
+
+ // (C) No corresponding item, lookup should fail.
+ // CHECK: C: 1 1 1
+ #pragma omp target data use_device_addr(h)
+ {
+ void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_h == nullptr, mapped_ptr_h != original_addr_h, (void*) &h == nullptr);
+ }
+
+ // (D) Lookup should succeed.
+ // CHECK: D: 1 1 1
+ #pragma omp target data map(h) use_device_addr(h)
+ {
+ void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_h != nullptr, mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+ }
+
+ // (E) No corresponding item, lookup should fail.
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ }
+
+ // (F) Lookup should succeed.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(ph) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ }
+
+ // (G) Maps pointee only, but use_device_addr operand is pointer.
+ // Lookup should fail.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(ph[0:1]) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ }
+
+ // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(ph[0:1]) map(ph) use_device_addr(ph)
+ {
+ void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ }
+
+ // (I) No corresponding item, lookup should fail.
+ // CHECK: I: 1 1 1
+ #pragma omp target data use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ }
+
+ // (J) Maps pointee only, but use_device_addr operand is pointer.
+ // Lookup should fail.
+ // CHECK: J: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ }
+
+ // (K) Lookup should succeed.
+ // CHECK: K: 1 1 1
+ #pragma omp target data map(paa) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ }
+
+ // (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+ // CHECK: L: 1 1 1
+ #pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+ {
+ void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c
similarity index 100%
rename from offload/test/mapping/target_use_device_addr.c
rename to offload/test/mapping/use_device_addr/target_use_device_addr.c
diff --git a/offload/test/mapping/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
similarity index 100%
rename from offload/test/mapping/target_wrong_use_device_addr.c
rename to offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
diff --git a/offload/test/mapping/array_section_use_device_ptr.c b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c
similarity index 100%
rename from offload/test/mapping/array_section_use_device_ptr.c
rename to offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
new file mode 100644
index 0000000000000..7cb7b57f1acf0
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
@@ -0,0 +1,102 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_ptr on a variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int h[10];
+int *ph = &h[0];
+
+struct S {
+ int (*paa)[10][10] = &aa;
+
+ void f1(int i) {
+ paa--;
+ void *original_ph3 = &ph[3];
+ void *original_paa102 = &paa[1][0][2];
+
+ #pragma omp target enter data map(to:ph[3:4], paa[1][0][2:5])
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+
+ // CHECK-COUNT-4: 1
+ printf("%d\n", mapped_ptr_ph3 != nullptr);
+ printf("%d\n", mapped_ptr_paa102 != nullptr);
+ printf("%d\n", original_ph3 != mapped_ptr_ph3);
+ printf("%d\n", original_paa102 != mapped_ptr_paa102);
+
+ // (A) Mapped data is within extended address range. Lookup should succeed.
+ // CHECK: A: 1
+ #pragma omp target data use_device_ptr(ph)
+ printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (B) use_device_ptr/map on pointer, and pointee already exists.
+ // Lookup should succeed.
+ // CHECK: B: 1
+ #pragma omp target data map(ph) use_device_ptr(ph)
+ printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: C: 1
+ #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+ printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (D) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: D: 1
+ #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+ printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (E) Mapped data is within extended address range. Lookup should succeed.
+ // Lookup should succeed.
+ // CHECK: E: 1
+ #pragma omp target data use_device_ptr(paa)
+ printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (F) use_device_ptr/map on pointer, and pointee already exists.
+ // &paa[0] should be in extended address-range of the existing paa[1][...]
+ // Lookup should succeed.
+ // FIXME: However, it currently does not. Might need an RT fix.
+ // EXPECTED: F: 1
+ // CHECK: F: 0
+ #pragma omp target data map(paa) use_device_ptr(paa)
+ printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+ printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (H) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: H: 1
+ #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+ printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+
+ #pragma omp target exit data map(release:ph[3:4], paa[1][0][2:5])
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
new file mode 100644
index 0000000000000..3b83c7f196784
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
@@ -0,0 +1,109 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_ptr on a variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int h[10];
+int *ph = &h[0];
+
+struct S {
+ int (*paa)[10][10] = &aa;
+
+ void f1(int i) {
+ paa--;
+ void *original_addr_ph3 = &ph[3];
+ void *original_addr_paa102 = &paa[1][0][2];
+
+ // (A) No corresponding item, lookup should fail.
+ // CHECK: A: 1 1 1
+ #pragma omp target data use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ }
+
+ // (B) use_device_ptr/map on pointer, and pointee does not exist.
+ // Lookup should fail.
+ // CHECK: B: 1 1 1
+ #pragma omp target data map(ph) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ }
+
+ // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: C: 1 1 1
+ #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (D) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: D: 1 1 1
+ #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (E) No corresponding item, lookup should fail.
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ }
+
+ // (F) use_device_ptr/map on pointer, and pointee does not exist.
+ // Lookup should fail.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(paa) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ }
+
+ // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ }
+
+ // (H) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
new file mode 100644
index 0000000000000..0d681d773c5a9
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
@@ -0,0 +1,113 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_ptr on a reference variable.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int (*paa_ptee)[10][10] = &aa;
+
+int h[10];
+int *ph_ptee = &h[0];
+int *&ph = ph_ptee;
+
+struct S {
+ int (*&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa--;
+ void *original_ph3 = &ph[3];
+ void *original_paa102 = &paa[1][0][2];
+
+ #pragma omp target enter data map(to:ph[3:4], paa[1][0][2:5])
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+
+ // CHECK-COUNT-4: 1
+ printf("%d\n", mapped_ptr_ph3 != nullptr);
+ printf("%d\n", mapped_ptr_paa102 != nullptr);
+ printf("%d\n", original_ph3 != mapped_ptr_ph3);
+ printf("%d\n", original_paa102 != mapped_ptr_paa102);
+
+ // (A) Mapped data is within extended address range. Lookup should succeed.
+ // EXPECTED: A: 1
+ // CHECK: A: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data use_device_ptr(ph)
+ printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (B) use_device_ptr/map on pointer, and pointee already exists.
+ // Lookup should succeed.
+ // EXPECTED: B: 1
+ // CHECK: B: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) use_device_ptr(ph)
+ printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // EXPECTED: C: 1
+ // CHECK: C: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+ printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (D) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // EXPECTED: D: 1
+ // CHECK: D: 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+ printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
+
+ // (E) Mapped data is within extended address range. Lookup should succeed.
+ // Lookup should succeed.
+ // CHECK: E: 1
+ #pragma omp target data use_device_ptr(paa)
+ printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (F) use_device_ptr/map on pointer, and pointee already exists.
+ // &paa[0] should be in extended address-range of the existing paa[1][...]
+ // Lookup should succeed.
+ // FIXME: However, it currently does not. Might need an RT fix.
+ // EXPECTED: F: 1
+ // CHECK: F: 0
+ #pragma omp target data map(paa) use_device_ptr(paa)
+ printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: G: 1
+ #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+ printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+ // (H) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: H: 1
+ #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+ printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
+
+
+ #pragma omp target exit data map(release:ph[3:4], paa[1][0][2:5])
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
new file mode 100644
index 0000000000000..141ccef52fb0b
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
@@ -0,0 +1,120 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <stdio.h>
+#include <omp.h>
+
+// Test for various cases of use_device_ptr on a reference variable.
+// The corresponding data is not previously mapped.
+
+// Note that this tests for the current behavior wherein if a lookup fails,
+// the runtime returns nullptr, instead of the original host-address.
+// That was compatible with OpenMP 5.0, where it was a user error if
+// corresponding storage didn't exist, but with 5.1+, the runtime needs to
+// return the host address, as it needs to assume that the host-address is
+// device-accessible, as the user has guaranteed it.
+// Once the runtime returns the original host-address when the lookup fails, the
+// test will need to be updated.
+
+int aa[10][10];
+int (*paa_ptee)[10][10] = &aa;
+
+int h[10];
+int *ph_ptee = &h[0];
+int *&ph = ph_ptee;
+
+struct S {
+ int (*&paa)[10][10] = paa_ptee;
+
+ void f1(int i) {
+ paa--;
+ void *original_addr_ph3 = &ph[3];
+ void *original_addr_paa102 = &paa[1][0][2];
+
+ // (A) No corresponding item, lookup should fail.
+ // EXPECTED: A: 1 1 1
+ // CHECK: A: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ }
+
+ // (B) use_device_ptr/map on pointer, and pointee does not exist.
+ // Lookup should fail.
+ // EXPECTED: B: 1 1 1
+ // CHECK: B: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ }
+
+ // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // EXPECTED: C: 1 1 1
+ // CHECK: C: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (D) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // EXPECTED: D: 1 1 1
+ // CHECK: D: 1 1 0
+ // FIXME: ph is not being privatized in the region.
+ #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+ {
+ void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ }
+
+ // (E) No corresponding item, lookup should fail.
+ // CHECK: E: 1 1 1
+ #pragma omp target data use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ }
+
+ // (F) use_device_ptr/map on pointer, and pointee does not exist.
+ // Lookup should fail.
+ // CHECK: F: 1 1 1
+ #pragma omp target data map(paa) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ }
+
+ // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: G: 1 1 1
+ #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ }
+
+ // (H) map on pointer and pointee. Base-pointer of map on pointee matches
+ // use_device_ptr operand.
+ // Lookup should succeed.
+ // CHECK: H: 1 1 1
+ #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+ {
+ void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ }
+ }
+};
+
+S s1;
+int main() {
+ s1.f1(1);
+}
>From 398a3a64892c827f3f331e9c0922a8f2ff3f3f8f Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 22 Aug 2025 06:06:34 -0700
Subject: [PATCH 2/3] Clang-format fixes
---
...t_data_use_device_addr_arrsec_existing.cpp | 60 +++----
...ta_use_device_addr_arrsec_not_existing.cpp | 141 +++++++++-------
...ta_use_device_addr_arrsec_ref_existing.cpp | 76 ++++-----
...se_device_addr_arrsec_ref_not_existing.cpp | 157 ++++++++++--------
...rget_data_use_device_addr_var_existing.cpp | 58 ++++---
..._data_use_device_addr_var_not_existing.cpp | 154 +++++++++--------
..._data_use_device_addr_var_ref_existing.cpp | 58 ++++---
...a_use_device_addr_var_ref_not_existing.cpp | 154 +++++++++--------
.../target_data_use_device_ptr_existing.cpp | 86 +++++-----
...arget_data_use_device_ptr_not_existing.cpp | 120 +++++++------
...arget_data_use_device_ptr_ref_existing.cpp | 102 ++++++------
...t_data_use_device_ptr_ref_not_existing.cpp | 136 ++++++++-------
12 files changed, 706 insertions(+), 596 deletions(-)
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
index eeb18ef57ca12..3b1a8192bf2cf 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on an array-section.
// The corresponding data is mapped on a previous enter_data directive.
@@ -29,9 +29,11 @@ struct S {
int *original_ph3 = &ph[3];
int **original_paa02 = &paa[0][2];
- #pragma omp target enter data map(to:ph[3:4], paa[0][2:5])
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(&ph[3], omp_get_default_device());
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
// CHECK-COUNT-4: 1
printf("%d\n", mapped_ptr_ph3 != nullptr);
@@ -39,47 +41,45 @@ struct S {
printf("%d\n", original_ph3 != mapped_ptr_ph3);
printf("%d\n", original_paa02 != mapped_ptr_paa02);
- // (A) use_device_addr operand within mapped address range.
- // CHECK: A: 1
- #pragma omp target data use_device_addr(ph[3:4])
+// (A) use_device_addr operand within mapped address range.
+// CHECK: A: 1
+#pragma omp target data use_device_addr(ph[3 : 4])
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (B) use_device_addr operand in extended address range, but not
- // mapped address range.
- // CHECK: B: 1
- #pragma omp target data use_device_addr(ph[2])
+// (B) use_device_addr operand in extended address range, but not
+// mapped address range.
+// CHECK: B: 1
+#pragma omp target data use_device_addr(ph[2])
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (C) use_device_addr/map: same base-array, different first-location.
- // CHECK: C: 1
- #pragma omp target data map(ph[3:2]) use_device_addr(ph[4:1])
+// (C) use_device_addr/map: same base-array, different first-location.
+// CHECK: C: 1
+#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (D) use_device_addr/map: different base-array/pointers.
- // CHECK: D: 1
- #pragma omp target data map(ph) use_device_addr(ph[3:4])
+// (D) use_device_addr/map: different base-array/pointers.
+// CHECK: D: 1
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (E) use_device_addr operand within mapped range of previous map.
- // CHECK: E: 1
- #pragma omp target data use_device_addr(paa[0])
+// (E) use_device_addr operand within mapped range of previous map.
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa[0])
printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- // (F) use_device_addr/map: different operands, same base-array.
- // CHECK: F: 1
- #pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+// (F) use_device_addr/map: different operands, same base-array.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- // (G) use_device_addr/map: different base-array/pointers.
- // CHECK: G: 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+// (G) use_device_addr/map: different base-array/pointers.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- #pragma omp target exit data map(release:ph[3:4], paa[0][2:5])
+#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
index 11543dffcce6e..22a31b9b0bd84 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on an array-section.
// The corresponding data is not previously mapped.
@@ -29,93 +29,114 @@ struct S {
int *original_ph3 = &ph[3];
int **original_paa02 = &paa[0][2];
- // (A) No corresponding map, lookup should fail.
- // CHECK: A: 1 1 1
- #pragma omp target data use_device_addr(ph[3:4])
+// (A) No corresponding map, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
}
- // (B) use_device_addr/map: different operands, same base-pointer.
- // use_device_addr operand within mapped address range.
- // CHECK: B: 1 1 1
- #pragma omp target data map(ph[2:3]) use_device_addr(ph[3:1])
+// (B) use_device_addr/map: different operands, same base-pointer.
+// use_device_addr operand within mapped address range.
+// CHECK: B: 1 1 1
+#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1])
{
- int *mapped_ptr_ph4 = (int*) omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+ int *mapped_ptr_ph4 =
+ (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr,
+ mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
}
- // (C) use_device_addr/map: different base-pointers.
- // No corresponding storage, lookup should fail.
- // CHECK: C: 1 1 1
- #pragma omp target data map(ph) use_device_addr(ph[3:4])
+// (C) use_device_addr/map: different base-pointers.
+// No corresponding storage, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
}
- // (D) use_device_addr/map: one of two maps with matching base-pointer.
- // use_device_addr operand within mapped address range of second map,
- // lookup should succeed.
- // CHECK: D: 1 1 1
- #pragma omp target data map(ph) map(ph[2:5]) use_device_addr(ph[3:4])
+// (D) use_device_addr/map: one of two maps with matching base-pointer.
+// use_device_addr operand within mapped address range of second map,
+// lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (E) No corresponding map, lookup should fail
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_addr(paa[0])
+// (E) No corresponding map, lookup should fail
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(paa[0])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == (int**) nullptr + 2);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == (int **)nullptr + 2);
}
- // (F) use_device_addr/map: different operands, same base-array.
- // use_device_addr within mapped address range. Lookup should succeed.
- // CHECK: F: 1 1 1
- #pragma omp target data map(paa) use_device_addr(paa[0])
+// (F) use_device_addr/map: different operands, same base-array.
+// use_device_addr within mapped address range. Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa[0])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == mapped_ptr_paa02);
}
- // (G) use_device_addr/map: different operands, same base-array.
- // use_device_addr extends beyond existing mapping. Not spec compliant.
- // But the lookup succeeds because we use the base-address for translation.
- // CHECK: G: 1 1 1
- #pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+// (G) use_device_addr/map: different operands, same base-array.
+// use_device_addr extends beyond existing mapping. Not spec compliant.
+// But the lookup succeeds because we use the base-address for translation.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
{
- int **mapped_ptr_paa04 = (int**) omp_get_mapped_ptr(original_paa02 + 2, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, mapped_ptr_paa04 != original_paa02 + 2, &paa[0][4] == mapped_ptr_paa04);
+ int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr(
+ original_paa02 + 2, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr,
+ mapped_ptr_paa04 != original_paa02 + 2,
+ &paa[0][4] == mapped_ptr_paa04);
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int**) &paa[0];
- // (H) use_device_addr/map: different base-pointers.
- // No corresponding storage for use_device_addr opnd, lookup should fail.
- // CHECK: H: 1 1 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+ int **original_paa0 = (int **)&paa[0];
+// (H) use_device_addr/map: different base-pointers.
+// No corresponding storage for use_device_addr opnd, lookup should fail.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
{
- int **mapped_ptr_paa020 = (int**) omp_get_mapped_ptr(original_paa020, omp_get_default_device());
- int **mapped_ptr_paa0 = (int**) omp_get_mapped_ptr(original_paa0, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ int **mapped_ptr_paa020 =
+ (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+ int **mapped_ptr_paa0 =
+ (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
+ mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
}
- // (I) use_device_addr/map: one map with different, one with same base-ptr.
- // Lookup should succeed.
- // CHECK: I: 1 1 1
- #pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+// (I) use_device_addr/map: one map with different, one with same base-ptr.
+// Lookup should succeed.
+// CHECK: I: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == mapped_ptr_paa02);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
index 502bcdf2ad2b8..e9a1124bc4612 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp
@@ -1,7 +1,7 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on an array-section on a reference.
// The corresponding data is mapped on a previous enter_data directive.
@@ -34,9 +34,11 @@ struct S {
int *original_ph3 = &ph[3];
int **original_paa02 = &paa[0][2];
- #pragma omp target enter data map(to:ph[3:4], paa[0][2:5])
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(&ph[3], omp_get_default_device());
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
+#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
// CHECK-COUNT-4: 1
printf("%d\n", mapped_ptr_ph3 != nullptr);
@@ -44,55 +46,53 @@ struct S {
printf("%d\n", original_ph3 != mapped_ptr_ph3);
printf("%d\n", original_paa02 != mapped_ptr_paa02);
- // (A) use_device_addr operand within mapped address range.
- // EXPECTED: A: 1
- // CHECK: A: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data use_device_addr(ph[3:4])
+// (A) use_device_addr operand within mapped address range.
+// EXPECTED: A: 1
+// CHECK: A: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[3 : 4])
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (B) use_device_addr operand in extended address range, but not
- // mapped address range.
- // EXPECTED: B: 1
- // CHECK: B: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data use_device_addr(ph[2])
+// (B) use_device_addr operand in extended address range, but not
+// mapped address range.
+// EXPECTED: B: 1
+// CHECK: B: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[2])
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (C) use_device_addr/map: same base-array, different first-location.
- // EXPECTED: C: 1
- // CHECK: C: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph[3:2]) use_device_addr(ph[4:1])
+// (C) use_device_addr/map: same base-array, different first-location.
+// EXPECTED: C: 1
+// CHECK: C: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (D) use_device_addr/map: different base-array/pointers.
- // EXPECTED: D: 1
- // CHECK: D: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) use_device_addr(ph[3:4])
+// (D) use_device_addr/map: different base-array/pointers.
+// EXPECTED: D: 1
+// CHECK: D: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (E) use_device_addr operand within mapped range of previous map.
- // CHECK: E: 1
- #pragma omp target data use_device_addr(paa[0])
+// (E) use_device_addr operand within mapped range of previous map.
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa[0])
printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- // (F) use_device_addr/map: different operands, same base-array.
- // CHECK: F: 1
- #pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
+// (F) use_device_addr/map: different operands, same base-array.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- // (G) use_device_addr/map: different base-array/pointers.
- // CHECK: G: 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
+// (G) use_device_addr/map: different base-array/pointers.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
- #pragma omp target exit data map(release:ph[3:4], paa[0][2:5])
+#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
index 18436dbee79a6..2bf803d7f5a6c 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on an array-section on a reference.
// The corresponding data is not previously mapped.
@@ -36,101 +36,122 @@ struct S {
int *original_ph3 = &ph[3];
int **original_paa02 = &paa[0][2];
- // (A) No corresponding map, lookup should fail.
- // EXPECTED: A: 1 1 1
- // CHECK: A: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data use_device_addr(ph[3:4])
+// (A) No corresponding map, lookup should fail.
+// EXPECTED: A: 1 1 1
+// CHECK: A: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
}
- // (B) use_device_addr/map: different operands, same base-pointer.
- // use_device_addr operand within mapped address range.
- // EXPECTED: B: 1 1 1
- // CHECK: B: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph[2:3]) use_device_addr(ph[3:1])
+// (B) use_device_addr/map: different operands, same base-pointer.
+// use_device_addr operand within mapped address range.
+// EXPECTED: B: 1 1 1
+// CHECK: B: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1])
{
- int *mapped_ptr_ph4 = (int*) omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
+ int *mapped_ptr_ph4 =
+ (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr,
+ mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
}
- // (C) use_device_addr/map: different base-pointers.
- // No corresponding storage, lookup should fail.
- // EXPECTED: C: 1 1 1
- // CHECK: C: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) use_device_addr(ph[3:4])
+// (C) use_device_addr/map: different base-pointers.
+// No corresponding storage, lookup should fail.
+// EXPECTED: C: 1 1 1
+// CHECK: C: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == (int*) nullptr + 3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
}
- // (D) use_device_addr/map: one of two maps with matching base-pointer.
- // use_device_addr operand within mapped address range of second map,
- // lookup should succeed.
- // EXPECTED: D: 1 1 1
- // CHECK: D: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) map(ph[2:5]) use_device_addr(ph[3:4])
+// (D) use_device_addr/map: one of two maps with matching base-pointer.
+// use_device_addr operand within mapped address range of second map,
+// lookup should succeed.
+// EXPECTED: D: 1 1 1
+// CHECK: D: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4])
{
- int *mapped_ptr_ph3 = (int*) omp_get_mapped_ptr(original_ph3, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
+ int *mapped_ptr_ph3 =
+ (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (E) No corresponding map, lookup should fail
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_addr(paa[0])
+// (E) No corresponding map, lookup should fail
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(paa[0])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == (int**) nullptr + 2);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == (int **)nullptr + 2);
}
- // (F) use_device_addr/map: different operands, same base-array.
- // use_device_addr within mapped address range. Lookup should succeed.
- // CHECK: F: 1 1 1
- #pragma omp target data map(paa) use_device_addr(paa[0])
+// (F) use_device_addr/map: different operands, same base-array.
+// use_device_addr within mapped address range. Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa[0])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == mapped_ptr_paa02);
}
- // (G) use_device_addr/map: different operands, same base-array.
- // use_device_addr extends beyond existing mapping. Not spec compliant.
- // But the lookup succeeds because we use the base-address for translation.
- // CHECK: G: 1 1 1
- #pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
+// (G) use_device_addr/map: different operands, same base-array.
+// use_device_addr extends beyond existing mapping. Not spec compliant.
+// But the lookup succeeds because we use the base-address for translation.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
{
- int **mapped_ptr_paa04 = (int**) omp_get_mapped_ptr(original_paa02 + 2, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, mapped_ptr_paa04 != original_paa02 + 2, &paa[0][4] == mapped_ptr_paa04);
+ int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr(
+ original_paa02 + 2, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr,
+ mapped_ptr_paa04 != original_paa02 + 2,
+ &paa[0][4] == mapped_ptr_paa04);
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int**) &paa[0];
- // (H) use_device_addr/map: different base-pointers.
- // No corresponding storage for use_device_addr opnd, lookup should fail.
- // CHECK: H: 1 1 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
+ int **original_paa0 = (int **)&paa[0];
+// (H) use_device_addr/map: different base-pointers.
+// No corresponding storage for use_device_addr opnd, lookup should fail.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
{
- int **mapped_ptr_paa020 = (int**) omp_get_mapped_ptr(original_paa020, omp_get_default_device());
- int **mapped_ptr_paa0 = (int**) omp_get_mapped_ptr(original_paa0, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ int **mapped_ptr_paa020 =
+ (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device());
+ int **mapped_ptr_paa0 =
+ (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
+ mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
}
- // (I) use_device_addr/map: one map with different, one with same base-ptr.
- // Lookup should succeed.
- // CHECK: I: 1 1 1
- #pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
+// (I) use_device_addr/map: one map with different, one with same base-ptr.
+// Lookup should succeed.
+// CHECK: I: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
{
- int **mapped_ptr_paa02 = (int**) omp_get_mapped_ptr(original_paa02, omp_get_default_device());
- printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, mapped_ptr_paa02 != original_paa02, &paa[0][2] == mapped_ptr_paa02);
+ int **mapped_ptr_paa02 =
+ (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr,
+ mapped_ptr_paa02 != original_paa02,
+ &paa[0][2] == mapped_ptr_paa02);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
index ae61142827652..883297f7e90cd 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on a variable (not a section).
// The corresponding data is mapped on a previous enter_data directive.
@@ -31,7 +31,7 @@ struct S {
void *original_addr_ph = &ph;
void *original_addr_paa = &paa;
- #pragma omp target enter data map(to:g, h, ph, paa)
+#pragma omp target enter data map(to : g, h, ph, paa)
void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
@@ -47,49 +47,47 @@ struct S {
printf("%d\n", original_addr_ph != mapped_ptr_ph);
printf("%d\n", original_addr_paa != mapped_ptr_paa);
- // (A)
- // CHECK: A: 1
- #pragma omp target data use_device_addr(g)
+// (A)
+// CHECK: A: 1
+#pragma omp target data use_device_addr(g)
printf("A: %d\n", mapped_ptr_g == &g);
- // (B)
- // CHECK: B: 1
- #pragma omp target data use_device_addr(h)
+// (B)
+// CHECK: B: 1
+#pragma omp target data use_device_addr(h)
printf("B: %d\n", mapped_ptr_h == &h);
- // (C)
- // CHECK: C: 1
- #pragma omp target data use_device_addr(ph)
+// (C)
+// CHECK: C: 1
+#pragma omp target data use_device_addr(ph)
printf("C: %d\n", mapped_ptr_ph == &ph);
- // (D) use_device_addr/map with different base-array/pointer.
- // Address translation should happen for &ph, not &ph[0/1].
- // CHECK: D: 1
- #pragma omp target data map(ph[1:2]) use_device_addr(ph)
+// (D) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &ph, not &ph[0/1].
+// CHECK: D: 1
+#pragma omp target data map(ph[1 : 2]) use_device_addr(ph)
printf("D: %d\n", mapped_ptr_ph == &ph);
- // (E)
- // CHECK: E: 1
- #pragma omp target data use_device_addr(paa)
+// (E)
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa)
printf("E: %d\n", mapped_ptr_paa == &paa);
- // (F) use_device_addr/map with same base-array, paa.
- // Address translation should happen for &paa.
- // CHECK: F: 1
- #pragma omp target data map(paa[0][2]) use_device_addr(paa)
+// (F) use_device_addr/map with same base-array, paa.
+// Address translation should happen for &paa.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][2]) use_device_addr(paa)
printf("F: %d\n", mapped_ptr_paa == &paa);
- // (G) use_device_addr/map with different base-array/pointer.
- // Address translation should happen for &paa.
- // CHECK: G: 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+// (G) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &paa.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
printf("G: %d\n", mapped_ptr_paa == &paa);
- #pragma omp target exit data map(release:g, h, ph, paa)
+#pragma omp target exit data map(release : g, h, ph, paa)
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
index 5fadd36eb36b0..79c6f69edba8e 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on a variable (not a section).
// The corresponding data is not previously mapped.
@@ -31,107 +31,129 @@ struct S {
void *original_addr_ph = &ph;
void *original_addr_paa = &paa;
- // (A) No corresponding item, lookup should fail.
- // CHECK: A: 1 1 1
- #pragma omp target data use_device_addr(g)
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(g)
{
- void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_g == nullptr, mapped_ptr_g != original_addr_g, (void*) &g == nullptr);
+ void *mapped_ptr_g =
+ omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
+ mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
}
- // (B) Lookup should succeed.
- // CHECK: B: 1 1 1
- #pragma omp target data map(g) use_device_addr(g)
+// (B) Lookup should succeed.
+// CHECK: B: 1 1 1
+#pragma omp target data map(g) use_device_addr(g)
{
- void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_g != nullptr, mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+ void *mapped_ptr_g =
+ omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_g != nullptr,
+ mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
}
- // (C) No corresponding item, lookup should fail.
- // CHECK: C: 1 1 1
- #pragma omp target data use_device_addr(h)
+// (C) No corresponding item, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data use_device_addr(h)
{
- void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_h == nullptr, mapped_ptr_h != original_addr_h, (void*) &h == nullptr);
+ void *mapped_ptr_h =
+ omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
+ mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
}
- // (D) Lookup should succeed.
- // CHECK: D: 1 1 1
- #pragma omp target data map(h) use_device_addr(h)
+// (D) Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(h) use_device_addr(h)
{
- void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_h != nullptr, mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+ void *mapped_ptr_h =
+ omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_h != nullptr,
+ mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
}
- // (E) No corresponding item, lookup should fail.
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_addr(ph)
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
+ mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
}
- // (F) Lookup should succeed.
- // CHECK: F: 1 1 1
- #pragma omp target data map(ph) use_device_addr(ph)
+// (F) Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_ph != nullptr,
+ mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
}
- // (G) Maps pointee only, but use_device_addr operand is pointer.
- // Lookup should fail.
- // CHECK: G: 1 1 1
- #pragma omp target data map(ph[0:1]) use_device_addr(ph)
+// (G) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: G: 1 1 1
+#pragma omp target data map(ph[0 : 1]) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
+ mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
}
- // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
- // CHECK: H: 1 1 1
- #pragma omp target data map(ph[0:1]) map(ph) use_device_addr(ph)
+// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_ph != nullptr,
+ mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
}
- // (I) No corresponding item, lookup should fail.
- // CHECK: I: 1 1 1
- #pragma omp target data use_device_addr(paa)
+// (I) No corresponding item, lookup should fail.
+// CHECK: I: 1 1 1
+#pragma omp target data use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
+ mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
}
- // (J) Maps pointee only, but use_device_addr operand is pointer.
- // Lookup should fail.
- // CHECK: J: 1 1 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+// (J) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: J: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
+ mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
}
- // (K) Lookup should succeed.
- // CHECK: K: 1 1 1
- #pragma omp target data map(paa) use_device_addr(paa)
+// (K) Lookup should succeed.
+// CHECK: K: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("K: %d %d %d\n", mapped_ptr_paa != nullptr,
+ mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
}
- // (L) Maps both pointee and pointer. Lookup for pointer should succeed.
- // CHECK: L: 1 1 1
- #pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+// (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: L: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("L: %d %d %d\n", mapped_ptr_paa != nullptr,
+ mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
index aad1afb265885..f018c65f36ec5 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on a reference variable.
// The corresponding data is mapped on a previous enter_data directive.
@@ -38,7 +38,7 @@ struct S {
void *original_addr_ph = &ph;
void *original_addr_paa = &paa;
- #pragma omp target enter data map(to:g, h, ph, paa)
+#pragma omp target enter data map(to : g, h, ph, paa)
void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device());
void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device());
void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device());
@@ -54,49 +54,47 @@ struct S {
printf("%d\n", original_addr_ph != mapped_ptr_ph);
printf("%d\n", original_addr_paa != mapped_ptr_paa);
- // (A)
- // CHECK: A: 1
- #pragma omp target data use_device_addr(g)
+// (A)
+// CHECK: A: 1
+#pragma omp target data use_device_addr(g)
printf("A: %d\n", mapped_ptr_g == &g);
- // (B)
- // CHECK: B: 1
- #pragma omp target data use_device_addr(h)
+// (B)
+// CHECK: B: 1
+#pragma omp target data use_device_addr(h)
printf("B: %d\n", mapped_ptr_h == &h);
- // (C)
- // CHECK: C: 1
- #pragma omp target data use_device_addr(ph)
+// (C)
+// CHECK: C: 1
+#pragma omp target data use_device_addr(ph)
printf("C: %d\n", mapped_ptr_ph == &ph);
- // (D) use_device_addr/map with different base-array/pointer.
- // Address translation should happen for &ph, not &ph[0/1].
- // CHECK: D: 1
- #pragma omp target data map(ph[1:2]) use_device_addr(ph)
+// (D) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &ph, not &ph[0/1].
+// CHECK: D: 1
+#pragma omp target data map(ph[1 : 2]) use_device_addr(ph)
printf("D: %d\n", mapped_ptr_ph == &ph);
- // (E)
- // CHECK: E: 1
- #pragma omp target data use_device_addr(paa)
+// (E)
+// CHECK: E: 1
+#pragma omp target data use_device_addr(paa)
printf("E: %d\n", mapped_ptr_paa == &paa);
- // (F) use_device_addr/map with same base-array, paa.
- // Address translation should happen for &paa.
- // CHECK: F: 1
- #pragma omp target data map(paa[0][2]) use_device_addr(paa)
+// (F) use_device_addr/map with same base-array, paa.
+// Address translation should happen for &paa.
+// CHECK: F: 1
+#pragma omp target data map(paa[0][2]) use_device_addr(paa)
printf("F: %d\n", mapped_ptr_paa == &paa);
- // (G) use_device_addr/map with different base-array/pointer.
- // Address translation should happen for &paa.
- // CHECK: G: 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+// (G) use_device_addr/map with different base-array/pointer.
+// Address translation should happen for &paa.
+// CHECK: G: 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
printf("G: %d\n", mapped_ptr_paa == &paa);
- #pragma omp target exit data map(release:g, h, ph, paa)
+#pragma omp target exit data map(release : g, h, ph, paa)
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
index 6fcdd220d4f37..9360db4195041 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_addr on a reference variable.
// The corresponding data is not previously mapped.
@@ -38,107 +38,129 @@ struct S {
void *original_addr_ph = &ph;
void *original_addr_paa = &paa;
- // (A) No corresponding item, lookup should fail.
- // CHECK: A: 1 1 1
- #pragma omp target data use_device_addr(g)
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_addr(g)
{
- void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_g == nullptr, mapped_ptr_g != original_addr_g, (void*) &g == nullptr);
+ void *mapped_ptr_g =
+ omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
+ mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
}
- // (B) Lookup should succeed.
- // CHECK: B: 1 1 1
- #pragma omp target data map(g) use_device_addr(g)
+// (B) Lookup should succeed.
+// CHECK: B: 1 1 1
+#pragma omp target data map(g) use_device_addr(g)
{
- void *mapped_ptr_g = omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_g != nullptr, mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
+ void *mapped_ptr_g =
+ omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_g != nullptr,
+ mapped_ptr_g != original_addr_g, &g == mapped_ptr_g);
}
- // (C) No corresponding item, lookup should fail.
- // CHECK: C: 1 1 1
- #pragma omp target data use_device_addr(h)
+// (C) No corresponding item, lookup should fail.
+// CHECK: C: 1 1 1
+#pragma omp target data use_device_addr(h)
{
- void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_h == nullptr, mapped_ptr_h != original_addr_h, (void*) &h == nullptr);
+ void *mapped_ptr_h =
+ omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
+ mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
}
- // (D) Lookup should succeed.
- // CHECK: D: 1 1 1
- #pragma omp target data map(h) use_device_addr(h)
+// (D) Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(h) use_device_addr(h)
{
- void *mapped_ptr_h = omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_h != nullptr, mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
+ void *mapped_ptr_h =
+ omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_h != nullptr,
+ mapped_ptr_h != original_addr_h, &h == mapped_ptr_h);
}
- // (E) No corresponding item, lookup should fail.
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_addr(ph)
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
+ mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
}
- // (F) Lookup should succeed.
- // CHECK: F: 1 1 1
- #pragma omp target data map(ph) use_device_addr(ph)
+// (F) Lookup should succeed.
+// CHECK: F: 1 1 1
+#pragma omp target data map(ph) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_ph != nullptr,
+ mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
}
- // (G) Maps pointee only, but use_device_addr operand is pointer.
- // Lookup should fail.
- // CHECK: G: 1 1 1
- #pragma omp target data map(ph[0:1]) use_device_addr(ph)
+// (G) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: G: 1 1 1
+#pragma omp target data map(ph[0 : 1]) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, mapped_ptr_ph != original_addr_ph, (void*) &ph == nullptr);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
+ mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
}
- // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
- // CHECK: H: 1 1 1
- #pragma omp target data map(ph[0:1]) map(ph) use_device_addr(ph)
+// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph)
{
- void *mapped_ptr_ph = omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
+ void *mapped_ptr_ph =
+ omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_ph != nullptr,
+ mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph);
}
- // (I) No corresponding item, lookup should fail.
- // CHECK: I: 1 1 1
- #pragma omp target data use_device_addr(paa)
+// (I) No corresponding item, lookup should fail.
+// CHECK: I: 1 1 1
+#pragma omp target data use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
+ mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
}
- // (J) Maps pointee only, but use_device_addr operand is pointer.
- // Lookup should fail.
- // CHECK: J: 1 1 1
- #pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
+// (J) Maps pointee only, but use_device_addr operand is pointer.
+// Lookup should fail.
+// CHECK: J: 1 1 1
+#pragma omp target data map(paa[0][2][0]) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, mapped_ptr_paa != original_addr_paa, (void*) &paa == nullptr);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
+ mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
}
- // (K) Lookup should succeed.
- // CHECK: K: 1 1 1
- #pragma omp target data map(paa) use_device_addr(paa)
+// (K) Lookup should succeed.
+// CHECK: K: 1 1 1
+#pragma omp target data map(paa) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("K: %d %d %d\n", mapped_ptr_paa != nullptr,
+ mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
}
- // (L) Maps both pointee and pointer. Lookup for pointer should succeed.
- // CHECK: L: 1 1 1
- #pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
+// (L) Maps both pointee and pointer. Lookup for pointer should succeed.
+// CHECK: L: 1 1 1
+#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa)
{
- void *mapped_ptr_paa = omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
- printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
+ void *mapped_ptr_paa =
+ omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
+ printf("L: %d %d %d\n", mapped_ptr_paa != nullptr,
+ mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
index 7cb7b57f1acf0..a7745de53298e 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_ptr on a variable.
// The corresponding data is mapped on a previous enter_data directive.
@@ -29,9 +29,10 @@ struct S {
void *original_ph3 = &ph[3];
void *original_paa102 = &paa[1][0][2];
- #pragma omp target enter data map(to:ph[3:4], paa[1][0][2:5])
+#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5])
void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
// CHECK-COUNT-4: 1
printf("%d\n", mapped_ptr_ph3 != nullptr);
@@ -39,64 +40,61 @@ struct S {
printf("%d\n", original_ph3 != mapped_ptr_ph3);
printf("%d\n", original_paa102 != mapped_ptr_paa102);
- // (A) Mapped data is within extended address range. Lookup should succeed.
- // CHECK: A: 1
- #pragma omp target data use_device_ptr(ph)
+// (A) Mapped data is within extended address range. Lookup should succeed.
+// CHECK: A: 1
+#pragma omp target data use_device_ptr(ph)
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (B) use_device_ptr/map on pointer, and pointee already exists.
- // Lookup should succeed.
- // CHECK: B: 1
- #pragma omp target data map(ph) use_device_ptr(ph)
+// (B) use_device_ptr/map on pointer, and pointee already exists.
+// Lookup should succeed.
+// CHECK: B: 1
+#pragma omp target data map(ph) use_device_ptr(ph)
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: C: 1
- #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: C: 1
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (D) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: D: 1
- #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: D: 1
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (E) Mapped data is within extended address range. Lookup should succeed.
- // Lookup should succeed.
- // CHECK: E: 1
- #pragma omp target data use_device_ptr(paa)
+// (E) Mapped data is within extended address range. Lookup should succeed.
+// Lookup should succeed.
+// CHECK: E: 1
+#pragma omp target data use_device_ptr(paa)
printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (F) use_device_ptr/map on pointer, and pointee already exists.
- // &paa[0] should be in extended address-range of the existing paa[1][...]
- // Lookup should succeed.
- // FIXME: However, it currently does not. Might need an RT fix.
- // EXPECTED: F: 1
- // CHECK: F: 0
- #pragma omp target data map(paa) use_device_ptr(paa)
+// (F) use_device_ptr/map on pointer, and pointee already exists.
+// &paa[0] should be in extended address-range of the existing paa[1][...]
+// Lookup should succeed.
+// FIXME: However, it currently does not. Might need an RT fix.
+// EXPECTED: F: 1
+// CHECK: F: 0
+#pragma omp target data map(paa) use_device_ptr(paa)
printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: G: 1
- #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (H) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: H: 1
- #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
-
- #pragma omp target exit data map(release:ph[3:4], paa[1][0][2:5])
+#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5])
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
index 3b83c7f196784..fe3cdb56e4baa 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_ptr on a variable.
// The corresponding data is not previously mapped.
@@ -29,81 +29,97 @@ struct S {
void *original_addr_ph3 = &ph[3];
void *original_addr_paa102 = &paa[1][0][2];
- // (A) No corresponding item, lookup should fail.
- // CHECK: A: 1 1 1
- #pragma omp target data use_device_ptr(ph)
+// (A) No corresponding item, lookup should fail.
+// CHECK: A: 1 1 1
+#pragma omp target data use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
}
- // (B) use_device_ptr/map on pointer, and pointee does not exist.
- // Lookup should fail.
- // CHECK: B: 1 1 1
- #pragma omp target data map(ph) use_device_ptr(ph)
+// (B) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: B: 1 1 1
+#pragma omp target data map(ph) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
}
- // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: C: 1 1 1
- #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: C: 1 1 1
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (D) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: D: 1 1 1
- #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: D: 1 1 1
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (E) No corresponding item, lookup should fail.
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_ptr(paa)
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+ mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
}
- // (F) use_device_ptr/map on pointer, and pointee does not exist.
- // Lookup should fail.
- // CHECK: F: 1 1 1
- #pragma omp target data map(paa) use_device_ptr(paa)
+// (F) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+ mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
}
- // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: G: 1 1 1
- #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+ mapped_ptr_paa102 != original_addr_paa102,
+ &paa[1][0][2] == mapped_ptr_paa102);
}
- // (H) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: H: 1 1 1
- #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+ mapped_ptr_paa102 != original_addr_paa102,
+ &paa[1][0][2] == mapped_ptr_paa102);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
index 0d681d773c5a9..66e65de4195a4 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_ptr on a reference variable.
// The corresponding data is mapped on a previous enter_data directive.
@@ -32,9 +32,10 @@ struct S {
void *original_ph3 = &ph[3];
void *original_paa102 = &paa[1][0][2];
- #pragma omp target enter data map(to:ph[3:4], paa[1][0][2:5])
+#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5])
void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device());
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device());
// CHECK-COUNT-4: 1
printf("%d\n", mapped_ptr_ph3 != nullptr);
@@ -42,72 +43,69 @@ struct S {
printf("%d\n", original_ph3 != mapped_ptr_ph3);
printf("%d\n", original_paa102 != mapped_ptr_paa102);
- // (A) Mapped data is within extended address range. Lookup should succeed.
- // EXPECTED: A: 1
- // CHECK: A: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data use_device_ptr(ph)
+// (A) Mapped data is within extended address range. Lookup should succeed.
+// EXPECTED: A: 1
+// CHECK: A: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_ptr(ph)
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (B) use_device_ptr/map on pointer, and pointee already exists.
- // Lookup should succeed.
- // EXPECTED: B: 1
- // CHECK: B: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) use_device_ptr(ph)
+// (B) use_device_ptr/map on pointer, and pointee already exists.
+// Lookup should succeed.
+// EXPECTED: B: 1
+// CHECK: B: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_ptr(ph)
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // EXPECTED: C: 1
- // CHECK: C: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: C: 1
+// CHECK: C: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (D) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // EXPECTED: D: 1
- // CHECK: D: 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: D: 1
+// CHECK: D: 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
- // (E) Mapped data is within extended address range. Lookup should succeed.
- // Lookup should succeed.
- // CHECK: E: 1
- #pragma omp target data use_device_ptr(paa)
+// (E) Mapped data is within extended address range. Lookup should succeed.
+// Lookup should succeed.
+// CHECK: E: 1
+#pragma omp target data use_device_ptr(paa)
printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (F) use_device_ptr/map on pointer, and pointee already exists.
- // &paa[0] should be in extended address-range of the existing paa[1][...]
- // Lookup should succeed.
- // FIXME: However, it currently does not. Might need an RT fix.
- // EXPECTED: F: 1
- // CHECK: F: 0
- #pragma omp target data map(paa) use_device_ptr(paa)
+// (F) use_device_ptr/map on pointer, and pointee already exists.
+// &paa[0] should be in extended address-range of the existing paa[1][...]
+// Lookup should succeed.
+// FIXME: However, it currently does not. Might need an RT fix.
+// EXPECTED: F: 1
+// CHECK: F: 0
+#pragma omp target data map(paa) use_device_ptr(paa)
printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: G: 1
- #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
- // (H) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: H: 1
- #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]);
-
- #pragma omp target exit data map(release:ph[3:4], paa[1][0][2:5])
+#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5])
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
index 141ccef52fb0b..419ab3eb33d4d 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
@@ -2,8 +2,8 @@
// XFAIL: *
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// Test for various cases of use_device_ptr on a reference variable.
// The corresponding data is not previously mapped.
@@ -32,89 +32,105 @@ struct S {
void *original_addr_ph3 = &ph[3];
void *original_addr_paa102 = &paa[1][0][2];
- // (A) No corresponding item, lookup should fail.
- // EXPECTED: A: 1 1 1
- // CHECK: A: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data use_device_ptr(ph)
+// (A) No corresponding item, lookup should fail.
+// EXPECTED: A: 1 1 1
+// CHECK: A: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
}
- // (B) use_device_ptr/map on pointer, and pointee does not exist.
- // Lookup should fail.
- // EXPECTED: B: 1 1 1
- // CHECK: B: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) use_device_ptr(ph)
+// (B) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// EXPECTED: B: 1 1 1
+// CHECK: B: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
}
- // (C) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // EXPECTED: C: 1 1 1
- // CHECK: C: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph[3:2]) use_device_ptr(ph)
+// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: C: 1 1 1
+// CHECK: C: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (D) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // EXPECTED: D: 1 1 1
- // CHECK: D: 1 1 0
- // FIXME: ph is not being privatized in the region.
- #pragma omp target data map(ph) map(ph[3:2]) use_device_ptr(ph)
+// (D) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// EXPECTED: D: 1 1 1
+// CHECK: D: 1 1 0
+// FIXME: ph is not being privatized in the region.
+#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph)
{
- void *mapped_ptr_ph3 = omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
- printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
+ void *mapped_ptr_ph3 =
+ omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
+ printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
+ mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3);
}
- // (E) No corresponding item, lookup should fail.
- // CHECK: E: 1 1 1
- #pragma omp target data use_device_ptr(paa)
+// (E) No corresponding item, lookup should fail.
+// CHECK: E: 1 1 1
+#pragma omp target data use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+ mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
}
- // (F) use_device_ptr/map on pointer, and pointee does not exist.
- // Lookup should fail.
- // CHECK: F: 1 1 1
- #pragma omp target data map(paa) use_device_ptr(paa)
+// (F) use_device_ptr/map on pointer, and pointee does not exist.
+// Lookup should fail.
+// CHECK: F: 1 1 1
+#pragma omp target data map(paa) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
+ mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
}
- // (G) map on pointee: base-pointer of map matches use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: G: 1 1 1
- #pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
+// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: G: 1 1 1
+#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+ mapped_ptr_paa102 != original_addr_paa102,
+ &paa[1][0][2] == mapped_ptr_paa102);
}
- // (H) map on pointer and pointee. Base-pointer of map on pointee matches
- // use_device_ptr operand.
- // Lookup should succeed.
- // CHECK: H: 1 1 1
- #pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
+// (H) map on pointer and pointee. Base-pointer of map on pointee matches
+// use_device_ptr operand.
+// Lookup should succeed.
+// CHECK: H: 1 1 1
+#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa)
{
- void *mapped_ptr_paa102 = omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
- printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, mapped_ptr_paa102 != original_addr_paa102, &paa[1][0][2] == mapped_ptr_paa102);
+ void *mapped_ptr_paa102 =
+ omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
+ printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr,
+ mapped_ptr_paa102 != original_addr_paa102,
+ &paa[1][0][2] == mapped_ptr_paa102);
}
}
};
S s1;
-int main() {
- s1.f1(1);
-}
+int main() { s1.f1(1); }
>From 3e4d99bb74cb4998ba68b11da931ce36b19e5fff Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 22 Aug 2025 06:13:10 -0700
Subject: [PATCH 3/3] Add two empty lines.
---
.../target_data_use_device_addr_arrsec_not_existing.cpp | 1 +
.../target_data_use_device_addr_arrsec_ref_not_existing.cpp | 1 +
2 files changed, 2 insertions(+)
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
index 22a31b9b0bd84..b9ebde431e7bf 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -111,6 +111,7 @@ struct S {
int *original_paa020 = &paa[0][2][0];
int **original_paa0 = (int **)&paa[0];
+
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
// CHECK: H: 1 1 1
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
index 2bf803d7f5a6c..0090cdb095366 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -126,6 +126,7 @@ struct S {
int *original_paa020 = &paa[0][2][0];
int **original_paa0 = (int **)&paa[0];
+
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
// CHECK: H: 1 1 1
More information about the llvm-commits
mailing list