[NFC][OpenMP] Add several use_device_ptr/addr tests. #154939
Open
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
abhinavgaba
changed the title
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
abhinavgaba
commented
Aug 22, 2025
| // 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) |
There was a problem hiding this comment.
@dreachem, please check if the expectation that this lookup should succeed, is valid. The pointee is mapped on the outer target_enter_data.
|
@llvm/pr-subscribers-offload Author: Abhinav Gaba (abhinavgaba) ChangesMost tests are either compfailing or runfailing. They should start passing once we start using ATTACH map-type based codegen. (#153683) Even after they start passing, there are a few places where the EXPECTED and actual CHECKs are different, due to two main issues:
The above should be fixed as separate standalone changes. Patch is 57.97 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/154939.diff 15 Files Affected:
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..3b1a8192bf2cf
--- /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 <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// 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..b9ebde431e7bf
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -0,0 +1,143 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section.
+// The corresponding data is not previously mapped.
+
+// 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..e9a1124bc4612
--- /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 <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is mapped on a previous enter_data directive.
+
+// 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..0090cdb095366
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -0,0 +1,158 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on an array-section on a reference.
+// The corresponding data is not previously mapped.
+
+// 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..883297f7e90cd
--- /dev/null
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp
@@ -0,0 +1,93 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// XFAIL: *
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test for various cases of use_device_addr on a variable (not a section).
+// The corresponding data is mapped on a previous enter_data directive.
+
+// 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
+//...
[truncated]
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Most tests are either compfailing or runfailing.
They should start passing once we start using ATTACH map-type based codegen. (#153683)
Even after they start passing, there are a few places where the EXPECTED and actual CHECKs are different, due to two main issues:
&p[0]is not succeeding in looking-up a previously mapped&p[1]The above should be fixed as separate standalone changes.