Index: openmp/libomptarget/src/device.h =================================================================== --- openmp/libomptarget/src/device.h +++ openmp/libomptarget/src/device.h @@ -49,6 +49,8 @@ const uintptr_t TgtPtrBegin; // target info. + const bool IsUSMAlloc; // used to track maps under USM mode (optional) + private: static const uint64_t INFRefCount = ~(uint64_t)0; static std::string refCountToStr(uint64_t RefCount) { @@ -90,14 +92,15 @@ public: HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB, bool UseHoldRefCount, map_var_info_t Name = nullptr, - bool IsINF = false) + bool IsINF = false, bool IsUSMAlloc = false) : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name), - TgtPtrBegin(TB), States(std::make_unique(UseHoldRefCount ? 0 - : IsINF ? INFRefCount - : 1, - !UseHoldRefCount ? 0 - : IsINF ? INFRefCount - : 1)) {} + TgtPtrBegin(TB), IsUSMAlloc(IsUSMAlloc), + States(std::make_unique(UseHoldRefCount ? 0 + : IsINF ? INFRefCount + : 1, + !UseHoldRefCount ? 0 + : IsINF ? INFRefCount + : 1)) {} /// Get the total reference count. This is smarter than just getDynRefCount() /// + getHoldRefCount() because it handles the case where at least one is Index: openmp/libomptarget/src/device.cpp =================================================================== --- openmp/libomptarget/src/device.cpp +++ openmp/libomptarget/src/device.cpp @@ -224,6 +224,17 @@ DPxPTR((uintptr_t)HstPtrBegin), Size); IsHostPtr = true; TargetPointer = HstPtrBegin; + + // If user requested map checks under USM mode, add map information + // to the table and mark as usm allocation. + // Infinite ref count to disable memcopies to/from device memory + if (PM->RTLs.USMMapChecks) + Entry = HostDataToTargetMap + .emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin, + (uintptr_t)HstPtrBegin + Size, + (uintptr_t)HstPtrBegin, HasHoldModifier, + HstPtrName, /*IsInf=*/true, /*IsUSMAlloc=*/true) + .first; } } else if (HasPresentModifier) { DP("Mapping required by 'present' map type modifier does not exist for " @@ -296,8 +307,11 @@ DataMapMtx.lock(); LookupResult lr = lookupMapping(HstPtrBegin, Size); - if (lr.Flags.IsContained || - (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) { + // When map checks are enabled under USM mode, mapped host pointers are + // tracked in the map table but should be treated as in the USM case + if ((lr.Flags.IsContained || + (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) && + !lr.Entry->IsUSMAlloc) { auto &HT = *lr.Entry; // We do not zero the total reference count here. deallocTgtPtr does that // atomically with removing the mapping. Otherwise, before this thread @@ -334,9 +348,10 @@ DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction); rc = (void *)tp; } else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) { - // If the value isn't found in the mapping and unified shared memory - // is on then it means we have stumbled upon a value which we need to - // use directly from the host. + // If the value isn't found in the mapping or if the value is found but it + // is related to a USM mapping and unified shared memory is on then it means + // we have stumbled upon a value which we need to use directly from the + // host. DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared " "memory\n", DPxPTR((uintptr_t)HstPtrBegin), Size); @@ -371,16 +386,19 @@ if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) { auto &HT = *lr.Entry; if (HT.decRefCount(HasHoldModifier) == 0) { - DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n", - DPxPTR(HT.TgtPtrBegin), Size); - deleteData((void *)HT.TgtPtrBegin); + if (!HT.IsUSMAlloc) { + DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n", + DPxPTR(HT.TgtPtrBegin), Size); + deleteData((void *)HT.TgtPtrBegin); + } INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID, "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%" PRId64 ", Name=%s\n", DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size, (HT.HstPtrName) ? getNameFromMapping(HT.HstPtrName).c_str() : "unknown"); - HostDataToTargetMap.erase(lr.Entry); + if (PM->RTLs.USMMapChecks) + HostDataToTargetMap.erase(lr.Entry); } rc = OFFLOAD_SUCCESS; } else { Index: openmp/libomptarget/src/rtl.h =================================================================== --- openmp/libomptarget/src/rtl.h +++ openmp/libomptarget/src/rtl.h @@ -126,6 +126,9 @@ int64_t RequiresFlags = OMP_REQ_UNDEFINED; + // When true, perform map checks under USM mode + bool USMMapChecks = false; + explicit RTLsTy() = default; // Register the clauses of the requires directive. Index: openmp/libomptarget/src/rtl.cpp =================================================================== --- openmp/libomptarget/src/rtl.cpp +++ openmp/libomptarget/src/rtl.cpp @@ -77,6 +77,18 @@ return; } + // Parse environment variable LIBOMPTARGET_USM_MAP_CHECKS to enable map checks + // under USM mode Expected values are: + // * "check" : enable map checks under usm mode + // * "none" : disable map checks under usm mode + // * not set : disable map checks under usm mode + // (no device allocations or host-to-device and device-to-host transfers) + if (const char *USMMapChecksStr = getenv("LIBOMPTARGET_USM_MAP_CHECKS")) { + const char *checkStr = "check"; + if (strncmp(USMMapChecksStr, checkStr, strlen(checkStr)) == 0) + USMMapChecks = true; + } + DP("Loading RTLs...\n"); // Attempt to open all the plugins and, if they exist, check if the interface Index: openmp/libomptarget/test/unified_shared_memory/check_maps.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/unified_shared_memory/check_maps.c @@ -0,0 +1,55 @@ +// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_USM_MAP_CHECKS=check +// %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty +// -check-prefix=CHECK-USM-CHECK RUN: %libomptarget-compile-generic && env +// LIBOMPTARGET_USM_MAP_CHECKS=none %libomptarget-run-generic 2>&1 | +// %fcheck-generic -allow-empty -check-prefix=CHECK-USM-NONE RUN: +// %libomptarget-compile-generic && %libomptarget-run-generic 2>&1 | +// %fcheck-generic -allow-empty -check-prefix=CHECK-USM-UNSPEC + +// CHECK-USM-CHECK: Libomptarget message: explicit extension not allowed: host +// address specified is 0x{{.*}}, but device allocation maps to host at 0x{{.*}} +// CHECK-USM-CHECK: Libomptarget error: Call to getTargetPointer returned null +// pointer (device failure or illegal mapping). CHECK-USM-CHECK: Libomptarget +// fatal error 1: failure of target construct while offloading is mandatory + +// CHECK-USM-NONE: No errors + +// CHECK-USM-UNSPEC: No errors + +// REQUIRES: unified_shared_memory + +#include +#include + +// --------------------------------------------------------------------------- +// Various definitions copied from OpenMP RTL + +extern void __tgt_register_requires(int64_t); + +// End of definitions copied from OpenMP RTL. +// --------------------------------------------------------------------------- + +#pragma omp requires unified_shared_memory + +#define N 1024 + +int main(int argc, char *argv[]) { + int data[N]; + + // Manual registration of requires flags for Clang versions + // that do not support requires. + __tgt_register_requires(8); + +#pragma omp target enter data map(to : data [0:N / 2]) + +// extending a map under unified_shared_memory mode is prohibited +// when map checks are on +#pragma omp target enter data map(to : data [0:N]) + + printf("No errors\n"); + +#pragma omp target + {} + + return 0; +}