[OpenMP] Preserve the original address when use_device_ptr/addr lookup fails.
#169438
+78
−118
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
…ddr lookup failure. 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, so we could do anything in that scenario.
Member
|
@llvm/pr-subscribers-offload Author: Abhinav Gaba (abhinavgaba) ChangesAs per OpenMP 5.1, we need to assume that when the lookup for Prior to 5.1, a lookup-failure meant a user-error, so we could do anything Patch is 27.69 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169438.diff 13 Files Affected:
diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index f7e6061044c6d..7cebf96cfe026 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/169438 |
++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | iterators in map clause or motion clauses | :none:`unclaimed` | |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | indirect clause on declare target directive | :part:`In Progress` | |
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 51f07256c5d9f..ed22cdb39068f 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -759,6 +759,8 @@ OpenMP Support
- Updated parsing and semantic analysis support for ``nowait`` clause to accept
optional argument in OpenMP >= 60.
- Added support for ``default`` clause on ``target`` directive.
+- ``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 69725e77bae00..287564f53101a 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -675,9 +675,39 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
- void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- DP("Returning device pointer " DPxMOD "\n", DPxPTR(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);
+ DP("Returning device pointer " DPxMOD "\n", DPxPTR(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);
+ DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n",
+ DPxPTR(TgtPtrBase));
+ }
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 b9ebde431e7bf..78e6bf7c070a0 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
@@ -8,15 +8,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];
@@ -36,7 +27,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.
@@ -58,7 +49,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.
@@ -80,8 +71,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.
@@ -110,7 +100,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.
@@ -122,7 +112,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 0090cdb095366..d981da925acc2 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
@@ -8,15 +8,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;
@@ -37,15 +28,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.
@@ -63,15 +52,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.
@@ -95,8 +82,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.
@@ -125,7 +111,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.
@@ -137,7 +123,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 79c6f69edba8e..e855b0dd82744 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
@@ -8,15 +8,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];
@@ -38,7 +29,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.
@@ -58,7 +49,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.
@@ -78,7 +69,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.
@@ -99,7 +90,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.
@@ -119,7 +110,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.
@@ -130,7 +121,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 9360db4195041..1a3ed148f288b 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
@@ -8,15 +8,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;
@@ -45,7 +36,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.
@@ -65,7 +56,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.
@@ -85,7 +76,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.
@@ -106,7 +97,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.
@@ -126,7 +117,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.
@@ -137,7 +128,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 ...
[truncated]
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
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.
As per OpenMP 5.1, we need to assume that when the lookup for
use_device_ptr/addrfails, the incoming pointer was already deviceaccessible.
Prior to 5.1, a lookup-failure meant a user-error, so we could do anything
in that scenario.