29 #ifdef XPTI_ENABLE_INSTRUMENTATION
30 #include <xpti/xpti_data_types.h>
31 #include <xpti/xpti_trace_framework.hpp>
35 inline namespace _V1 {
38 #ifdef XPTI_ENABLE_INSTRUMENTATION
39 uint8_t GMemAllocStreamID;
40 xpti::trace_event_data_t *GMemAllocEvent;
48 uint64_t CorrelationID = 0;
49 #ifdef XPTI_ENABLE_INSTRUMENTATION
50 constexpr uint16_t NotificationTraceType =
51 static_cast<uint16_t
>(xpti::trace_point_type_t::mem_alloc_begin);
52 if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) {
53 xpti::mem_alloc_data_t MemAlloc{ObjHandle, 0 , AllocSize,
56 CorrelationID = xptiGetUniqueId();
57 xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType,
58 GMemAllocEvent,
nullptr, CorrelationID, &MemAlloc);
65 size_t AllocSize,
size_t GuardZone,
66 uint64_t CorrelationID) {
72 #ifdef XPTI_ENABLE_INSTRUMENTATION
73 constexpr uint16_t NotificationTraceType =
74 static_cast<uint16_t
>(xpti::trace_point_type_t::mem_alloc_end);
75 if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) {
76 xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, AllocSize, GuardZone};
78 xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType,
79 GMemAllocEvent,
nullptr, CorrelationID, &MemAlloc);
87 uint64_t CorrelationID = 0;
88 #ifdef XPTI_ENABLE_INSTRUMENTATION
89 constexpr uint16_t NotificationTraceType =
90 static_cast<uint16_t
>(xpti::trace_point_type_t::mem_release_begin);
91 if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) {
92 xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, 0 ,
95 CorrelationID = xptiGetUniqueId();
96 xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType,
97 GMemAllocEvent,
nullptr, CorrelationID, &MemAlloc);
100 return CorrelationID;
104 uint64_t CorrelationID) {
108 #ifdef XPTI_ENABLE_INSTRUMENTATION
109 constexpr uint16_t NotificationTraceType =
110 static_cast<uint16_t
>(xpti::trace_point_type_t::mem_release_end);
111 if (xptiCheckTraceEnabled(GMemAllocStreamID, NotificationTraceType)) {
112 xpti::mem_alloc_data_t MemAlloc{ObjHandle, AllocPtr, 0 ,
115 xptiNotifySubscribers(GMemAllocStreamID, NotificationTraceType,
116 GMemAllocEvent,
nullptr, CorrelationID, &MemAlloc);
124 if (!Events.empty()) {
125 const PluginPtr &Plugin = Events[0]->getPlugin();
126 std::vector<ur_event_handle_t> UrEvents(Events.size());
127 std::transform(Events.begin(), Events.end(), UrEvents.begin(),
129 return EventImpl->getHandleRef();
131 if (!UrEvents.empty() && UrEvents[0]) {
132 Plugin->call(urEventWait, UrEvents.size(), &UrEvents[0]);
138 ur_mem_flags_t Flags,
size_t Size,
139 ur_mem_handle_t *RetMem,
140 const ur_buffer_properties_t *Props) {
141 #ifdef XPTI_ENABLE_INSTRUMENTATION
146 #ifdef XPTI_ENABLE_INSTRUMENTATION
149 xpti::utils::finally _{[&] {
151 uintptr_t MemObjID = (uintptr_t)(*RetMem);
152 ur_native_handle_t Ptr = 0;
159 Plugin->call_nocheck(urMemGetNativeHandle, *RetMem,
nullptr,
166 Plugin->call(urMemBufferCreate, Ctx, Flags, Size, Props, RetMem);
174 #ifdef XPTI_ENABLE_INSTRUMENTATION
177 uintptr_t MemObjID = (uintptr_t)(Mem);
180 if (xptiTraceEnabled()) {
181 ur_native_handle_t PtrHandle = 0;
185 Plugin->call_nocheck(urMemGetNativeHandle, Mem,
nullptr,
187 Ptr = (uintptr_t)(PtrHandle);
192 #ifdef XPTI_ENABLE_INSTRUMENTATION
194 xpti::utils::finally _{
197 Plugin->call(urMemRelease, Mem);
202 ur_mem_handle_t Buffer,
bool Blocking,
203 ur_map_flags_t Flags,
size_t Offset,
size_t Size,
204 uint32_t NumEvents,
const ur_event_handle_t *WaitList,
205 ur_event_handle_t *Event,
void **RetMap) {
206 #ifdef XPTI_ENABLE_INSTRUMENTATION
208 uintptr_t MemObjID = (uintptr_t)(Buffer);
212 #ifdef XPTI_ENABLE_INSTRUMENTATION
214 xpti::utils::finally _{[&] {
219 Plugin->call(urEnqueueMemBufferMap, Queue, Buffer, Blocking, Flags, Offset,
220 Size, NumEvents, WaitList, Event, RetMap);
224 ur_mem_handle_t Mem,
void *MappedPtr, uint32_t NumEvents,
225 const ur_event_handle_t *WaitList,
226 ur_event_handle_t *Event) {
227 #ifdef XPTI_ENABLE_INSTRUMENTATION
229 uintptr_t MemObjID = (uintptr_t)(Mem);
230 uintptr_t Ptr = (uintptr_t)(MappedPtr);
234 #ifdef XPTI_ENABLE_INSTRUMENTATION
236 xpti::utils::finally _{[&] {
243 Plugin->call_nocheck(urEventWait, 1, Event);
247 Plugin->call(urEnqueueMemUnmap, Queue, Mem, MappedPtr, NumEvents, WaitList,
254 std::vector<EventImplPtr> DepEvents,
255 ur_event_handle_t &OutEvent) {
261 MemObj->
releaseMem(TargetContext, MemAllocation);
267 if (UserPtr == MemAllocation) {
272 if (!TargetContext) {
277 const PluginPtr &Plugin = TargetContext->getPlugin();
282 bool InitFromUserData,
void *HostPtr,
283 std::vector<EventImplPtr> DepEvents,
284 ur_event_handle_t &OutEvent) {
290 return MemObj->
allocateMem(TargetContext, InitFromUserData, HostPtr,
295 bool HostPtrReadOnly,
size_t Size,
297 std::ignore = HostPtrReadOnly;
312 (void)InteropContext;
314 assert(TargetContext == InteropContext &&
"Expected matching contexts");
316 OutEventToWait = InteropEvent->getHandleRef();
319 if (
nullptr != OutEventToWait) {
320 const PluginPtr &Plugin = InteropEvent->getPlugin();
321 Plugin->call(urEventRetain, OutEventToWait);
327 bool HostPtrReadOnly) {
329 ur_mem_flags_t Result =
330 HostPtrReadOnly ? UR_MEM_FLAG_READ_ONLY : UR_MEM_FLAG_READ_WRITE;
332 Result |= UR_MEM_FLAG_USE_HOST_POINTER;
337 void *UserPtr,
bool HostPtrReadOnly,
338 const ur_image_desc_t &Desc,
339 const ur_image_format_t &Format,
341 ur_mem_flags_t CreationFlags =
344 ur_mem_handle_t NewMem =
nullptr;
345 const PluginPtr &Plugin = TargetContext->getPlugin();
346 Plugin->call(urMemImageCreate, TargetContext->getHandleRef(), CreationFlags,
347 &Format, &Desc, UserPtr, &NewMem);
353 bool HostPtrReadOnly,
const size_t Size,
355 ur_mem_flags_t CreationFlags =
359 CreationFlags |= UR_MEM_FLAG_ALLOC_HOST_POINTER;
361 ur_mem_handle_t NewMem =
nullptr;
362 const PluginPtr &Plugin = TargetContext->getPlugin();
364 ur_buffer_properties_t AllocProps = {UR_STRUCTURE_TYPE_BUFFER_PROPERTIES,
367 void **Next = &AllocProps.pNext;
368 ur_buffer_alloc_location_properties_t LocationProperties = {
369 UR_STRUCTURE_TYPE_BUFFER_ALLOC_LOCATION_PROPERTIES,
nullptr, 0};
371 TargetContext->isBufferLocationSupported()) {
372 LocationProperties.location =
374 .get_buffer_location();
375 *Next = &LocationProperties;
376 Next = &LocationProperties.pNext;
379 ur_buffer_channel_properties_t ChannelProperties = {
380 UR_STRUCTURE_TYPE_BUFFER_CHANNEL_PROPERTIES,
nullptr, 0};
382 ChannelProperties.channel =
384 *Next = &ChannelProperties;
388 Size, &NewMem, &AllocProps);
394 bool HostPtrReadOnly,
size_t Size,
398 ur_event_handle_t &OutEventToWait) {
403 else if (UserPtr && InteropContext)
406 InteropContext, PropsList, OutEventToWait);
416 bool HostPtrReadOnly,
size_t Size,
const ur_image_desc_t &Desc,
417 const ur_image_format_t &Format,
const EventImplPtr &InteropEvent,
419 ur_event_handle_t &OutEventToWait) {
423 if (UserPtr && InteropContext)
425 InteropContext, PropsList, OutEventToWait);
431 void *ParentMemObj,
size_t ElemSize,
433 std::vector<EventImplPtr> DepEvents,
434 ur_event_handle_t &OutEvent) {
439 return static_cast<void *
>(
static_cast<char *
>(ParentMemObj) + Offset);
441 size_t SizeInBytes = ElemSize;
442 for (
size_t I = 0; I < 3; ++I)
443 SizeInBytes *= Range[I];
445 ur_result_t Error = UR_RESULT_SUCCESS;
446 ur_buffer_region_t Region = {UR_STRUCTURE_TYPE_BUFFER_REGION,
nullptr, Offset,
448 ur_mem_handle_t NewMem;
449 const PluginPtr &Plugin = TargetContext->getPlugin();
450 Error = Plugin->call_nocheck(
451 urMemBufferPartition, ur::cast<ur_mem_handle_t>(ParentMemObj),
452 UR_MEM_FLAG_READ_WRITE, UR_BUFFER_CREATE_TYPE_REGION, &Region, &NewMem);
453 if (Error == UR_RESULT_ERROR_MISALIGNED_SUB_BUFFER_OFFSET)
456 "Specified offset of the sub-buffer being constructed is not "
457 "a multiple of the memory base address alignment"),
460 Plugin->checkUrResult(Error);
479 if (Type == detail::SYCLMemObjI::MemObjType::Buffer) {
498 unsigned int SrcElemSize, ur_mem_handle_t DstMem,
501 unsigned int DstElemSize, std::vector<ur_event_handle_t> DepEvents,
502 ur_event_handle_t &OutEvent,
504 (void)SrcAccessRange;
505 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
506 assert(TgtQueue &&
"Destination mem object queue must be not nullptr");
508 const ur_queue_handle_t Queue = TgtQueue->getHandleRef();
509 const PluginPtr &Plugin = TgtQueue->getPlugin();
516 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
517 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
518 size_t DstAccessRangeWidthBytes = DstAccessRange[DstPos.
XTerm] * DstElemSize;
519 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
520 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
522 if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) {
523 if (1 == DimDst && 1 == DimSrc) {
524 if (OutEventImpl !=
nullptr)
525 OutEventImpl->setHostEnqueueTime();
526 Plugin->call(urEnqueueMemBufferWrite, Queue, DstMem,
528 DstAccessRangeWidthBytes, SrcMem + SrcXOffBytes,
529 DepEvents.size(), DepEvents.data(), &OutEvent);
531 size_t BufferRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes;
532 size_t BufferSlicePitch =
533 (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.
YTerm] : 0;
534 size_t HostRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes;
535 size_t HostSlicePitch =
536 (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.
YTerm] : 0;
538 ur_rect_offset_t BufferOffset{DstXOffBytes, DstOffset[DstPos.
YTerm],
539 DstOffset[DstPos.
ZTerm]};
540 ur_rect_offset_t HostOffset{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
541 SrcOffset[SrcPos.
ZTerm]};
542 ur_rect_region_t RectRegion{DstAccessRangeWidthBytes,
543 DstAccessRange[DstPos.
YTerm],
544 DstAccessRange[DstPos.
ZTerm]};
545 if (OutEventImpl !=
nullptr)
546 OutEventImpl->setHostEnqueueTime();
547 Plugin->call(urEnqueueMemBufferWriteRect, Queue, DstMem,
548 false, BufferOffset, HostOffset,
549 RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch,
550 HostSlicePitch, SrcMem, DepEvents.size(), DepEvents.data(),
554 size_t InputRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes;
555 size_t InputSlicePitch =
556 (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.YTerm] : 0;
558 ur_rect_offset_t Origin{DstOffset[DstPos.XTerm], DstOffset[DstPos.YTerm],
559 DstOffset[DstPos.ZTerm]};
560 ur_rect_region_t Region{DstAccessRange[DstPos.XTerm],
561 DstAccessRange[DstPos.YTerm],
562 DstAccessRange[DstPos.ZTerm]};
563 if (OutEventImpl !=
nullptr)
564 OutEventImpl->setHostEnqueueTime();
565 Plugin->call(urEnqueueMemImageWrite, Queue, DstMem,
566 false, Origin, Region, InputRowPitch,
567 InputSlicePitch, SrcMem, DepEvents.size(), DepEvents.data(),
578 unsigned int DstElemSize, std::vector<ur_event_handle_t> DepEvents,
579 ur_event_handle_t &OutEvent,
581 (void)DstAccessRange;
582 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
583 assert(SrcQueue &&
"Source mem object queue is expected to be not nullptr");
585 const ur_queue_handle_t Queue = SrcQueue->getHandleRef();
586 const PluginPtr &Plugin = SrcQueue->getPlugin();
599 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
600 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
601 size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.
XTerm] * SrcElemSize;
602 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
603 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
605 if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) {
606 if (1 == DimDst && 1 == DimSrc) {
607 if (OutEventImpl !=
nullptr)
608 OutEventImpl->setHostEnqueueTime();
609 Plugin->call(urEnqueueMemBufferRead, Queue, SrcMem,
611 SrcAccessRangeWidthBytes, DstMem + DstXOffBytes,
612 DepEvents.size(), DepEvents.data(), &OutEvent);
614 size_t BufferRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes;
615 size_t BufferSlicePitch =
616 (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.
YTerm] : 0;
617 size_t HostRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes;
618 size_t HostSlicePitch =
619 (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.
YTerm] : 0;
621 ur_rect_offset_t BufferOffset{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
622 SrcOffset[SrcPos.
ZTerm]};
623 ur_rect_offset_t HostOffset{DstXOffBytes, DstOffset[DstPos.
YTerm],
624 DstOffset[DstPos.
ZTerm]};
625 ur_rect_region_t RectRegion{SrcAccessRangeWidthBytes,
626 SrcAccessRange[SrcPos.
YTerm],
627 SrcAccessRange[SrcPos.
ZTerm]};
628 if (OutEventImpl !=
nullptr)
629 OutEventImpl->setHostEnqueueTime();
630 Plugin->call(urEnqueueMemBufferReadRect, Queue, SrcMem,
631 false, BufferOffset, HostOffset,
632 RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch,
633 HostSlicePitch, DstMem, DepEvents.size(), DepEvents.data(),
637 size_t RowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes;
639 (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.YTerm] : 0;
641 ur_rect_offset_t Offset{SrcOffset[SrcPos.XTerm], SrcOffset[SrcPos.YTerm],
642 SrcOffset[SrcPos.ZTerm]};
643 ur_rect_region_t Region{SrcAccessRange[SrcPos.XTerm],
644 SrcAccessRange[SrcPos.YTerm],
645 SrcAccessRange[SrcPos.ZTerm]};
646 if (OutEventImpl !=
nullptr)
647 OutEventImpl->setHostEnqueueTime();
648 Plugin->call(urEnqueueMemImageRead, Queue, SrcMem,
false, Offset, Region,
649 RowPitch, SlicePitch, DstMem, DepEvents.size(),
650 DepEvents.data(), &OutEvent);
657 unsigned int SrcElemSize, ur_mem_handle_t DstMem,
QueueImplPtr,
660 std::vector<ur_event_handle_t> DepEvents,
661 ur_event_handle_t &OutEvent,
663 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
664 assert(SrcQueue &&
"Source mem object and target mem object queues are "
665 "expected to be not nullptr");
667 const ur_queue_handle_t Queue = SrcQueue->getHandleRef();
668 const PluginPtr &Plugin = SrcQueue->getPlugin();
675 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
676 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
677 size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.
XTerm] * SrcElemSize;
678 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
679 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
681 if (MemType == detail::SYCLMemObjI::MemObjType::Buffer) {
682 if (1 == DimDst && 1 == DimSrc) {
683 if (OutEventImpl !=
nullptr)
684 OutEventImpl->setHostEnqueueTime();
685 Plugin->call(urEnqueueMemBufferCopy, Queue, SrcMem, DstMem, SrcXOffBytes,
686 DstXOffBytes, SrcAccessRangeWidthBytes, DepEvents.size(),
687 DepEvents.data(), &OutEvent);
692 size_t SrcRowPitch = SrcSzWidthBytes;
693 size_t SrcSlicePitch = (DimSrc <= 1)
695 : SrcSzWidthBytes * SrcSize[SrcPos.
YTerm];
696 size_t DstRowPitch = DstSzWidthBytes;
697 size_t DstSlicePitch = (DimDst <= 1)
699 : DstSzWidthBytes * DstSize[DstPos.
YTerm];
701 ur_rect_offset_t SrcOrigin{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
702 SrcOffset[SrcPos.
ZTerm]};
703 ur_rect_offset_t DstOrigin{DstXOffBytes, DstOffset[DstPos.
YTerm],
704 DstOffset[DstPos.
ZTerm]};
705 ur_rect_region_t Region{SrcAccessRangeWidthBytes,
706 SrcAccessRange[SrcPos.
YTerm],
707 SrcAccessRange[SrcPos.
ZTerm]};
708 if (OutEventImpl !=
nullptr)
709 OutEventImpl->setHostEnqueueTime();
710 Plugin->call(urEnqueueMemBufferCopyRect, Queue, SrcMem, DstMem, SrcOrigin,
711 DstOrigin, Region, SrcRowPitch, SrcSlicePitch, DstRowPitch,
712 DstSlicePitch, DepEvents.size(), DepEvents.data(),
716 ur_rect_offset_t SrcOrigin{SrcOffset[SrcPos.
XTerm], SrcOffset[SrcPos.
YTerm],
717 SrcOffset[SrcPos.
ZTerm]};
718 ur_rect_offset_t DstOrigin{DstOffset[DstPos.
XTerm], DstOffset[DstPos.
YTerm],
719 DstOffset[DstPos.
ZTerm]};
720 ur_rect_region_t Region{SrcAccessRange[SrcPos.
XTerm],
721 SrcAccessRange[SrcPos.
YTerm],
722 SrcAccessRange[SrcPos.
ZTerm]};
723 if (OutEventImpl !=
nullptr)
724 OutEventImpl->setHostEnqueueTime();
725 Plugin->call(urEnqueueMemImageCopy, Queue, SrcMem, DstMem, SrcOrigin,
726 DstOrigin, Region, DepEvents.size(), DepEvents.data(),
737 unsigned int DstElemSize, std::vector<ur_event_handle_t>,
739 if ((DimSrc != 1 || DimDst != 1) &&
740 (SrcOffset !=
id<3>{0, 0, 0} || DstOffset !=
id<3>{0, 0, 0} ||
741 SrcSize != SrcAccessRange || DstSize != DstAccessRange)) {
743 "Not supported configuration of memcpy requested");
746 SrcMem += SrcOffset[0] * SrcElemSize;
747 DstMem += DstOffset[0] * DstElemSize;
749 if (SrcMem == DstMem)
753 SrcAccessRange[0] * SrcElemSize * SrcAccessRange[1] * SrcAccessRange[2];
754 std::memcpy(DstMem, SrcMem, BytesToCopy);
762 sycl::id<3> SrcOffset,
unsigned int SrcElemSize,
void *DstMem,
765 unsigned int DstElemSize, std::vector<ur_event_handle_t> DepEvents,
770 copyH2H(SYCLMemObj, (
char *)SrcMem,
nullptr, DimSrc, SrcSize,
771 SrcAccessRange, SrcOffset, SrcElemSize, (
char *)DstMem,
nullptr,
772 DimDst, DstSize, DstAccessRange, DstOffset, DstElemSize,
773 std::move(DepEvents), OutEvent, OutEventImpl);
775 copyH2D(SYCLMemObj, (
char *)SrcMem,
nullptr, DimSrc, SrcSize,
776 SrcAccessRange, SrcOffset, SrcElemSize,
777 ur::cast<ur_mem_handle_t>(DstMem), std::move(TgtQueue), DimDst,
778 DstSize, DstAccessRange, DstOffset, DstElemSize,
779 std::move(DepEvents), OutEvent, OutEventImpl);
782 copyD2H(SYCLMemObj, ur::cast<ur_mem_handle_t>(SrcMem),
783 std::move(SrcQueue), DimSrc, SrcSize, SrcAccessRange, SrcOffset,
784 SrcElemSize, (
char *)DstMem,
nullptr, DimDst, DstSize,
785 DstAccessRange, DstOffset, DstElemSize, std::move(DepEvents),
786 OutEvent, OutEventImpl);
788 copyD2D(SYCLMemObj, ur::cast<ur_mem_handle_t>(SrcMem),
789 std::move(SrcQueue), DimSrc, SrcSize, SrcAccessRange, SrcOffset,
790 SrcElemSize, ur::cast<ur_mem_handle_t>(DstMem),
791 std::move(TgtQueue), DimDst, DstSize, DstAccessRange, DstOffset,
792 DstElemSize, std::move(DepEvents), OutEvent, OutEventImpl);
797 size_t PatternSize,
const unsigned char *Pattern,
800 unsigned int ElementSize,
801 std::vector<ur_event_handle_t> DepEvents,
802 ur_event_handle_t &OutEvent,
804 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
805 assert(Queue &&
"Fill should be called only with a valid device queue");
807 const PluginPtr &Plugin = Queue->getPlugin();
809 if (SYCLMemObj->
getType() == detail::SYCLMemObjI::MemObjType::Buffer) {
810 if (OutEventImpl !=
nullptr)
811 OutEventImpl->setHostEnqueueTime();
815 bool RangesUsable = (Dim <= 1) || (MemRange == AccRange);
818 bool OffsetUsable = (Dim <= 1) || (Offset ==
sycl::id<3>{0, 0, 0});
819 size_t RangeMultiplier = AccRange[0] * AccRange[1] * AccRange[2];
821 if (RangesUsable && OffsetUsable) {
822 Plugin->call(urEnqueueMemBufferFill, Queue->getHandleRef(),
823 ur::cast<ur_mem_handle_t>(Mem), Pattern, PatternSize,
824 Offset[0] * ElementSize, RangeMultiplier * ElementSize,
825 DepEvents.size(), DepEvents.data(), &OutEvent);
831 "Not supported configuration of fill requested");
833 if (OutEventImpl !=
nullptr)
834 OutEventImpl->setHostEnqueueTime();
839 "Fill operation not supported for the given mem object");
846 unsigned int ElementSize,
847 std::vector<ur_event_handle_t> DepEvents,
848 ur_event_handle_t &OutEvent) {
851 "Not supported configuration of map requested");
854 ur_map_flags_t Flags = 0;
857 case access::mode::read:
858 Flags |= UR_MAP_FLAG_READ;
861 Flags |= UR_MAP_FLAG_WRITE;
864 case access::mode::atomic:
865 Flags = UR_MAP_FLAG_WRITE | UR_MAP_FLAG_READ;
867 case access::mode::discard_write:
868 case access::mode::discard_read_write:
869 Flags |= UR_MAP_FLAG_WRITE_INVALIDATE_REGION;
873 AccessOffset[0] *= ElementSize;
874 AccessRange[0] *= ElementSize;
877 assert(AccessOffset[0] == 0 &&
"Handle offset");
879 void *MappedPtr =
nullptr;
880 const size_t BytesToMap = AccessRange[0] * AccessRange[1] * AccessRange[2];
881 const PluginPtr &Plugin = Queue->getPlugin();
883 ur::cast<ur_mem_handle_t>(Mem),
false, Flags,
884 AccessOffset[0], BytesToMap, DepEvents.size(),
885 DepEvents.data(), &OutEvent, &MappedPtr);
891 std::vector<ur_event_handle_t> DepEvents,
892 ur_event_handle_t &OutEvent) {
897 "Not supported configuration of unmap requested");
902 const PluginPtr &Plugin = Queue->getPlugin();
903 memUnmapHelper(Plugin, Queue->getHandleRef(), ur::cast<ur_mem_handle_t>(Mem),
904 MappedPtr, DepEvents.size(), DepEvents.data(), &OutEvent);
907 void MemoryManager::copy_usm(
const void *SrcMem,
QueueImplPtr SrcQueue,
908 size_t Len,
void *DstMem,
909 std::vector<ur_event_handle_t> DepEvents,
910 ur_event_handle_t *OutEvent,
912 assert(SrcQueue &&
"USM copy must be called with a valid device queue");
914 if (!DepEvents.empty()) {
915 if (OutEventImpl !=
nullptr)
916 OutEventImpl->setHostEnqueueTime();
917 SrcQueue->getPlugin()->call(urEnqueueEventsWait, SrcQueue->getHandleRef(),
918 DepEvents.size(), DepEvents.data(), OutEvent);
923 if (!SrcMem || !DstMem)
925 "NULL pointer argument in memory copy operation.");
927 const PluginPtr &Plugin = SrcQueue->getPlugin();
928 if (OutEventImpl !=
nullptr)
929 OutEventImpl->setHostEnqueueTime();
930 Plugin->call(urEnqueueUSMMemcpy, SrcQueue->getHandleRef(),
931 false, DstMem, SrcMem, Len, DepEvents.size(),
932 DepEvents.data(), OutEvent);
935 void MemoryManager::fill_usm(
void *Mem,
QueueImplPtr Queue,
size_t Length,
936 const std::vector<unsigned char> &Pattern,
937 std::vector<ur_event_handle_t> DepEvents,
938 ur_event_handle_t *OutEvent,
940 assert(Queue &&
"USM fill must be called with a valid device queue");
942 if (!DepEvents.empty()) {
943 if (OutEventImpl !=
nullptr)
944 OutEventImpl->setHostEnqueueTime();
945 Queue->getPlugin()->call(urEnqueueEventsWait, Queue->getHandleRef(),
946 DepEvents.size(), DepEvents.data(), OutEvent);
953 "NULL pointer argument in memory fill operation.");
954 if (OutEventImpl !=
nullptr)
955 OutEventImpl->setHostEnqueueTime();
956 const PluginPtr &Plugin = Queue->getPlugin();
957 Plugin->call(urEnqueueUSMFill, Queue->getHandleRef(), Mem, Pattern.size(),
958 Pattern.data(), Length, DepEvents.size(), DepEvents.data(),
962 void MemoryManager::prefetch_usm(
void *Mem,
QueueImplPtr Queue,
size_t Length,
963 std::vector<ur_event_handle_t> DepEvents,
964 ur_event_handle_t *OutEvent,
966 assert(Queue &&
"USM prefetch must be called with a valid device queue");
967 const PluginPtr &Plugin = Queue->getPlugin();
968 if (OutEventImpl !=
nullptr)
969 OutEventImpl->setHostEnqueueTime();
970 Plugin->call(urEnqueueUSMPrefetch, Queue->getHandleRef(), Mem, Length, 0,
971 DepEvents.size(), DepEvents.data(), OutEvent);
975 size_t Length, ur_usm_advice_flags_t Advice,
976 std::vector<ur_event_handle_t> ,
977 ur_event_handle_t *OutEvent,
979 assert(Queue &&
"USM advise must be called with a valid device queue");
980 const PluginPtr &Plugin = Queue->getPlugin();
981 if (OutEventImpl !=
nullptr)
982 OutEventImpl->setHostEnqueueTime();
983 Plugin->call(urEnqueueUSMAdvise, Queue->getHandleRef(), Mem, Length, Advice,
987 void MemoryManager::copy_2d_usm(
const void *SrcMem,
size_t SrcPitch,
989 size_t DstPitch,
size_t Width,
size_t Height,
990 std::vector<ur_event_handle_t> DepEvents,
991 ur_event_handle_t *OutEvent,
993 assert(Queue &&
"USM copy 2d must be called with a valid device queue");
994 if (Width == 0 || Height == 0) {
996 if (!DepEvents.empty()) {
997 if (OutEventImpl !=
nullptr)
998 OutEventImpl->setHostEnqueueTime();
999 Queue->getPlugin()->call(urEnqueueEventsWait, Queue->getHandleRef(),
1000 DepEvents.size(), DepEvents.data(), OutEvent);
1005 if (!DstMem || !SrcMem)
1007 "NULL pointer argument in 2D memory copy operation.");
1009 const PluginPtr &Plugin = Queue->getPlugin();
1011 bool SupportsUSMMemcpy2D =
false;
1012 Plugin->call(urContextGetInfo, Queue->getContextImplPtr()->getHandleRef(),
1013 UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT,
sizeof(
bool),
1014 &SupportsUSMMemcpy2D,
nullptr);
1016 if (SupportsUSMMemcpy2D) {
1017 if (OutEventImpl !=
nullptr)
1018 OutEventImpl->setHostEnqueueTime();
1020 Plugin->call(urEnqueueUSMMemcpy2D, Queue->getHandleRef(),
1021 false, DstMem, DstPitch, SrcMem, SrcPitch, Width,
1022 Height, DepEvents.size(), DepEvents.data(), OutEvent);
1028 context Ctx = createSyclObjFromImpl<context>(Queue->getContextImplPtr());
1032 SrcAllocType == usm::alloc::unknown || SrcAllocType == usm::alloc::host;
1034 DstAllocType == usm::alloc::unknown || DstAllocType == usm::alloc::host;
1035 assert((SrcIsHost || DstIsHost) &&
"In fallback path for copy_2d_usm either "
1036 "source or destination must be on host.");
1040 std::vector<OwnedUrEvent> CopyEventsManaged;
1041 CopyEventsManaged.reserve(Height);
1043 std::vector<ur_event_handle_t> CopyEvents(Height);
1044 if (OutEventImpl !=
nullptr)
1045 OutEventImpl->setHostEnqueueTime();
1047 for (
size_t I = 0; I < Height; ++I) {
1048 char *DstItBegin =
static_cast<char *
>(DstMem) + I * DstPitch;
1049 const char *SrcItBegin =
static_cast<const char *
>(SrcMem) + I * SrcPitch;
1050 Plugin->call(urEnqueueUSMMemcpy, Queue->getHandleRef(),
1051 false, DstItBegin, SrcItBegin, Width,
1052 DepEvents.size(), DepEvents.data(), CopyEvents.data() + I);
1053 CopyEventsManaged.emplace_back(CopyEvents[I], Plugin,
1056 if (OutEventImpl !=
nullptr)
1057 OutEventImpl->setHostEnqueueTime();
1059 Queue->getPlugin()->call(urEnqueueEventsWait, Queue->getHandleRef(),
1060 CopyEvents.size(), CopyEvents.data(), OutEvent);
1063 void MemoryManager::fill_2d_usm(
void *DstMem,
QueueImplPtr Queue,
size_t Pitch,
1064 size_t Width,
size_t Height,
1065 const std::vector<unsigned char> &Pattern,
1066 std::vector<ur_event_handle_t> DepEvents,
1067 ur_event_handle_t *OutEvent,
1069 assert(Queue &&
"USM fill 2d must be called with a valid device queue");
1070 if (Width == 0 || Height == 0) {
1072 if (!DepEvents.empty()) {
1073 if (OutEventImpl !=
nullptr)
1074 OutEventImpl->setHostEnqueueTime();
1075 Queue->getPlugin()->call(urEnqueueEventsWait, Queue->getHandleRef(),
1076 DepEvents.size(), DepEvents.data(), OutEvent);
1083 "NULL pointer argument in 2D memory fill operation.");
1084 if (OutEventImpl !=
nullptr)
1085 OutEventImpl->setHostEnqueueTime();
1086 const PluginPtr &Plugin = Queue->getPlugin();
1087 Plugin->call(urEnqueueUSMFill2D, Queue->getHandleRef(), DstMem, Pitch,
1088 Pattern.size(), Pattern.data(), Width, Height, DepEvents.size(),
1089 DepEvents.data(), OutEvent);
1093 size_t Pitch,
size_t Width,
size_t Height,
1095 std::vector<ur_event_handle_t> DepEvents,
1096 ur_event_handle_t *OutEvent,
1098 assert(Queue &&
"USM memset 2d must be called with a valid device queue");
1099 if (Width == 0 || Height == 0) {
1101 if (!DepEvents.empty()) {
1102 if (OutEventImpl !=
nullptr)
1103 OutEventImpl->setHostEnqueueTime();
1104 Queue->getPlugin()->call(urEnqueueEventsWait, Queue->getHandleRef(),
1105 DepEvents.size(), DepEvents.data(), OutEvent);
1113 "NULL pointer argument in 2D memory memset operation.");
1114 if (OutEventImpl !=
nullptr)
1115 OutEventImpl->setHostEnqueueTime();
1116 MemoryManager::fill_2d_usm(DstMem, Queue, Pitch, Width, Height,
1117 {
static_cast<unsigned char>(Value)}, DepEvents,
1123 size_t NumBytes,
size_t Offset,
const void *Src,
1124 const std::vector<ur_event_handle_t> &DepEvents,
1127 "Copy to device global USM must be called with a valid device queue");
1131 void *Dest = DeviceGlobalUSM.
getPtr();
1139 std::vector<ur_event_handle_t> AuxDepEventsStorage;
1140 const std::vector<ur_event_handle_t> &ActualDepEvents =
1141 ZIEvent ? AuxDepEventsStorage : DepEvents;
1146 AuxDepEventsStorage = DepEvents;
1147 AuxDepEventsStorage.push_back(ZIEvent.
GetEvent());
1150 MemoryManager::copy_usm(Src, Queue, NumBytes,
1151 reinterpret_cast<char *
>(Dest) + Offset,
1152 ActualDepEvents, OutEvent, OutEventImpl);
1157 size_t NumBytes,
size_t Offset,
void *Dest,
1158 const std::vector<ur_event_handle_t> &DepEvents,
1160 assert(Queue &&
"Copying from device global USM must be called with a valid "
1166 void *Src = DeviceGlobalUSM.
getPtr();
1174 std::vector<ur_event_handle_t> AuxDepEventsStorage;
1175 const std::vector<ur_event_handle_t> &ActualDepEvents =
1176 ZIEvent ? AuxDepEventsStorage : DepEvents;
1181 AuxDepEventsStorage = DepEvents;
1182 AuxDepEventsStorage.push_back(ZIEvent.
GetEvent());
1185 MemoryManager::copy_usm(
reinterpret_cast<const char *
>(Src) + Offset, Queue,
1186 NumBytes, Dest, ActualDepEvents, OutEvent,
1190 static ur_program_handle_t
1194 "device_global is not device image scope decorated.");
1199 "More than one image exists with the device_global.");
1204 "No image exists with the device_global.");
1207 device Device = Queue->get_device();
1209 std::optional<ur_program_handle_t> CachedProgram =
1210 ContextImpl->getProgramForDeviceGlobal(Device, DeviceGlobalEntry);
1212 return *CachedProgram;
1215 auto Context = createSyclObjFromImpl<context>(ContextImpl);
1228 size_t NumBytes,
size_t Offset,
const void *Src,
1229 const std::vector<ur_event_handle_t> &DepEvents,
1230 ur_event_handle_t *OutEvent) {
1233 "Direct copy to device global must be called with a valid device queue");
1234 ur_program_handle_t Program =
1236 const PluginPtr &Plugin = Queue->getPlugin();
1237 Plugin->call(urEnqueueDeviceGlobalVariableWrite, Queue->getHandleRef(),
1238 Program, DeviceGlobalEntry->
MUniqueId.c_str(),
false, NumBytes,
1239 Offset, Src, DepEvents.size(), DepEvents.data(), OutEvent);
1245 size_t NumBytes,
size_t Offset,
void *Dest,
1246 const std::vector<ur_event_handle_t> &DepEvents,
1247 ur_event_handle_t *OutEvent) {
1248 assert(Queue &&
"Direct copy from device global must be called with a valid "
1250 ur_program_handle_t Program =
1252 const PluginPtr &Plugin = Queue->getPlugin();
1253 Plugin->call(urEnqueueDeviceGlobalVariableRead, Queue->getHandleRef(),
1254 Program, DeviceGlobalEntry->
MUniqueId.c_str(),
false, NumBytes,
1255 Offset, Dest, DepEvents.size(), DepEvents.data(), OutEvent);
1258 void MemoryManager::copy_to_device_global(
1259 const void *DeviceGlobalPtr,
bool IsDeviceImageScoped,
QueueImplPtr Queue,
1260 size_t NumBytes,
size_t Offset,
const void *SrcMem,
1261 const std::vector<ur_event_handle_t> &DepEvents,
1264 detail::ProgramManager::getInstance().getDeviceGlobalEntry(
1268 "Invalid copy operation for device_global.");
1270 "Copy to device_global is out of bounds.");
1272 if (IsDeviceImageScoped)
1274 DepEvents, OutEvent);
1277 OutEvent, OutEventImpl);
1280 void MemoryManager::copy_from_device_global(
1281 const void *DeviceGlobalPtr,
bool IsDeviceImageScoped,
QueueImplPtr Queue,
1282 size_t NumBytes,
size_t Offset,
void *DstMem,
1283 const std::vector<ur_event_handle_t> &DepEvents,
1286 detail::ProgramManager::getInstance().getDeviceGlobalEntry(
1290 "Invalid copy operation for device_global.");
1292 "Copy from device_global is out of bounds.");
1294 if (IsDeviceImageScoped)
1296 DepEvents, OutEvent);
1299 DepEvents, OutEvent, OutEventImpl);
1303 void MemoryManager::ext_oneapi_copyD2D_cmd_buffer(
1305 ur_exp_command_buffer_handle_t CommandBuffer,
SYCLMemObjI *SYCLMemObj,
1308 unsigned int SrcElemSize,
void *DstMem,
unsigned int DimDst,
1311 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1312 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1313 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
1314 (void)DstAccessRange;
1316 const PluginPtr &Plugin = Context->getPlugin();
1323 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
1324 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
1325 size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.
XTerm] * SrcElemSize;
1326 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
1327 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
1329 if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) {
1331 "Images are not supported in Graphs");
1334 if (1 == DimDst && 1 == DimSrc) {
1335 Plugin->call(urCommandBufferAppendMemBufferCopyExp, CommandBuffer,
1336 sycl::detail::ur::cast<ur_mem_handle_t>(SrcMem),
1337 sycl::detail::ur::cast<ur_mem_handle_t>(DstMem), SrcXOffBytes,
1338 DstXOffBytes, SrcAccessRangeWidthBytes, Deps.size(),
1339 Deps.data(), OutSyncPoint);
1344 size_t SrcRowPitch = SrcSzWidthBytes;
1345 size_t SrcSlicePitch = (DimSrc <= 1)
1347 : SrcSzWidthBytes * SrcSize[SrcPos.
YTerm];
1348 size_t DstRowPitch = DstSzWidthBytes;
1349 size_t DstSlicePitch = (DimDst <= 1)
1351 : DstSzWidthBytes * DstSize[DstPos.
YTerm];
1353 ur_rect_offset_t SrcOrigin{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
1354 SrcOffset[SrcPos.
ZTerm]};
1355 ur_rect_offset_t DstOrigin{DstXOffBytes, DstOffset[DstPos.
YTerm],
1356 DstOffset[DstPos.
ZTerm]};
1357 ur_rect_region_t Region{SrcAccessRangeWidthBytes,
1358 SrcAccessRange[SrcPos.
YTerm],
1359 SrcAccessRange[SrcPos.
ZTerm]};
1361 Plugin->call(urCommandBufferAppendMemBufferCopyRectExp, CommandBuffer,
1362 sycl::detail::ur::cast<ur_mem_handle_t>(SrcMem),
1363 sycl::detail::ur::cast<ur_mem_handle_t>(DstMem), SrcOrigin,
1364 DstOrigin, Region, SrcRowPitch, SrcSlicePitch, DstRowPitch,
1365 DstSlicePitch, Deps.size(), Deps.data(), OutSyncPoint);
1369 void MemoryManager::ext_oneapi_copyD2H_cmd_buffer(
1371 ur_exp_command_buffer_handle_t CommandBuffer,
SYCLMemObjI *SYCLMemObj,
1374 unsigned int SrcElemSize,
char *DstMem,
unsigned int DimDst,
1376 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1377 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1378 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
1380 const PluginPtr &Plugin = Context->getPlugin();
1387 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
1388 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
1389 size_t SrcAccessRangeWidthBytes = SrcAccessRange[SrcPos.
XTerm] * SrcElemSize;
1390 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
1391 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
1393 if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) {
1395 "Images are not supported in Graphs");
1398 if (1 == DimDst && 1 == DimSrc) {
1399 ur_result_t Result = Plugin->call_nocheck(
1400 urCommandBufferAppendMemBufferReadExp, CommandBuffer,
1401 sycl::detail::ur::cast<ur_mem_handle_t>(SrcMem), SrcXOffBytes,
1402 SrcAccessRangeWidthBytes, DstMem + DstXOffBytes, Deps.size(),
1403 Deps.data(), OutSyncPoint);
1405 if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
1408 "Device-to-host buffer copy command not supported by graph backend");
1410 Plugin->checkUrResult(Result);
1413 size_t BufferRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes;
1414 size_t BufferSlicePitch =
1415 (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.
YTerm] : 0;
1416 size_t HostRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes;
1417 size_t HostSlicePitch =
1418 (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.
YTerm] : 0;
1420 ur_rect_offset_t BufferOffset{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
1421 SrcOffset[SrcPos.
ZTerm]};
1422 ur_rect_offset_t HostOffset{DstXOffBytes, DstOffset[DstPos.
YTerm],
1423 DstOffset[DstPos.
ZTerm]};
1424 ur_rect_region_t RectRegion{SrcAccessRangeWidthBytes,
1425 SrcAccessRange[SrcPos.
YTerm],
1426 SrcAccessRange[SrcPos.
ZTerm]};
1428 ur_result_t Result = Plugin->call_nocheck(
1429 urCommandBufferAppendMemBufferReadRectExp, CommandBuffer,
1430 sycl::detail::ur::cast<ur_mem_handle_t>(SrcMem), BufferOffset,
1431 HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch,
1432 HostSlicePitch, DstMem, Deps.size(), Deps.data(), OutSyncPoint);
1433 if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
1436 "Device-to-host buffer copy command not supported by graph backend");
1438 Plugin->checkUrResult(Result);
1443 void MemoryManager::ext_oneapi_copyH2D_cmd_buffer(
1445 ur_exp_command_buffer_handle_t CommandBuffer,
SYCLMemObjI *SYCLMemObj,
1447 sycl::id<3> SrcOffset,
unsigned int SrcElemSize,
void *DstMem,
1450 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1451 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1452 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
1454 const PluginPtr &Plugin = Context->getPlugin();
1461 size_t DstXOffBytes = DstOffset[DstPos.
XTerm] * DstElemSize;
1462 size_t SrcXOffBytes = SrcOffset[SrcPos.
XTerm] * SrcElemSize;
1463 size_t DstAccessRangeWidthBytes = DstAccessRange[DstPos.
XTerm] * DstElemSize;
1464 size_t DstSzWidthBytes = DstSize[DstPos.
XTerm] * DstElemSize;
1465 size_t SrcSzWidthBytes = SrcSize[SrcPos.
XTerm] * SrcElemSize;
1467 if (MemType != detail::SYCLMemObjI::MemObjType::Buffer) {
1469 "Images are not supported in Graphs");
1472 if (1 == DimDst && 1 == DimSrc) {
1473 ur_result_t Result = Plugin->call_nocheck(
1474 urCommandBufferAppendMemBufferWriteExp, CommandBuffer,
1475 sycl::detail::ur::cast<ur_mem_handle_t>(DstMem), DstXOffBytes,
1476 DstAccessRangeWidthBytes, SrcMem + SrcXOffBytes, Deps.size(),
1477 Deps.data(), OutSyncPoint);
1479 if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
1482 "Host-to-device buffer copy command not supported by graph backend");
1484 Plugin->checkUrResult(Result);
1487 size_t BufferRowPitch = (1 == DimDst) ? 0 : DstSzWidthBytes;
1488 size_t BufferSlicePitch =
1489 (3 == DimDst) ? DstSzWidthBytes * DstSize[DstPos.
YTerm] : 0;
1490 size_t HostRowPitch = (1 == DimSrc) ? 0 : SrcSzWidthBytes;
1491 size_t HostSlicePitch =
1492 (3 == DimSrc) ? SrcSzWidthBytes * SrcSize[SrcPos.
YTerm] : 0;
1494 ur_rect_offset_t BufferOffset{DstXOffBytes, DstOffset[DstPos.
YTerm],
1495 DstOffset[DstPos.
ZTerm]};
1496 ur_rect_offset_t HostOffset{SrcXOffBytes, SrcOffset[SrcPos.
YTerm],
1497 SrcOffset[SrcPos.
ZTerm]};
1498 ur_rect_region_t RectRegion{DstAccessRangeWidthBytes,
1499 DstAccessRange[DstPos.
YTerm],
1500 DstAccessRange[DstPos.
ZTerm]};
1502 ur_result_t Result = Plugin->call_nocheck(
1503 urCommandBufferAppendMemBufferWriteRectExp, CommandBuffer,
1504 sycl::detail::ur::cast<ur_mem_handle_t>(DstMem), BufferOffset,
1505 HostOffset, RectRegion, BufferRowPitch, BufferSlicePitch, HostRowPitch,
1506 HostSlicePitch, SrcMem, Deps.size(), Deps.data(), OutSyncPoint);
1508 if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
1511 "Host-to-device buffer copy command not supported by graph backend");
1513 Plugin->checkUrResult(Result);
1518 void MemoryManager::ext_oneapi_copy_usm_cmd_buffer(
1520 ur_exp_command_buffer_handle_t CommandBuffer,
size_t Len,
void *DstMem,
1521 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1522 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1523 if (!SrcMem || !DstMem)
1525 "NULL pointer argument in memory copy operation.");
1527 const PluginPtr &Plugin = Context->getPlugin();
1528 ur_result_t Result = Plugin->call_nocheck(
1529 urCommandBufferAppendUSMMemcpyExp, CommandBuffer, DstMem, SrcMem, Len,
1530 Deps.size(), Deps.data(), OutSyncPoint);
1531 if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
1534 "USM copy command not supported by graph backend");
1536 Plugin->checkUrResult(Result);
1540 void MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
1542 ur_exp_command_buffer_handle_t CommandBuffer,
void *DstMem,
size_t Len,
1543 const std::vector<unsigned char> &Pattern,
1544 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1545 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1549 "NULL pointer argument in memory fill operation.");
1551 const PluginPtr &Plugin = Context->getPlugin();
1552 Plugin->call(urCommandBufferAppendUSMFillExp, CommandBuffer, DstMem,
1553 Pattern.data(), Pattern.size(), Len, Deps.size(), Deps.data(),
1557 void MemoryManager::ext_oneapi_fill_cmd_buffer(
1559 ur_exp_command_buffer_handle_t CommandBuffer,
SYCLMemObjI *SYCLMemObj,
1560 void *Mem,
size_t PatternSize,
const unsigned char *Pattern,
1562 sycl::id<3> AccessOffset,
unsigned int ElementSize,
1563 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1564 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1565 assert(SYCLMemObj &&
"The SYCLMemObj is nullptr");
1567 const PluginPtr &Plugin = Context->getPlugin();
1568 if (SYCLMemObj->
getType() != detail::SYCLMemObjI::MemObjType::Buffer) {
1570 "Images are not supported in Graphs");
1575 bool RangesUsable = (Dim <= 1) || (Size == AccessRange);
1578 bool OffsetUsable = (Dim <= 1) || (AccessOffset ==
sycl::id<3>{0, 0, 0});
1579 size_t RangeMultiplier = AccessRange[0] * AccessRange[1] * AccessRange[2];
1581 if (RangesUsable && OffsetUsable) {
1582 Plugin->call(urCommandBufferAppendMemBufferFillExp, CommandBuffer,
1583 ur::cast<ur_mem_handle_t>(Mem), Pattern, PatternSize,
1584 AccessOffset[0] * ElementSize, RangeMultiplier * ElementSize,
1585 Deps.size(), Deps.data(), OutSyncPoint);
1591 "Not supported configuration of fill requested");
1594 void MemoryManager::ext_oneapi_prefetch_usm_cmd_buffer(
1596 ur_exp_command_buffer_handle_t CommandBuffer,
void *Mem,
size_t Length,
1597 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1598 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1599 const PluginPtr &Plugin = Context->getPlugin();
1600 Plugin->call(urCommandBufferAppendUSMPrefetchExp, CommandBuffer, Mem, Length,
1601 ur_usm_migration_flags_t(0), Deps.size(), Deps.data(),
1605 void MemoryManager::ext_oneapi_advise_usm_cmd_buffer(
1607 ur_exp_command_buffer_handle_t CommandBuffer,
const void *Mem,
1608 size_t Length, ur_usm_advice_flags_t Advice,
1609 std::vector<ur_exp_command_buffer_sync_point_t> Deps,
1610 ur_exp_command_buffer_sync_point_t *OutSyncPoint) {
1611 const PluginPtr &Plugin = Context->getPlugin();
1612 Plugin->call(urCommandBufferAppendUSMAdviseExp, CommandBuffer, Mem, Length,
1613 Advice, Deps.size(), Deps.data(), OutSyncPoint);
1616 void MemoryManager::copy_image_bindless(
1618 const ur_image_desc_t &SrcDesc,
const ur_image_desc_t &DstDesc,
1619 const ur_image_format_t &SrcFormat,
const ur_image_format_t &DstFormat,
1620 const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset,
1621 ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent,
1622 const std::vector<ur_event_handle_t> &DepEvents,
1623 ur_event_handle_t *OutEvent) {
1625 "Copy image bindless must be called with a valid device queue");
1626 assert((Flags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE ||
1627 Flags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST ||
1628 Flags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE) &&
1629 "Invalid flags passed to copy_image_bindless.");
1633 "NULL pointer argument in bindless image copy operation.");
1637 ur_exp_image_copy_region_t CopyRegion{};
1638 CopyRegion.stype = UR_STRUCTURE_TYPE_EXP_IMAGE_COPY_REGION;
1639 CopyRegion.copyExtent = CopyExtent;
1640 CopyRegion.srcOffset = SrcOffset;
1641 CopyRegion.dstOffset = DstOffset;
1643 Plugin->call(urBindlessImagesImageCopyExp, Queue->getHandleRef(), Src, Dst,
1644 &SrcDesc, &DstDesc, &SrcFormat, &DstFormat, &CopyRegion, Flags,
1645 DepEvents.size(), DepEvents.data(), OutEvent);
The context class represents a SYCL context on which kernel functions may be executed.
static void * allocateImageObject(ContextImplPtr TargetContext, void *UserPtr, bool HostPtrReadOnly, const ur_image_desc_t &Desc, const ur_image_format_t &Format, const sycl::property_list &PropsList)
static void * 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)
static void * allocateHostMemory(SYCLMemObjI *MemObj, void *UserPtr, bool HostPtrReadOnly, size_t Size, const sycl::property_list &PropsList)
static void * allocateMemSubBuffer(ContextImplPtr TargetContext, void *ParentMemObj, size_t ElemSize, size_t Offset, range< 3 > Range, std::vector< EventImplPtr > DepEvents, ur_event_handle_t &OutEvent)
static void * allocateBufferObject(ContextImplPtr TargetContext, void *UserPtr, bool HostPtrReadOnly, const size_t Size, const sycl::property_list &PropsList)
static void * 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)
static void * allocate(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, bool InitFromUserData, void *HostPtr, std::vector< EventImplPtr > DepEvents, ur_event_handle_t &OutEvent)
static void * allocateInteropMemObject(ContextImplPtr TargetContext, void *UserPtr, const EventImplPtr &InteropEvent, const ContextImplPtr &InteropContext, const sycl::property_list &PropsList, ur_event_handle_t &OutEventToWait)
static void releaseMemObj(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *MemAllocation, void *UserPtr)
static void release(ContextImplPtr TargetContext, SYCLMemObjI *MemObj, void *MemAllocation, std::vector< EventImplPtr > DepEvents, ur_event_handle_t &OutEvent)
device_image_plain getDeviceImageFromBinaryImage(RTDeviceBinaryImage *BinImage, const context &Ctx, const device &Dev)
RTDeviceBinaryImage & getDeviceImage(const std::string &KernelName, const context &Context, const device &Device, bool JITCompilationIsRequired=false)
device_image_plain build(const device_image_plain &DeviceImage, const std::vector< device > &Devs, const property_list &PropList)
virtual void * allocateHostMem()=0
virtual void releaseMem(ContextImplPtr Context, void *Ptr)=0
virtual MemObjType getType() const =0
virtual void releaseHostMem(void *Ptr)=0
virtual void * allocateMem(ContextImplPtr Context, bool InitFromUserData, void *HostPtr, ur_event_handle_t &InteropEvent)=0
static void bufferAssociateNotification(const void *, const void *)
static void bufferReleaseNotification(const void *, const void *)
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
A unique identifier of an item in an index space.
Objects of the property_list class are containers for the SYCL properties.
bool has_property() const noexcept
PropT get_property() const
uint64_t emitMemAllocBeginTrace(uintptr_t ObjHandle, size_t AllocSize, size_t GuardZone)
decltype(Obj::impl) const & getSyclObjImpl(const Obj &SyclObject)
void memReleaseHelper(const PluginPtr &Plugin, ur_mem_handle_t Mem)
static void memcpyFromDeviceGlobalUSM(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, void *Dest, const std::vector< ur_event_handle_t > &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl)
void prepTermPositions(TermPositions &pos, int Dimensions, detail::SYCLMemObjI::MemObjType Type)
static void waitForEvents(const std::vector< EventImplPtr > &Events)
static ur_program_handle_t getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry)
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< ur_event_handle_t > DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl)
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
void memUnmapHelper(const PluginPtr &Plugin, ur_queue_handle_t command_queue, ur_mem_handle_t memobj, void *mapped_ptr, uint32_t num_events_in_wait_list, const ur_event_handle_t *event_wait_list, ur_event_handle_t *event)
static void memcpyToDeviceGlobalUSM(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, const void *Src, const std::vector< ur_event_handle_t > &DepEvents, ur_event_handle_t *OutEvent, const detail::EventImplPtr &OutEventImpl)
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< ur_event_handle_t > DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl)
std::shared_ptr< event_impl > EventImplPtr
void memBufferCreateHelper(const PluginPtr &Plugin, ur_context_handle_t Ctx, ur_mem_flags_t Flags, size_t Size, ur_mem_handle_t *RetMem, const ur_buffer_properties_t *Props=nullptr)
static void memcpyFromDeviceGlobalDirect(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, void *Dest, const std::vector< ur_event_handle_t > &DepEvents, ur_event_handle_t *OutEvent)
std::shared_ptr< plugin > PluginPtr
uint64_t emitMemReleaseBeginTrace(uintptr_t ObjHandle, uintptr_t AllocPtr)
void emitMemAllocEndTrace(uintptr_t ObjHandle, uintptr_t AllocPtr, size_t AllocSize, size_t GuardZone, uint64_t CorrelationID)
void emitMemReleaseEndTrace(uintptr_t ObjHandle, uintptr_t AllocPtr, uint64_t CorrelationID)
static void memcpyToDeviceGlobalDirect(QueueImplPtr Queue, DeviceGlobalMapEntry *DeviceGlobalEntry, size_t NumBytes, size_t Offset, const void *Src, const std::vector< ur_event_handle_t > &DepEvents, ur_event_handle_t *OutEvent)
std::shared_ptr< sycl::detail::queue_impl > QueueImplPtr
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 >, ur_event_handle_t &, const detail::EventImplPtr &)
static ur_mem_flags_t getMemObjCreationFlags(void *UserPtr, bool HostPtrReadOnly)
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
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< ur_event_handle_t > DepEvents, ur_event_handle_t &OutEvent, const detail::EventImplPtr &OutEventImpl)
exception set_ur_error(exception &&e, int32_t ur_err)
void memBufferMapHelper(const PluginPtr &Plugin, ur_queue_handle_t command_queue, ur_mem_handle_t buffer, bool blocking_map, ur_map_flags_t map_flags, size_t offset, size_t size, uint32_t num_events_in_wait_list, const ur_event_handle_t *event_wait_list, ur_event_handle_t *event, void **ret_map)
void copy(handler &CGH, const T *Src, T *Dest, size_t Count)
void unmap(const void *Ptr, size_t NumBytes, const context &SyclContext)
void fill(sycl::handler &CGH, T *Ptr, const T &Pattern, size_t Count)
class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
usm::alloc get_pointer_type(const void *ptr, const context &ctxt)
Query the allocation type from a USM pointer.
constexpr mode_tag_t< access_mode::read_write > read_write
class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
bool MIsDeviceImageScopeDecorated
std::set< std::uintptr_t > MImageIdentifiers
DeviceGlobalUSMMem & getOrAllocateDeviceGlobalUSM(const std::shared_ptr< queue_impl > &QueueImpl)
std::uint32_t MDeviceGlobalTSize
std::unordered_set< RTDeviceBinaryImage * > MImages
OwnedUrEvent getInitEvent(const PluginPtr &Plugin)
void *const & getPtr() const noexcept
ur_event_handle_t GetEvent()
C++ utilities for Unified Runtime integration.