llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Abhinav Gaba (abhinavgaba)
<details>
<summary>Changes</summary>
As per OpenMP 5.1, we need to assume that when the lookup for
`use_device_ptr/addr` fails, the incoming pointer was already device
accessible.
Prior to 5.1, a lookup-failure meant a user-error (for `use_device_ptr`),
so we could do anything in that scenario. For `use_device_addr`,
it was always incorrect to set the address to null.
OpenMP 6.1 adds a way to retain the previous behavior of nullifying a pointer
when the lookup fails. That will be tackled by the PR stack
starting with https://github.com/llvm/llvm-project/pull/169603.
---
Patch is 27.66 KiB, truncated to 20.00 KiB below, full version:
https://github.com/llvm/llvm-project/pull/174659.diff
13 Files Affected:
- (modified) clang/docs/OpenMPSupport.rst (+2)
- (modified) clang/docs/ReleaseNotes.rst (+2)
- (modified) offload/libomptarget/omptarget.cpp (+33-3)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
(-2)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
(+5-15)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
(+7-21)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
(-2)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
(+6-15)
- (modified)
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
(+6-15)
- (modified)
offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c (+3-2)
- (modified)
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
(+6-13)
- (modified)
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
(+8-19)
- (modified)
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
(-11)
``````````diff
diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index ab3f2c48983ca..7941c2e439ed6 100644
--- a/clang/docs/OpenMPSupport.rst
+++ b/clang/docs/OpenMPSupport.rst
@@ -266,6 +266,8 @@ implementation.
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | has_device_addr clause on target construct
| :none:`unclaimed` |
|
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
+| device | use_device_ptr/addr preserve host address
when lookup fails | :good:`done` |
https://github.com/llvm/llvm-project/pull/174659 |
++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | iterators in map clause or motion clauses
| :none:`done` |
https://github.com/llvm/llvm-project/pull/159112 |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | indirect clause on declare target directive
| :part:`In Progress` |
|
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 2319ff13f7864..829eec6184afe 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -860,6 +860,8 @@ OpenMP Support
- Added parsing and semantic analysis support for ``need_device_ptr`` modifier
to accept an optional fallback argument (``fb_nullify`` or ``fb_preserve``)
with OpenMP >= 61.
+- ``use_device_ptr`` and ``use_device_addr`` now preserve the original host
+ address when lookup fails.
Improvements
^^^^^^^^^^^^
diff --git a/offload/libomptarget/omptarget.cpp
b/offload/libomptarget/omptarget.cpp
index 35c2c662a3884..676fda5fc8671 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -683,9 +683,39 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device,
int32_t ArgNum,
<< " new";
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
- void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- ODBG(ODT_Mapping) << "Returning device pointer " << TgtPtrBase;
+ uintptr_t Delta = reinterpret_cast<uintptr_t>(HstPtrBegin) -
+ reinterpret_cast<uintptr_t>(HstPtrBase);
+ void *TgtPtrBase;
+ if (TgtPtrBegin) {
+ // Lookup succeeded, return device pointer adjusted by delta
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<uintptr_t>(TgtPtrBegin) - Delta);
+ ODBG(ODT_Mapping) << "Returning device pointer " << TgtPtrBase;
+ } else {
+ // Lookup failed. So we have to decide what to do based on the
+ // requested fallback behavior.
+ //
+ // Treat "preserve" as the default fallback behavior, since as per
+ // OpenMP 5.1, for use_device_ptr/addr, when there's no corresponding
+ // device pointer to translate into, it's the user's responsibility to
+ // ensure that the host address is device-accessible.
+ //
+ // OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
+ // If a list item that appears in a use_device_ptr clause ... does not
+ // point to a mapped object, it must contain a valid device address for
+ // the target device, and the list item references are instead
converted
+ // to references to a local device pointer that refers to this device
+ // address.
+ //
+ // TODO: Add a new map-type bit to support OpenMP 6.1's `fb_nullify`
+ // and set the result to `nullptr - Delta`. Note that `fb_nullify` is
+ // already the default for `need_device_ptr`, but clang/flang do not
+ // support its codegen yet.
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
+ ODBG(ODT_Mapping) << "Returning host pointer " << TgtPtrBase
+ << " as fallback (lookup failed)";
+ }
ArgsBase[I] = TgtPtrBase;
}
diff --git
a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
index 4b67a3bc2aa7f..118b664fb6e53 100644
---
a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
+++
b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
@@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int h[10];
int *ph = &h[0];
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 833cc2bf262d4..7c2e3fda81a75 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
@@ -6,15 +6,6 @@
// 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];
@@ -34,7 +25,7 @@ struct S {
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);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@@ -56,7 +47,7 @@ struct S {
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);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -78,8 +69,7 @@ struct S {
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);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@@ -108,7 +98,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -120,7 +110,7 @@ struct S {
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);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.
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 1f89ab1575d01..e74c215593511 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
@@ -6,15 +6,6 @@
// 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;
@@ -35,15 +26,13 @@ struct S {
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.
+// 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);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@@ -61,15 +50,13 @@ struct S {
// (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.
+// 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);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -93,8 +80,7 @@ struct S {
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);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@@ -123,7 +109,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -135,7 +121,7 @@ struct S {
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);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.
diff --git
a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
index 4495a46b6d204..4b0819ef6a9fe 100644
---
a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
+++
b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
@@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int x;
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 b802857c013ae..7e79b6ec3b13e 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
@@ -6,15 +6,6 @@
// 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];
@@ -36,7 +27,7 @@ struct S {
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);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@@ -56,7 +47,7 @@ struct S {
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);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@@ -76,7 +67,7 @@ struct S {
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);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@@ -97,7 +88,7 @@ struct S {
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);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -117,7 +108,7 @@ struct S {
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);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -128,7 +119,7 @@ struct S {
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);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.
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 ded8a03d57997..41e0abae123d2 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
@@ -6,15 +6,6 @@
// 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;
@@ -43,7 +34,7 @@ struct S {
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);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@@ -63,7 +54,7 @@ struct S {
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);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@@ -83,7 +74,7 @@ struct S {
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);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@@ -104,7 +95,7 @@ struct S {
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);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -124,7 +115,7 @@ struct S {
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);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -135,7 +126,7 @@ struct S {
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);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.
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 28ec6857fa1a8..f8c9d7c1fe7df 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
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=51 -g
-// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \
+// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// FIXME: Fails due to optimized debugging in 'ptxas'
@@ -20,7 +20,8 @@ in...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/174659
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
