From b86bd0addcf626eac8a4b8922916926207e6ec01 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba Date: Fri, 22 Aug 2025 04:22:53 -0700 Subject: [PATCH 1/4] [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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 +#include + +// 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 Date: Fri, 22 Aug 2025 06:06:34 -0700 Subject: [PATCH 2/4] 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 #include +#include // 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 Date: Fri, 22 Aug 2025 06:13:10 -0700 Subject: [PATCH 3/4] 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 From a49143a926c2c42704fa053f0c86236b020c192e Mon Sep 17 00:00:00 2001 From: Abhinav Gaba Date: Fri, 22 Aug 2025 06:14:08 -0700 Subject: [PATCH 4/4] Clang-format renamed files as well. --- offload/test/mapping/use_device_addr/target_use_device_addr.c | 4 +++- .../mapping/use_device_addr/target_wrong_use_device_addr.c | 3 +-- .../mapping/use_device_ptr/array_section_use_device_ptr.c | 4 +++- 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/offload/test/mapping/use_device_addr/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c index 5c2bb8a48f6e6..4a9dbe252f761 100644 --- a/offload/test/mapping/use_device_addr/target_use_device_addr.c +++ b/offload/test/mapping/use_device_addr/target_use_device_addr.c @@ -12,7 +12,9 @@ int main() { printf("%d, %p\n", xp[1], &xp[1]); #pragma omp target data use_device_addr(xp[1 : 3]) map(tofrom : x) #pragma omp target is_device_ptr(xp) - { xp[1] = 222; } + { + xp[1] = 222; + } // CHECK: 222 printf("%d, %p\n", xp[1], &xp[1]); } diff --git a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c index 7a5babd692530..28ec6857fa1a8 100644 --- a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c +++ b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c @@ -14,7 +14,7 @@ int main() { // CHECK: host addr=0x[[#%x,HOST_ADDR:]] fprintf(stderr, "host addr=%p\n", x); -#pragma omp target data map(to : x [0:10]) +#pragma omp target data map(to : x[0 : 10]) { // CHECK: omptarget device 0 info: variable x does not have a valid device // counterpart @@ -27,4 +27,3 @@ int main() { return 0; } - diff --git a/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c index 86e2875c35c4a..4cfcce28c1127 100644 --- a/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c +++ b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c @@ -20,7 +20,9 @@ int main() { float *A_dev = NULL; #pragma omp target data use_device_ptr(A) - { A_dev = A; } + { + A_dev = A; + } #pragma omp target exit data map(delete : A[FROM : LENGTH]) // CHECK: Success