[Openmp-commits] [openmp] d11bab0 - [OpenMP] Use IsHostPtr where needed for targetDataBegin
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Wed Sep 1 14:38:12 PDT 2021
Author: Joel E. Denny
Date: 2021-09-01T17:31:42-04:00
New Revision: d11bab0b73ec485f0b1e8dc38af0be72fcda1e34
URL: https://github.com/llvm/llvm-project/commit/d11bab0b73ec485f0b1e8dc38af0be72fcda1e34
DIFF: https://github.com/llvm/llvm-project/commit/d11bab0b73ec485f0b1e8dc38af0be72fcda1e34.diff
LOG: [OpenMP] Use IsHostPtr where needed for targetDataBegin
As discussed in D105990, without this patch, `targetDataBegin`
determines whether to transfer data (as opposed to assuming it's in
shared memory) using the condition `!UseUSM || HasCloseModifier`.
However, this condition is broken if use of discrete memory was forced
by `omp_target_associate_ptr`. This patch extends
`unified_shared_memory/associate_ptr.c` to reveal this case, and it
fixes it using `!IsHostPtr` in `DeviceTy::getTargetPointer` to replace
this condition.
Reviewed By: grokos
Differential Revision: https://reviews.llvm.org/D107927
Added:
Modified:
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/device.h
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/test/unified_shared_memory/associate_ptr.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index c70d9e8975726..62d694ec8529e 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -177,10 +177,11 @@ LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) {
TargetPointerResultTy
DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
- map_var_info_t HstPtrName, MoveDataStateTy MoveData,
- bool IsImplicit, bool UpdateRefCount,
- bool HasCloseModifier, bool HasPresentModifier,
- bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
+ map_var_info_t HstPtrName, bool HasFlagTo,
+ bool HasFlagAlways, bool IsImplicit,
+ bool UpdateRefCount, bool HasCloseModifier,
+ bool HasPresentModifier, bool HasHoldModifier,
+ AsyncInfoTy &AsyncInfo) {
void *TargetPointer = nullptr;
bool IsHostPtr = false;
bool IsNew = false;
@@ -272,12 +273,9 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
TargetPointer = (void *)Ptr;
}
- if (IsNew && MoveData == MoveDataStateTy::UNKNOWN)
- MoveData = MoveDataStateTy::REQUIRED;
-
// If the target pointer is valid, and we need to transfer data, issue the
// data transfer.
- if (TargetPointer && (MoveData == MoveDataStateTy::REQUIRED)) {
+ if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {
// Lock the entry before releasing the mapping table lock such that another
// thread that could issue data movement will get the right result.
Entry->lock();
diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h
index ea87a4b270f0d..16120537c1f61 100644
--- a/openmp/libomptarget/src/device.h
+++ b/openmp/libomptarget/src/device.h
@@ -226,8 +226,6 @@ struct PendingCtorDtorListsTy {
typedef std::map<__tgt_bin_desc *, PendingCtorDtorListsTy>
PendingCtorsDtorsPerLibrary;
-enum class MoveDataStateTy : uint32_t { REQUIRED, NONE, UNKNOWN };
-
struct DeviceTy {
int32_t DeviceID;
RTLInfoTy *RTL;
@@ -264,20 +262,20 @@ struct DeviceTy {
LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
/// Get the target pointer based on host pointer begin and base. If the
/// mapping already exists, the target pointer will be returned directly. In
- /// addition, if \p MoveData is true, the memory region pointed by \p
- /// HstPtrBegin of size \p Size will also be transferred to the device. If the
- /// mapping doesn't exist, and if unified memory is not enabled, a new mapping
- /// will be created and the data will also be transferred accordingly. nullptr
- /// will be returned because of any of following reasons:
+ /// addition, if required, the memory region pointed by \p HstPtrBegin of size
+ /// \p Size will also be transferred to the device. If the mapping doesn't
+ /// exist, and if unified shared memory is not enabled, a new mapping will be
+ /// created and the data will also be transferred accordingly. nullptr will be
+ /// returned because of any of following reasons:
/// - Data allocation failed;
/// - The user tried to do an illegal mapping;
/// - Data transfer issue fails.
TargetPointerResultTy
getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
- map_var_info_t HstPtrName, MoveDataStateTy MoveData,
- bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
- bool HasPresentModifier, bool HasHoldModifier,
- AsyncInfoTy &AsyncInfo);
+ map_var_info_t HstPtrName, bool HasFlagTo,
+ bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
+ bool HasCloseModifier, bool HasPresentModifier,
+ bool HasHoldModifier, AsyncInfoTy &AsyncInfo);
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
bool UpdateRefCount, bool UseHoldRefCount,
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index fe50563d94a6c..55b840a611fed 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -490,9 +490,9 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
// PTR_AND_OBJ entry is handled below, and so the allocation might fail
// when HasPresentModifier.
Pointer_TPR = Device.getTargetPointer(
- HstPtrBase, HstPtrBase, sizeof(void *), nullptr,
- MoveDataStateTy::NONE, IsImplicit, UpdateRef, HasCloseModifier,
- HasPresentModifier, HasHoldModifier, AsyncInfo);
+ HstPtrBase, HstPtrBase, sizeof(void *), /*HstPtrName=*/nullptr,
+ /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,
+ HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo);
PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
if (!PointerTgtPtrBegin) {
@@ -514,18 +514,13 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
(!FromMapper || i != 0); // subsequently update ref count of pointee
}
- MoveDataStateTy MoveData = MoveDataStateTy::NONE;
- const bool UseUSM = PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY;
const bool HasFlagTo = arg_types[i] & OMP_TGT_MAPTYPE_TO;
const bool HasFlagAlways = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
- if (HasFlagTo && (!UseUSM || HasCloseModifier))
- MoveData = HasFlagAlways ? MoveDataStateTy::REQUIRED
- : MoveDataStateTy::UNKNOWN;
-
- auto TPR = Device.getTargetPointer(
- HstPtrBegin, HstPtrBase, data_size, HstPtrName, MoveData, IsImplicit,
- UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier,
- AsyncInfo);
+ auto TPR = Device.getTargetPointer(HstPtrBegin, HstPtrBase, data_size,
+ HstPtrName, HasFlagTo, HasFlagAlways,
+ IsImplicit, UpdateRef, HasCloseModifier,
+ HasPresentModifier, HasHoldModifier,
+ AsyncInfo);
void *TgtPtrBegin = TPR.TargetPointer;
IsHostPtr = TPR.Flags.IsHostPointer;
// If data_size==0, then the argument could be a zero-length pointer to
diff --git a/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c b/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c
index 7911046f5f3b1..c25c557b3b6d3 100644
--- a/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c
+++ b/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c
@@ -25,9 +25,9 @@ int main(int argc, char *argv[]) {
// specified. It must check whether x was previously placed in device memory
// by, for example, omp_target_associate_ptr.
#pragma omp target map(always, tofrom: x)
- x = 20;
+ x += 1;
- // CHECK: x=20
+ // CHECK: x=11
printf("x=%d\n", x);
// CHECK: present: 1
printf("present: %d\n", omp_target_is_present(&x, dev));
More information about the Openmp-commits
mailing list