//==-------------- memory_manager.cpp --------------------------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include "ur_api.h" #include #include #include #include #include #include #include #include #include #include #include #include #include #include #include #include #ifdef XPTI_ENABLE_INSTRUMENTATION #include #include #endif namespace sycl { inline namespace _V1 { namespace detail { #ifdef XPTI_ENABLE_INSTRUMENTATION uint8_t GMemAllocStreamID; xpti::trace_event_data_t *GMemAllocEvent; #endif uint64_t emitMemAllocBeginTrace(uintptr_t ObjHandle, size_t AllocSize, size_t GuardZone) { (void)ObjHandle; (void)AllocSize; (void)GuardZone; uint64_t CorrelationID = 0; #ifdef XPTI_ENABLE_INSTRUMENTATION constexpr uint16_t NotificationTraceType = static_cast(xpti::trace_point_type_t::mem_alloc_begin); if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) { xpti::mem_alloc_data_t MemAlloc{ObjHandle, 0 /* alloc ptr */, AllocSize, GuardZone}; CorrelationID = xptiGetUniqueId(); xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType, GMemAllocEvent, nullptr, CorrelationID, &MemAlloc); } #endif return CorrelationID; } void emitMemAllocEndTrace(uintptr_t ObjHandle, uintptr_t AllocPtr, size_t AllocSize, size_t GuardZone, uint64_t CorrelationID) { (void)ObjHandle; (void)AllocPtr; (void)AllocSize; (void)GuardZone; (void)CorrelationID; #ifdef XPTI_ENABLE_INSTRUMENTATION constexpr uint16_t NotificationTraceType = static_cast(xpti::trace_point_type_t::mem_alloc_end); if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) { xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, AllocSize, GuardZone}; xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType, GMemAllocEvent, nullptr, CorrelationID, &MemAlloc); } #endif } uint64_t emitMemReleaseBeginTrace(uintptr_t ObjHandle, uintptr_t AllocPtr) { (void)ObjHandle; (void)AllocPtr; uint64_t CorrelationID = 0; #ifdef XPTI_ENABLE_INSTRUMENTATION constexpr uint16_t NotificationTraceType = static_cast(xpti::trace_point_type_t::mem_release_begin); if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) { xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, 0 /* alloc size */, 0 /* guard zone */}; CorrelationID = xptiGetUniqueId(); xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType, GMemAllocEvent, nullptr, CorrelationID, &MemAlloc); } #endif return CorrelationID; } void emitMemReleaseEndTrace(uintptr_t ObjHandle, uintptr_t AllocPtr, uint64_t CorrelationID) { (void)ObjHandle; (void)AllocPtr; (void)CorrelationID; #ifdef XPTI_ENABLE_INSTRUMENTATION constexpr uint16_t NotificationTraceType = static_cast(xpti::trace_point_type_t::mem_release_end); if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) { xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, 0 /* alloc size */, 0 /* guard zone */}; xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType, GMemAllocEvent, nullptr, CorrelationID, &MemAlloc); } #endif } static void waitForEvents(const std::vector &Events) { // Assuming all events will be on the same device or // devices associated with the same Backend. if (!Events.empty()) { const AdapterPtr &Adapter = Events[0]->getAdapter(); std::vector UrEvents(Events.size()); std::transform( Events.begin(), Events.end(), UrEvents.begin(), [](const EventImplPtr &EventImpl) { return EventImpl->getHandle(); }); if (!UrEvents.empty() && UrEvents[0]) { Adapter->call(UrEvents.size(), &UrEvents[0]); } } } void memBufferCreateHelper(const AdapterPtr &Adapter, ur_context_handle_t Ctx, ur_mem_flags_t Flags, size_t Size, ur_mem_handle_t *RetMem, const ur_buffer_properties_t *Props) { #ifdef XPTI_ENABLE_INSTRUMENTATION uint64_t CorrID = 0; #endif // We only want to instrument urMemBufferCreate { #ifdef XPTI_ENABLE_INSTRUMENTATION CorrID = emitMemAllocBeginTrace(0 /* mem object */, Size, 0 /* guard zone */); xpti::utils::finally _{[&] { // C-style cast is required for MSVC uintptr_t MemObjID = (uintptr_t)(*RetMem); ur_native_handle_t Ptr = 0; // Always use call_nocheck here, because call may throw an exception, // and this lambda will be called from destructor, which in combination // rewards us with UB. // When doing buffer interop we don't know what device the memory should // be resident on, so pass nullptr for Device param. Buffer interop may // not be supported by all backends. Adapter->call_nocheck( *RetMem, /*Dev*/ nullptr, &Ptr); emitMemAllocEndTrace(MemObjID, (uintptr_t)(Ptr), Size, 0 /* guard zone */, CorrID); }}; #endif if (Size) Adapter->call(Ctx, Flags, Size, Props, RetMem); } } void memReleaseHelper(const AdapterPtr &Adapter, ur_mem_handle_t Mem) { // FIXME urMemRelease does not guarante memory release. It is only true if // reference counter is 1. However, SYCL runtime currently only calls // urMemRetain only for OpenCL interop #ifdef XPTI_ENABLE_INSTRUMENTATION uint64_t CorrID = 0; // C-style cast is required for MSVC uintptr_t MemObjID = (uintptr_t)(Mem); uintptr_t Ptr = 0; // Do not make unnecessary UR calls without instrumentation enabled if (xptiTraceEnabled()) { ur_native_handle_t PtrHandle = 0; // When doing buffer interop we don't know what device the memory should be // resident on, so pass nullptr for Device param. Buffer interop may not be // supported by all backends. Adapter->call_nocheck(Mem, /*Dev*/ nullptr, &PtrHandle); Ptr = (uintptr_t)(PtrHandle); } #endif // We only want to instrument urMemRelease { #ifdef XPTI_ENABLE_INSTRUMENTATION CorrID = emitMemReleaseBeginTrace(MemObjID, Ptr); xpti::utils::finally _{ [&] { emitMemReleaseEndTrace(MemObjID, Ptr, CorrID); }}; #endif Adapter->call(Mem); } } void memBufferMapHelper(const AdapterPtr &Adapter, ur_queue_handle_t Queue, ur_mem_handle_t Buffer, bool Blocking, ur_map_flags_t Flags, size_t Offset, size_t Size, uint32_t NumEvents, const ur_event_handle_t *WaitList, ur_event_handle_t *Event, void **RetMap) { #ifdef XPTI_ENABLE_INSTRUMENTATION uint64_t CorrID = 0; uintptr_t MemObjID = (uintptr_t)(Buffer); #endif // We only want to instrument urEnqueueMemBufferMap #ifdef XPTI_ENABLE_INSTRUMENTATION CorrID = emitMemAllocBeginTrace(MemObjID, Size, 0 /* guard zone */); xpti::utils::finally _{[&] { emitMemAllocEndTrace(MemObjID, (uintptr_t)(*RetMap), Size, 0 /* guard zone */, CorrID); }}; #endif Adapter->call( Queue, Buffer, Blocking, Flags, Offset, Size, NumEvents, WaitList, Event, RetMap); } void memUnmapHelper(const AdapterPtr &Adapter, ur_queue_handle_t Queue, ur_mem_handle_t Mem, void *MappedPtr, uint32_t NumEvents, const ur_event_handle_t *WaitList, ur_event_handle_t *Event) { #ifdef XPTI_ENABLE_INSTRUMENTATION uint64_t CorrID = 0; uintptr_t MemObjID = (uintptr_t)(Mem); uintptr_t Ptr = (uintptr_t)(MappedPtr); #endif // We only want to instrument urEnqueueMemUnmap { #ifdef XPTI_ENABLE_INSTRUMENTATION CorrID = emitMemReleaseBeginTrace(MemObjID, Ptr); xpti::utils::finally _{[&] { // There's no way for SYCL to know, when the pointer is freed, so we have // to explicitly wait for the end of data transfers here in order to // report correct events. // Always use call_nocheck here, because call may throw an exception, // and this lambda will be called from destructor, which in combination // rewards us with UB. Adapter->call_nocheck(1, Event); emitMemReleaseEndTrace(MemObjID, Ptr, CorrID); }}; #endif Adapter->call(Queue, Mem, MappedPtr, NumEvents, WaitList, Event); } } void MemoryManager::release(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *MemAllocation, std::vector DepEvents, ur_event_handle_t &OutEvent) { // There is no async API for memory releasing. Explicitly wait for all // dependency events and return empty event. waitForEvents(DepEvents); OutEvent = nullptr; XPTIRegistry::bufferReleaseNotification(MemObj, MemAllocation); MemObj->releaseMem(TargetContext, MemAllocation); } void MemoryManager::releaseMemObj(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *MemAllocation, void *UserPtr) { if (UserPtr == MemAllocation) { // Do nothing as it's user provided memory. return; } if (!TargetContext) { MemObj->releaseHostMem(MemAllocation); return; } const AdapterPtr &Adapter = TargetContext->getAdapter(); memReleaseHelper(Adapter, ur::cast(MemAllocation)); } void *MemoryManager::allocate(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, bool InitFromUserData, void *HostPtr, std::vector DepEvents, ur_event_handle_t &OutEvent) { // There is no async API for memory allocation. Explicitly wait for all // dependency events and return empty event. waitForEvents(DepEvents); OutEvent = nullptr; return MemObj->allocateMem(TargetContext, InitFromUserData, HostPtr, OutEvent); } void *MemoryManager::allocateHostMemory(SYCLMemObjI *MemObj, void *UserPtr, bool HostPtrReadOnly, size_t Size, const sycl::property_list &) { std::ignore = HostPtrReadOnly; std::ignore = Size; // Can return user pointer directly if it is not a nullptr. if (UserPtr) return UserPtr; return MemObj->allocateHostMem(); } void *MemoryManager::allocateInteropMemObject( ContextImplPtr TargetContext, void *UserPtr, const EventImplPtr &InteropEvent, const ContextImplPtr &InteropContext, const sycl::property_list &, ur_event_handle_t &OutEventToWait) { (void)TargetContext; (void)InteropContext; // If memory object is created with interop c'tor return cl_mem as is. assert(TargetContext == InteropContext && "Expected matching contexts"); OutEventToWait = InteropEvent->getHandle(); // Retain the event since it will be released during alloca command // destruction if (nullptr != OutEventToWait) { const AdapterPtr &Adapter = InteropEvent->getAdapter(); Adapter->call(OutEventToWait); } return UserPtr; } static ur_mem_flags_t getMemObjCreationFlags(void *UserPtr, bool HostPtrReadOnly) { // Create read_write mem object to handle arbitrary uses. ur_mem_flags_t Result = HostPtrReadOnly ? UR_MEM_FLAG_READ_ONLY : UR_MEM_FLAG_READ_WRITE; if (UserPtr) Result |= UR_MEM_FLAG_USE_HOST_POINTER; return Result; } void *MemoryManager::allocateImageObject(ContextImplPtr TargetContext, void *UserPtr, bool HostPtrReadOnly, const ur_image_desc_t &Desc, const ur_image_format_t &Format, const sycl::property_list &) { ur_mem_flags_t CreationFlags = getMemObjCreationFlags(UserPtr, HostPtrReadOnly); ur_mem_handle_t NewMem = nullptr; const AdapterPtr &Adapter = TargetContext->getAdapter(); Adapter->call(TargetContext->getHandleRef(), CreationFlags, &Format, &Desc, UserPtr, &NewMem); return NewMem; } void * MemoryManager::allocateBufferObject(ContextImplPtr TargetContext, void *UserPtr, bool HostPtrReadOnly, const size_t Size, const sycl::property_list &PropsList) { ur_mem_flags_t CreationFlags = getMemObjCreationFlags(UserPtr, HostPtrReadOnly); if (PropsList.has_property< sycl::ext::oneapi::property::buffer::use_pinned_host_memory>()) CreationFlags |= UR_MEM_FLAG_ALLOC_HOST_POINTER; ur_mem_handle_t NewMem = nullptr; const AdapterPtr &Adapter = TargetContext->getAdapter(); ur_buffer_properties_t AllocProps = {UR_STRUCTURE_TYPE_BUFFER_PROPERTIES, nullptr, UserPtr}; void **Next = &AllocProps.pNext; ur_buffer_alloc_location_properties_t LocationProperties = { UR_STRUCTURE_TYPE_BUFFER_ALLOC_LOCATION_PROPERTIES, nullptr, 0}; if (PropsList.has_property() && TargetContext->isBufferLocationSupported()) { LocationProperties.location = PropsList.get_property() .get_buffer_location(); *Next = &LocationProperties; Next = &LocationProperties.pNext; } ur_buffer_channel_properties_t ChannelProperties = { UR_STRUCTURE_TYPE_BUFFER_CHANNEL_PROPERTIES, nullptr, 0}; if (PropsList.has_property()) { ChannelProperties.channel = PropsList.get_property().get_channel(); *Next = &ChannelProperties; } memBufferCreateHelper(Adapter, TargetContext->getHandleRef(), CreationFlags, Size, &NewMem, &AllocProps); return NewMem; } void *MemoryManager::allocateMemBuffer(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *UserPtr, bool HostPtrReadOnly, size_t Size, const EventImplPtr &InteropEvent, const ContextImplPtr &InteropContext, const sycl::property_list &PropsList, ur_event_handle_t &OutEventToWait) { void *MemPtr; if (!TargetContext) MemPtr = allocateHostMemory(MemObj, UserPtr, HostPtrReadOnly, Size, PropsList); else if (UserPtr && InteropContext) MemPtr = allocateInteropMemObject(TargetContext, UserPtr, InteropEvent, InteropContext, PropsList, OutEventToWait); else MemPtr = allocateBufferObject(TargetContext, UserPtr, HostPtrReadOnly, Size, PropsList); XPTIRegistry::bufferAssociateNotification(MemObj, MemPtr); return MemPtr; } void *MemoryManager::allocateMemImage( ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *UserPtr, bool HostPtrReadOnly, size_t Size, const ur_image_desc_t &Desc, const ur_image_format_t &Format, const EventImplPtr &InteropEvent, const ContextImplPtr &InteropContext, const sycl::property_list &PropsList, ur_event_handle_t &OutEventToWait) { if (!TargetContext) return allocateHostMemory(MemObj, UserPtr, HostPtrReadOnly, Size, PropsList); if (UserPtr && InteropContext) return allocateInteropMemObject(TargetContext, UserPtr, InteropEvent, InteropContext, PropsList, OutEventToWait); return allocateImageObject(TargetContext, UserPtr, HostPtrReadOnly, Desc, Format, PropsList); } void *MemoryManager::allocateMemSubBuffer(ContextImplPtr TargetContext, void *ParentMemObj, size_t ElemSize, size_t Offset, range<3> Range, std::vector DepEvents, ur_event_handle_t &OutEvent) { waitForEvents(DepEvents); OutEvent = nullptr; if (!TargetContext) return static_cast(static_cast(ParentMemObj) + Offset); size_t SizeInBytes = ElemSize; for (size_t I = 0; I < 3; ++I) SizeInBytes *= Range[I]; ur_result_t Error = UR_RESULT_SUCCESS; ur_buffer_region_t Region = {UR_STRUCTURE_TYPE_BUFFER_REGION, nullptr, Offset, SizeInBytes}; ur_mem_handle_t NewMem; const AdapterPtr &Adapter = TargetContext->getAdapter(); Error = Adapter->call_nocheck( ur::cast(ParentMemObj), UR_MEM_FLAG_READ_WRITE, UR_BUFFER_CREATE_TYPE_REGION, &Region, &NewMem); if (Error == UR_RESULT_ERROR_MISALIGNED_SUB_BUFFER_OFFSET) throw detail::set_ur_error( exception(make_error_code(errc::invalid), "Specified offset of the sub-buffer being constructed is not " "a multiple of the memory base address alignment"), Error); Adapter->checkUrResult(Error); return NewMem; } struct TermPositions { int XTerm; int YTerm; int ZTerm; }; void prepTermPositions(TermPositions &pos, int Dimensions, detail::SYCLMemObjI::MemObjType Type) { // For buffers, the offsets/ranges coming from accessor are always // id<3>/range<3> But their organization varies by dimension: // 1 ==> {width, 1, 1} // 2 ==> {height, width, 1} // 3 ==> {depth, height, width} // Some callers schedule 0 as DimDst/DimSrc. if (Type == detail::SYCLMemObjI::MemObjType::Buffer) { if (Dimensions == 3) { pos.XTerm = 2, pos.YTerm = 1, pos.ZTerm = 0; } else if (Dimensions == 2) { pos.XTerm = 1, pos.YTerm = 0, pos.ZTerm = 2; } else { // Dimension is 1 or 0 pos.XTerm = 0, pos.YTerm = 1, pos.ZTerm = 2; } } else { // While range<>/id<> use by images is different than buffers, it's // consistent with their accessors. pos.XTerm = 0; pos.YTerm = 1; pos.ZTerm = 2; } } void copyH2D(SYCLMemObjI *SYCLMemObj, char *SrcMem, QueueImplPtr, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, ur_mem_handle_t DstMem, QueueImplPtr TgtQueue, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl) { (void)SrcAccessRange; assert(SYCLMemObj && "The SYCLMemObj is nullptr"); assert(TgtQueue && "Destination mem object queue must be not nullptr"); const ur_queue_handle_t Queue = TgtQueue->getHandleRef(); const AdapterPtr &Adapter = TgtQueue->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t DstAccessRangeWidthBytes = DstAccessRange[DstPos.XTerm] * DstElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) { if (1 == DimDst && 1 == DimSrc) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, DstMem, /*blocking_write=*/false, DstXOffBytes, DstAccessRangeWidthBytes, SrcMem + SrcXOffBytes, DepEvents.size(), DepEvents.data(), &OutEvent); } else { size_t BufferRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes; size_t BufferSlicePitch = (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0; size_t HostRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes; size_t HostSlicePitch = (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0; ur_rect_offset_t BufferOffset{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_offset_t HostOffset{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_region_t RectRegion{DstAccessRangeWidthBytes, DstAccessRange[DstPos.YTerm], DstAccessRange[DstPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, DstMem, /*blocking_write=*/false, BufferOffset, HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch, HostSlicePitch, SrcMem, DepEvents.size(), DepEvents.data(), &OutEvent); } } else { size_t InputRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes; size_t InputSlicePitch = (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0; ur_rect_offset_t Origin{DstOffset[DstPos.XTerm], DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t Region{DstAccessRange[DstPos.XTerm], DstAccessRange[DstPos.YTerm], DstAccessRange[DstPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, DstMem, /*blocking_write=*/false, Origin, Region, InputRowPitch, InputSlicePitch, SrcMem, DepEvents.size(), DepEvents.data(), &OutEvent); } } void copyD2H(SYCLMemObjI *SYCLMemObj, ur_mem_handle_t SrcMem, QueueImplPtr SrcQueue, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, char *DstMem, QueueImplPtr, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl) { (void)DstAccessRange; assert(SYCLMemObj && "The SYCLMemObj is nullptr"); assert(SrcQueue && "Source mem object queue is expected to be not nullptr"); const ur_queue_handle_t Queue = SrcQueue->getHandleRef(); const AdapterPtr &Adapter = SrcQueue->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); // For a given buffer, the various mem copy routines (copyD2H, copyH2D, // copyD2D) will usually have the same values for AccessRange, Size, // Dimension, Offset, etc. EXCEPT when the dtor for ~SYCLMemObjT is called. // Essentially, it schedules a copyBack of chars thus in copyD2H the // Dimension will then be 1 and DstAccessRange[0] and DstSize[0] will be // sized to bytes with a DstElemSize of 1. size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.XTerm] * SrcElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) { if (1 == DimDst && 1 == DimSrc) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, /*blocking_read=*/false, SrcXOffBytes, SrcAccessRangeWidthBytes, DstMem + DstXOffBytes, DepEvents.size(), DepEvents.data(), &OutEvent); } else { size_t BufferRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes; size_t BufferSlicePitch = (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0; size_t HostRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes; size_t HostSlicePitch = (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0; ur_rect_offset_t BufferOffset{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_offset_t HostOffset{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t RectRegion{SrcAccessRangeWidthBytes, SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, /*blocking_read=*/false, BufferOffset, HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch, HostSlicePitch, DstMem, DepEvents.size(), DepEvents.data(), &OutEvent); } } else { size_t RowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes; size_t SlicePitch = (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0; ur_rect_offset_t Offset{SrcOffset[SrcPos.XTerm], SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_region_t Region{SrcAccessRange[SrcPos.XTerm], SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, false, Offset, Region, RowPitch, SlicePitch, DstMem, DepEvents.size(), DepEvents.data(), &OutEvent); } } void copyD2D(SYCLMemObjI *SYCLMemObj, ur_mem_handle_t SrcMem, QueueImplPtr SrcQueue, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, ur_mem_handle_t DstMem, QueueImplPtr, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3>, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); assert(SrcQueue && "Source mem object and target mem object queues are " "expected to be not nullptr"); const ur_queue_handle_t Queue = SrcQueue->getHandleRef(); const AdapterPtr &Adapter = SrcQueue->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.XTerm] * SrcElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) { if (1 == DimDst && 1 == DimSrc) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, DstMem, SrcXOffBytes, DstXOffBytes, SrcAccessRangeWidthBytes, DepEvents.size(), DepEvents.data(), &OutEvent); } else { // passing 0 for pitches not allowed. Because clEnqueueCopyBufferRect will // calculate both src and dest pitch using region[0], which is not correct // if src and dest are not the same size. size_t SrcRowPitch = SrcSzWidthBytes; size_t SrcSlicePitch = (DimSrc <= 1) ? SrcSzWidthBytes : SrcSzWidthBytes * SrcSize[SrcPos.YTerm]; size_t DstRowPitch = DstSzWidthBytes; size_t DstSlicePitch = (DimDst <= 1) ? DstSzWidthBytes : DstSzWidthBytes * DstSize[DstPos.YTerm]; ur_rect_offset_t SrcOrigin{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_offset_t DstOrigin{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t Region{SrcAccessRangeWidthBytes, SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, DstMem, SrcOrigin, DstOrigin, Region, SrcRowPitch, SrcSlicePitch, DstRowPitch, DstSlicePitch, DepEvents.size(), DepEvents.data(), &OutEvent); } } else { ur_rect_offset_t SrcOrigin{SrcOffset[SrcPos.XTerm], SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_offset_t DstOrigin{DstOffset[DstPos.XTerm], DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t Region{SrcAccessRange[SrcPos.XTerm], SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call( Queue, SrcMem, DstMem, SrcOrigin, DstOrigin, Region, DepEvents.size(), DepEvents.data(), &OutEvent); } } static void copyH2H(SYCLMemObjI *, char *SrcMem, QueueImplPtr, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, char *DstMem, QueueImplPtr, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector, ur_event_handle_t &, const detail::EventImplPtr &) { if ((DimSrc != 1 || DimDst != 1) && (SrcOffset != id<3>{0, 0, 0} || DstOffset != id<3>{0, 0, 0} || SrcSize != SrcAccessRange || DstSize != DstAccessRange)) { throw exception(make_error_code(errc::feature_not_supported), "Not supported configuration of memcpy requested"); } SrcMem += SrcOffset[0] * SrcElemSize; DstMem += DstOffset[0] * DstElemSize; if (SrcMem == DstMem) return; size_t BytesToCopy = SrcAccessRange[0] * SrcElemSize * SrcAccessRange[1] * SrcAccessRange[2]; std::memcpy(DstMem, SrcMem, BytesToCopy); } // Copies memory between: host and device, host and host, // device and device if memory objects bound to the one context. void MemoryManager::copy( SYCLMemObjI *SYCLMemObj, void *SrcMem, QueueImplPtr SrcQueue, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, void *DstMem, QueueImplPtr TgtQueue, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl) { if (!SrcQueue) { if (!TgtQueue) copyH2H(SYCLMemObj, (char *)SrcMem, nullptr, DimSrc, SrcSize, SrcAccessRange, SrcOffset, SrcElemSize, (char *)DstMem, nullptr, DimDst, DstSize, DstAccessRange, DstOffset, DstElemSize, std::move(DepEvents), OutEvent, OutEventImpl); else copyH2D(SYCLMemObj, (char *)SrcMem, nullptr, DimSrc, SrcSize, SrcAccessRange, SrcOffset, SrcElemSize, ur::cast(DstMem), std::move(TgtQueue), DimDst, DstSize, DstAccessRange, DstOffset, DstElemSize, std::move(DepEvents), OutEvent, OutEventImpl); } else { if (!TgtQueue) copyD2H(SYCLMemObj, ur::cast(SrcMem), std::move(SrcQueue), DimSrc, SrcSize, SrcAccessRange, SrcOffset, SrcElemSize, (char *)DstMem, nullptr, DimDst, DstSize, DstAccessRange, DstOffset, DstElemSize, std::move(DepEvents), OutEvent, OutEventImpl); else copyD2D(SYCLMemObj, ur::cast(SrcMem), std::move(SrcQueue), DimSrc, SrcSize, SrcAccessRange, SrcOffset, SrcElemSize, ur::cast(DstMem), std::move(TgtQueue), DimDst, DstSize, DstAccessRange, DstOffset, DstElemSize, std::move(DepEvents), OutEvent, OutEventImpl); } } void MemoryManager::fill(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue, size_t PatternSize, const unsigned char *Pattern, unsigned int Dim, sycl::range<3> MemRange, sycl::range<3> AccRange, sycl::id<3> Offset, unsigned int ElementSize, std::vector DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); assert(Queue && "Fill should be called only with a valid device queue"); const AdapterPtr &Adapter = Queue->getAdapter(); if (SYCLMemObj->getType() == detail::SYCLMemObjI::MemObjType::Buffer) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); // 2D and 3D buffers accessors can't have custom range or the data will // likely be discontiguous. bool RangesUsable = (Dim <= 1) || (MemRange == AccRange); // For 2D and 3D buffers, the offset must be 0, or the data will be // discontiguous. bool OffsetUsable = (Dim <= 1) || (Offset == sycl::id<3>{0, 0, 0}); size_t RangeMultiplier = AccRange[0] * AccRange[1] * AccRange[2]; if (RangesUsable && OffsetUsable) { Adapter->call( Queue->getHandleRef(), ur::cast(Mem), Pattern, PatternSize, Offset[0] * ElementSize, RangeMultiplier * ElementSize, DepEvents.size(), DepEvents.data(), &OutEvent); return; } // The sycl::handler uses a parallel_for kernel in the case of unusable // Range or Offset, not CG:Fill. So we should not be here. throw exception(make_error_code(errc::runtime), "Not supported configuration of fill requested"); } else { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); // We don't have any backend implementations that support enqueueing a fill // on non-buffer mem objects like this. The old UR function was a stub with // an abort. throw exception(make_error_code(errc::runtime), "Fill operation not supported for the given mem object"); } } void *MemoryManager::map(SYCLMemObjI *, void *Mem, QueueImplPtr Queue, access::mode AccessMode, unsigned int, sycl::range<3>, sycl::range<3> AccessRange, sycl::id<3> AccessOffset, unsigned int ElementSize, std::vector DepEvents, ur_event_handle_t &OutEvent) { if (!Queue) { throw exception(make_error_code(errc::runtime), "Not supported configuration of map requested"); } ur_map_flags_t Flags = 0; switch (AccessMode) { case access::mode::read: Flags |= UR_MAP_FLAG_READ; break; case access::mode::write: Flags |= UR_MAP_FLAG_WRITE; break; case access::mode::read_write: case access::mode::atomic: Flags = UR_MAP_FLAG_WRITE | UR_MAP_FLAG_READ; break; case access::mode::discard_write: case access::mode::discard_read_write: Flags |= UR_MAP_FLAG_WRITE_INVALIDATE_REGION; break; } AccessOffset[0] *= ElementSize; AccessRange[0] *= ElementSize; // TODO: Handle offset assert(AccessOffset[0] == 0 && "Handle offset"); void *MappedPtr = nullptr; const size_t BytesToMap = AccessRange[0] * AccessRange[1] * AccessRange[2]; const AdapterPtr &Adapter = Queue->getAdapter(); memBufferMapHelper(Adapter, Queue->getHandleRef(), ur::cast(Mem), false, Flags, AccessOffset[0], BytesToMap, DepEvents.size(), DepEvents.data(), &OutEvent, &MappedPtr); return MappedPtr; } void MemoryManager::unmap(SYCLMemObjI *, void *Mem, QueueImplPtr Queue, void *MappedPtr, std::vector DepEvents, ur_event_handle_t &OutEvent) { // Execution on host is not supported here. if (!Queue) { throw exception(make_error_code(errc::runtime), "Not supported configuration of unmap requested"); } // All DepEvents are to the same Context. // Using the adapter of the Queue. const AdapterPtr &Adapter = Queue->getAdapter(); memUnmapHelper(Adapter, Queue->getHandleRef(), ur::cast(Mem), MappedPtr, DepEvents.size(), DepEvents.data(), &OutEvent); } void MemoryManager::copy_usm(const void *SrcMem, QueueImplPtr SrcQueue, size_t Len, void *DstMem, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(SrcQueue && "USM copy must be called with a valid device queue"); if (!Len) { // no-op, but ensure DepEvents will still be waited on if (!DepEvents.empty()) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); SrcQueue->getAdapter()->call( SrcQueue->getHandleRef(), DepEvents.size(), DepEvents.data(), OutEvent); } return; } if (!SrcMem || !DstMem) throw exception(make_error_code(errc::invalid), "NULL pointer argument in memory copy operation."); const AdapterPtr &Adapter = SrcQueue->getAdapter(); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call(SrcQueue->getHandleRef(), /* blocking */ false, DstMem, SrcMem, Len, DepEvents.size(), DepEvents.data(), OutEvent); } void MemoryManager::fill_usm(void *Mem, QueueImplPtr Queue, size_t Length, const std::vector &Pattern, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM fill must be called with a valid device queue"); if (!Length) { // no-op, but ensure DepEvents will still be waited on if (!DepEvents.empty()) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Queue->getAdapter()->call( Queue->getHandleRef(), DepEvents.size(), DepEvents.data(), OutEvent); } return; } if (!Mem) throw exception(make_error_code(errc::invalid), "NULL pointer argument in memory fill operation."); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); const AdapterPtr &Adapter = Queue->getAdapter(); Adapter->call( Queue->getHandleRef(), Mem, Pattern.size(), Pattern.data(), Length, DepEvents.size(), DepEvents.data(), OutEvent); } void MemoryManager::prefetch_usm(void *Mem, QueueImplPtr Queue, size_t Length, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM prefetch must be called with a valid device queue"); const AdapterPtr &Adapter = Queue->getAdapter(); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call(Queue->getHandleRef(), Mem, Length, 0, DepEvents.size(), DepEvents.data(), OutEvent); } void MemoryManager::advise_usm(const void *Mem, QueueImplPtr Queue, size_t Length, ur_usm_advice_flags_t Advice, std::vector /*DepEvents*/, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM advise must be called with a valid device queue"); const AdapterPtr &Adapter = Queue->getAdapter(); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Adapter->call(Queue->getHandleRef(), Mem, Length, Advice, OutEvent); } void MemoryManager::copy_2d_usm(const void *SrcMem, size_t SrcPitch, QueueImplPtr Queue, void *DstMem, size_t DstPitch, size_t Width, size_t Height, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM copy 2d must be called with a valid device queue"); if (Width == 0 || Height == 0) { // no-op, but ensure DepEvents will still be waited on if (!DepEvents.empty()) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Queue->getAdapter()->call( Queue->getHandleRef(), DepEvents.size(), DepEvents.data(), OutEvent); } return; } if (!DstMem || !SrcMem) throw sycl::exception(sycl::make_error_code(errc::invalid), "NULL pointer argument in 2D memory copy operation."); const AdapterPtr &Adapter = Queue->getAdapter(); bool SupportsUSMMemcpy2D = false; Adapter->call( Queue->getContextImplPtr()->getHandleRef(), UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT, sizeof(bool), &SupportsUSMMemcpy2D, nullptr); if (SupportsUSMMemcpy2D) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); // Direct memcpy2D is supported so we use this function. Adapter->call( Queue->getHandleRef(), /*blocking=*/false, DstMem, DstPitch, SrcMem, SrcPitch, Width, Height, DepEvents.size(), DepEvents.data(), OutEvent); return; } // Otherwise we allow the special case where the copy is to or from host. #ifndef NDEBUG context Ctx = createSyclObjFromImpl(Queue->getContextImplPtr()); usm::alloc SrcAllocType = get_pointer_type(SrcMem, Ctx); usm::alloc DstAllocType = get_pointer_type(DstMem, Ctx); bool SrcIsHost = SrcAllocType == usm::alloc::unknown || SrcAllocType == usm::alloc::host; bool DstIsHost = DstAllocType == usm::alloc::unknown || DstAllocType == usm::alloc::host; assert((SrcIsHost || DstIsHost) && "In fallback path for copy_2d_usm either " "source or destination must be on host."); #endif // NDEBUG // The fallback in this case is to insert a copy per row. std::vector CopyEventsManaged; CopyEventsManaged.reserve(Height); // We'll need continuous range of events for a wait later as well. std::vector CopyEvents(Height); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); for (size_t I = 0; I < Height; ++I) { char *DstItBegin = static_cast(DstMem) + I * DstPitch; const char *SrcItBegin = static_cast(SrcMem) + I * SrcPitch; Adapter->call( Queue->getHandleRef(), /* blocking */ false, DstItBegin, SrcItBegin, Width, DepEvents.size(), DepEvents.data(), CopyEvents.data() + I); CopyEventsManaged.emplace_back(CopyEvents[I], Adapter, /*TakeOwnership=*/true); } if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); // Then insert a wait to coalesce the copy events. Queue->getAdapter()->call( Queue->getHandleRef(), CopyEvents.size(), CopyEvents.data(), OutEvent); } void MemoryManager::fill_2d_usm(void *DstMem, QueueImplPtr Queue, size_t Pitch, size_t Width, size_t Height, const std::vector &Pattern, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM fill 2d must be called with a valid device queue"); if (Width == 0 || Height == 0) { // no-op, but ensure DepEvents will still be waited on if (!DepEvents.empty()) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Queue->getAdapter()->call( Queue->getHandleRef(), DepEvents.size(), DepEvents.data(), OutEvent); } return; } if (!DstMem) throw sycl::exception(sycl::make_error_code(errc::invalid), "NULL pointer argument in 2D memory fill operation."); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); const AdapterPtr &Adapter = Queue->getAdapter(); Adapter->call( Queue->getHandleRef(), DstMem, Pitch, Pattern.size(), Pattern.data(), Width, Height, DepEvents.size(), DepEvents.data(), OutEvent); } void MemoryManager::memset_2d_usm(void *DstMem, QueueImplPtr Queue, size_t Pitch, size_t Width, size_t Height, char Value, std::vector DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "USM memset 2d must be called with a valid device queue"); if (Width == 0 || Height == 0) { // no-op, but ensure DepEvents will still be waited on if (!DepEvents.empty()) { if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); Queue->getAdapter()->call( Queue->getHandleRef(), DepEvents.size(), DepEvents.data(), OutEvent); } return; } if (!DstMem) throw sycl::exception( sycl::make_error_code(errc::invalid), "NULL pointer argument in 2D memory memset operation."); if (OutEventImpl != nullptr) OutEventImpl->setHostEnqueueTime(); MemoryManager::fill_2d_usm(DstMem, Queue, Pitch, Width, Height, {static_cast(Value)}, DepEvents, OutEvent, nullptr); } static void memcpyToDeviceGlobalUSM( QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, const void *Src, const std::vector &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "Copy to device global USM must be called with a valid device queue"); // Get or allocate USM memory for the device_global. DeviceGlobalUSMMem &DeviceGlobalUSM = DeviceGlobalEntry->getOrAllocateDeviceGlobalUSM(Queue); void *Dest = DeviceGlobalUSM.getPtr(); // OwnedPiEvent will keep the initialization event alive for the duration // of this function call. OwnedUrEvent ZIEvent = DeviceGlobalUSM.getInitEvent(Queue->getAdapter()); // We may need addtional events, so create a non-const dependency events list // to use if we need to modify it. std::vector AuxDepEventsStorage; const std::vector &ActualDepEvents = ZIEvent ? AuxDepEventsStorage : DepEvents; // If there is a zero-initializer event the memory operation should wait for // it. if (ZIEvent) { AuxDepEventsStorage = DepEvents; AuxDepEventsStorage.push_back(ZIEvent.GetEvent()); } MemoryManager::copy_usm(Src, Queue, NumBytes, reinterpret_cast(Dest) + Offset, ActualDepEvents, OutEvent, OutEventImpl); } static void memcpyFromDeviceGlobalUSM( QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, void *Dest, const std::vector &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { assert(Queue && "Copying from device global USM must be called with a valid " "device queue"); // Get or allocate USM memory for the device_global. Since we are reading from // it, we need it initialized if it has not been yet. DeviceGlobalUSMMem &DeviceGlobalUSM = DeviceGlobalEntry->getOrAllocateDeviceGlobalUSM(Queue); void *Src = DeviceGlobalUSM.getPtr(); // OwnedPiEvent will keep the initialization event alive for the duration // of this function call. OwnedUrEvent ZIEvent = DeviceGlobalUSM.getInitEvent(Queue->getAdapter()); // We may need addtional events, so create a non-const dependency events list // to use if we need to modify it. std::vector AuxDepEventsStorage; const std::vector &ActualDepEvents = ZIEvent ? AuxDepEventsStorage : DepEvents; // If there is a zero-initializer event the memory operation should wait for // it. if (ZIEvent) { AuxDepEventsStorage = DepEvents; AuxDepEventsStorage.push_back(ZIEvent.GetEvent()); } MemoryManager::copy_usm(reinterpret_cast(Src) + Offset, Queue, NumBytes, Dest, ActualDepEvents, OutEvent, OutEventImpl); } static ur_program_handle_t getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry) { assert(DeviceGlobalEntry->MIsDeviceImageScopeDecorated && "device_global is not device image scope decorated."); // If the device global is used in multiple device images we cannot proceed. if (DeviceGlobalEntry->MImageIdentifiers.size() > 1) throw sycl::exception(make_error_code(errc::invalid), "More than one image exists with the device_global."); // If there are no kernels using the device_global we cannot proceed. if (DeviceGlobalEntry->MImageIdentifiers.size() == 0) throw sycl::exception(make_error_code(errc::invalid), "No image exists with the device_global."); // Look for cached programs with the device_global. device Device = Queue->get_device(); ContextImplPtr ContextImpl = Queue->getContextImplPtr(); std::optional CachedProgram = ContextImpl->getProgramForDeviceGlobal(Device, DeviceGlobalEntry); if (CachedProgram) return *CachedProgram; // If there was no cached program, build one. auto Context = createSyclObjFromImpl(ContextImpl); ProgramManager &PM = ProgramManager::getInstance(); RTDeviceBinaryImage &Img = PM.getDeviceImage(DeviceGlobalEntry->MImages, Context, Device); device_image_plain DeviceImage = PM.getDeviceImageFromBinaryImage(&Img, Context, Device); device_image_plain BuiltImage = PM.build(std::move(DeviceImage), {Device}, {}); return getSyclObjImpl(BuiltImage)->get_ur_program_ref(); } static void memcpyToDeviceGlobalDirect(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, const void *Src, const std::vector &DepEvents, ur_event_handle_t *OutEvent) { assert( Queue && "Direct copy to device global must be called with a valid device queue"); ur_program_handle_t Program = getOrBuildProgramForDeviceGlobal(Queue, DeviceGlobalEntry); const AdapterPtr &Adapter = Queue->getAdapter(); Adapter->call( Queue->getHandleRef(), Program, DeviceGlobalEntry->MUniqueId.c_str(), false, NumBytes, Offset, Src, DepEvents.size(), DepEvents.data(), OutEvent); } static void memcpyFromDeviceGlobalDirect(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, void *Dest, const std::vector &DepEvents, ur_event_handle_t *OutEvent) { assert(Queue && "Direct copy from device global must be called with a valid " "device queue"); ur_program_handle_t Program = getOrBuildProgramForDeviceGlobal(Queue, DeviceGlobalEntry); const AdapterPtr &Adapter = Queue->getAdapter(); Adapter->call( Queue->getHandleRef(), Program, DeviceGlobalEntry->MUniqueId.c_str(), false, NumBytes, Offset, Dest, DepEvents.size(), DepEvents.data(), OutEvent); } void MemoryManager::copy_to_device_global( const void *DeviceGlobalPtr, bool IsDeviceImageScoped, QueueImplPtr Queue, size_t NumBytes, size_t Offset, const void *SrcMem, const std::vector &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { DeviceGlobalMapEntry *DGEntry = detail::ProgramManager::getInstance().getDeviceGlobalEntry( DeviceGlobalPtr); assert(DGEntry && DGEntry->MIsDeviceImageScopeDecorated == IsDeviceImageScoped && "Invalid copy operation for device_global."); assert(DGEntry->MDeviceGlobalTSize >= Offset + NumBytes && "Copy to device_global is out of bounds."); if (IsDeviceImageScoped) memcpyToDeviceGlobalDirect(Queue, DGEntry, NumBytes, Offset, SrcMem, DepEvents, OutEvent); else memcpyToDeviceGlobalUSM(Queue, DGEntry, NumBytes, Offset, SrcMem, DepEvents, OutEvent, OutEventImpl); } void MemoryManager::copy_from_device_global( const void *DeviceGlobalPtr, bool IsDeviceImageScoped, QueueImplPtr Queue, size_t NumBytes, size_t Offset, void *DstMem, const std::vector &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl) { DeviceGlobalMapEntry *DGEntry = detail::ProgramManager::getInstance().getDeviceGlobalEntry( DeviceGlobalPtr); assert(DGEntry && DGEntry->MIsDeviceImageScopeDecorated == IsDeviceImageScoped && "Invalid copy operation for device_global."); assert(DGEntry->MDeviceGlobalTSize >= Offset + NumBytes && "Copy from device_global is out of bounds."); if (IsDeviceImageScoped) memcpyFromDeviceGlobalDirect(Queue, DGEntry, NumBytes, Offset, DstMem, DepEvents, OutEvent); else memcpyFromDeviceGlobalUSM(Queue, DGEntry, NumBytes, Offset, DstMem, DepEvents, OutEvent, OutEventImpl); } // Command buffer methods void MemoryManager::ext_oneapi_copyD2D_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, SYCLMemObjI *SYCLMemObj, void *SrcMem, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, void *DstMem, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); (void)DstAccessRange; const AdapterPtr &Adapter = Context->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.XTerm] * SrcElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "Images are not supported in Graphs"); } if (1 == DimDst && 1 == DimSrc) { Adapter->call( CommandBuffer, sycl::detail::ur::cast(SrcMem), sycl::detail::ur::cast(DstMem), SrcXOffBytes, DstXOffBytes, SrcAccessRangeWidthBytes, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); } else { // passing 0 for pitches not allowed. Because clEnqueueCopyBufferRect will // calculate both src and dest pitch using region[0], which is not correct // if src and dest are not the same size. size_t SrcRowPitch = SrcSzWidthBytes; size_t SrcSlicePitch = (DimSrc <= 1) ? SrcSzWidthBytes : SrcSzWidthBytes * SrcSize[SrcPos.YTerm]; size_t DstRowPitch = DstSzWidthBytes; size_t DstSlicePitch = (DimDst <= 1) ? DstSzWidthBytes : DstSzWidthBytes * DstSize[DstPos.YTerm]; ur_rect_offset_t SrcOrigin{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_offset_t DstOrigin{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t Region{SrcAccessRangeWidthBytes, SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; Adapter->call( CommandBuffer, sycl::detail::ur::cast(SrcMem), sycl::detail::ur::cast(DstMem), SrcOrigin, DstOrigin, Region, SrcRowPitch, SrcSlicePitch, DstRowPitch, DstSlicePitch, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); } } void MemoryManager::ext_oneapi_copyD2H_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, SYCLMemObjI *SYCLMemObj, void *SrcMem, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::range<3> SrcAccessRange, sycl::id<3> SrcOffset, unsigned int SrcElemSize, char *DstMem, unsigned int DimDst, sycl::range<3> DstSize, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); const AdapterPtr &Adapter = Context->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.XTerm] * SrcElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "Images are not supported in Graphs"); } if (1 == DimDst && 1 == DimSrc) { ur_result_t Result = Adapter->call_nocheck( CommandBuffer, sycl::detail::ur::cast(SrcMem), SrcXOffBytes, SrcAccessRangeWidthBytes, DstMem + DstXOffBytes, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Device-to-host buffer copy command not supported by graph backend"); } else { Adapter->checkUrResult(Result); } } else { size_t BufferRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes; size_t BufferSlicePitch = (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0; size_t HostRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes; size_t HostSlicePitch = (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0; ur_rect_offset_t BufferOffset{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_offset_t HostOffset{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_region_t RectRegion{SrcAccessRangeWidthBytes, SrcAccessRange[SrcPos.YTerm], SrcAccessRange[SrcPos.ZTerm]}; ur_result_t Result = Adapter->call_nocheck< UrApiKind::urCommandBufferAppendMemBufferReadRectExp>( CommandBuffer, sycl::detail::ur::cast(SrcMem), BufferOffset, HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch, HostSlicePitch, DstMem, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Device-to-host buffer copy command not supported by graph backend"); } else { Adapter->checkUrResult(Result); } } } void MemoryManager::ext_oneapi_copyH2D_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, SYCLMemObjI *SYCLMemObj, char *SrcMem, unsigned int DimSrc, sycl::range<3> SrcSize, sycl::id<3> SrcOffset, unsigned int SrcElemSize, void *DstMem, unsigned int DimDst, sycl::range<3> DstSize, sycl::range<3> DstAccessRange, sycl::id<3> DstOffset, unsigned int DstElemSize, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); const AdapterPtr &Adapter = Context->getAdapter(); detail::SYCLMemObjI::MemObjType MemType = SYCLMemObj->getType(); TermPositions SrcPos, DstPos; prepTermPositions(SrcPos, DimSrc, MemType); prepTermPositions(DstPos, DimDst, MemType); size_t DstXOffBytes = DstOffset[DstPos.XTerm] * DstElemSize; size_t SrcXOffBytes = SrcOffset[SrcPos.XTerm] * SrcElemSize; size_t DstAccessRangeWidthBytes = DstAccessRange[DstPos.XTerm] * DstElemSize; size_t DstSzWidthBytes = DstSize[DstPos.XTerm] * DstElemSize; size_t SrcSzWidthBytes = SrcSize[SrcPos.XTerm] * SrcElemSize; if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "Images are not supported in Graphs"); } if (1 == DimDst && 1 == DimSrc) { ur_result_t Result = Adapter ->call_nocheck( CommandBuffer, sycl::detail::ur::cast(DstMem), DstXOffBytes, DstAccessRangeWidthBytes, SrcMem + SrcXOffBytes, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Host-to-device buffer copy command not supported by graph backend"); } else { Adapter->checkUrResult(Result); } } else { size_t BufferRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes; size_t BufferSlicePitch = (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0; size_t HostRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes; size_t HostSlicePitch = (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0; ur_rect_offset_t BufferOffset{DstXOffBytes, DstOffset[DstPos.YTerm], DstOffset[DstPos.ZTerm]}; ur_rect_offset_t HostOffset{SrcXOffBytes, SrcOffset[SrcPos.YTerm], SrcOffset[SrcPos.ZTerm]}; ur_rect_region_t RectRegion{DstAccessRangeWidthBytes, DstAccessRange[DstPos.YTerm], DstAccessRange[DstPos.ZTerm]}; ur_result_t Result = Adapter->call_nocheck< UrApiKind::urCommandBufferAppendMemBufferWriteRectExp>( CommandBuffer, sycl::detail::ur::cast(DstMem), BufferOffset, HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch, HostSlicePitch, SrcMem, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "Host-to-device buffer copy command not supported by graph backend"); } else { Adapter->checkUrResult(Result); } } } void MemoryManager::ext_oneapi_copy_usm_cmd_buffer( ContextImplPtr Context, const void *SrcMem, ur_exp_command_buffer_handle_t CommandBuffer, size_t Len, void *DstMem, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { if (!SrcMem || !DstMem) throw exception(make_error_code(errc::invalid), "NULL pointer argument in memory copy operation."); const AdapterPtr &Adapter = Context->getAdapter(); ur_result_t Result = Adapter->call_nocheck( CommandBuffer, DstMem, SrcMem, Len, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), "USM copy command not supported by graph backend"); } else { Adapter->checkUrResult(Result); } } void MemoryManager::ext_oneapi_fill_usm_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, void *DstMem, size_t Len, const std::vector &Pattern, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { if (!DstMem) throw exception(make_error_code(errc::invalid), "NULL pointer argument in memory fill operation."); const AdapterPtr &Adapter = Context->getAdapter(); Adapter->call( CommandBuffer, DstMem, Pattern.data(), Pattern.size(), Len, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); } void MemoryManager::ext_oneapi_fill_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, SYCLMemObjI *SYCLMemObj, void *Mem, size_t PatternSize, const unsigned char *Pattern, unsigned int Dim, sycl::range<3> Size, sycl::range<3> AccessRange, sycl::id<3> AccessOffset, unsigned int ElementSize, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { assert(SYCLMemObj && "The SYCLMemObj is nullptr"); const AdapterPtr &Adapter = Context->getAdapter(); if (SYCLMemObj->getType() != detail::SYCLMemObjI::MemObjType::Buffer) { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "Images are not supported in Graphs"); } // 2D and 3D buffers accessors can't have custom range or the data will // likely be discontiguous. bool RangesUsable = (Dim <= 1) || (Size == AccessRange); // For 2D and 3D buffers, the offset must be 0, or the data will be // discontiguous. bool OffsetUsable = (Dim <= 1) || (AccessOffset == sycl::id<3>{0, 0, 0}); size_t RangeMultiplier = AccessRange[0] * AccessRange[1] * AccessRange[2]; if (RangesUsable && OffsetUsable) { Adapter->call( CommandBuffer, ur::cast(Mem), Pattern, PatternSize, AccessOffset[0] * ElementSize, RangeMultiplier * ElementSize, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); return; } // The sycl::handler uses a parallel_for kernel in the case of unusable // Range or Offset, not CG:Fill. So we should not be here. throw exception(make_error_code(errc::runtime), "Not supported configuration of fill requested"); } void MemoryManager::ext_oneapi_prefetch_usm_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, void *Mem, size_t Length, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { const AdapterPtr &Adapter = Context->getAdapter(); Adapter->call( CommandBuffer, Mem, Length, ur_usm_migration_flags_t(0), Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); } void MemoryManager::ext_oneapi_advise_usm_cmd_buffer( sycl::detail::ContextImplPtr Context, ur_exp_command_buffer_handle_t CommandBuffer, const void *Mem, size_t Length, ur_usm_advice_flags_t Advice, std::vector Deps, ur_exp_command_buffer_sync_point_t *OutSyncPoint) { const AdapterPtr &Adapter = Context->getAdapter(); Adapter->call( CommandBuffer, Mem, Length, Advice, Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr); } void MemoryManager::copy_image_bindless( QueueImplPtr Queue, const void *Src, void *Dst, const ur_image_desc_t &SrcDesc, const ur_image_desc_t &DstDesc, const ur_image_format_t &SrcFormat, const ur_image_format_t &DstFormat, const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent, const std::vector &DepEvents, ur_event_handle_t *OutEvent) { assert(Queue && "Copy image bindless must be called with a valid device queue"); assert((Flags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE || Flags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST || Flags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE || Flags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST) && "Invalid flags passed to copy_image_bindless."); if (!Dst || !Src) throw sycl::exception( sycl::make_error_code(errc::invalid), "NULL pointer argument in bindless image copy operation."); const detail::AdapterPtr &Adapter = Queue->getAdapter(); ur_exp_image_copy_region_t CopyRegion{}; CopyRegion.stype = UR_STRUCTURE_TYPE_EXP_IMAGE_COPY_REGION; CopyRegion.copyExtent = CopyExtent; CopyRegion.srcOffset = SrcOffset; CopyRegion.dstOffset = DstOffset; Adapter->call( Queue->getHandleRef(), Src, Dst, &SrcDesc, &DstDesc, &SrcFormat, &DstFormat, &CopyRegion, Flags, DepEvents.size(), DepEvents.data(), OutEvent); } } // namespace detail } // namespace _V1 } // namespace sycl