32 handler::handler(std::shared_ptr<detail::queue_impl> Queue,
bool IsHost)
33 : handler(Queue, Queue, nullptr, IsHost) {}
35 handler::handler(std::shared_ptr<detail::queue_impl> Queue,
36 std::shared_ptr<detail::queue_impl> PrimaryQueue,
37 std::shared_ptr<detail::queue_impl> SecondaryQueue,
39 : MImpl(
std::make_shared<detail::handler_impl>(
std::move(PrimaryQueue),
40 std::move(SecondaryQueue))),
41 MQueue(
std::move(Queue)), MIsHost(IsHost) {}
46 void handler::setStateExplicitKernelBundle() {
47 MImpl->setStateExplicitKernelBundle();
53 void handler::setStateSpecConstSet() { MImpl->setStateSpecConstSet(); }
57 bool handler::isStateExplicitKernelBundle()
const {
58 return MImpl->isStateExplicitKernelBundle();
65 std::shared_ptr<detail::kernel_bundle_impl>
66 handler::getOrInsertHandlerKernelBundle(
bool Insert)
const {
67 if (!MImpl->MKernelBundle && Insert) {
68 MImpl->MKernelBundle =
70 MQueue->get_context(), {MQueue->get_device()}, {}));
72 return MImpl->MKernelBundle;
76 void handler::setHandlerKernelBundle(
77 const std::shared_ptr<detail::kernel_bundle_impl> &NewKernelBundleImpPtr) {
78 MImpl->MKernelBundle = NewKernelBundleImpPtr;
81 void handler::setHandlerKernelBundle(kernel Kernel) {
85 std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpl =
87 setHandlerKernelBundle(KernelBundleImpl);
90 event handler::finalize() {
97 const auto &type = getType();
98 if (type == detail::CG::Kernel) {
100 std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr =
101 getOrInsertHandlerKernelBundle(
false);
102 if (KernelBundleImpPtr) {
104 if (!KernelBundleImpPtr->isInterop() &&
105 !MImpl->isStateExplicitKernelBundle()) {
107 detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName);
108 bool KernelInserted =
109 KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
112 if (!KernelInserted &&
113 KernelBundleImpPtr->get_bundle_state() == bundle_state::input) {
115 detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
117 kernel_bundle<bundle_state::executable> ExecKernelBundle =
120 setHandlerKernelBundle(KernelBundleImpPtr);
122 KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
127 "Failed to add kernel to kernel bundle.");
130 switch (KernelBundleImpPtr->get_bundle_state()) {
131 case bundle_state::input: {
133 kernel_bundle<bundle_state::executable> ExecBundle =
build(
135 KernelBundleImpPtr));
137 setHandlerKernelBundle(KernelBundleImpPtr);
140 case bundle_state::executable:
143 case bundle_state::object:
144 assert(0 &&
"Expected that the bundle is either in input or executable "
150 if (!MQueue->is_in_fusion_mode() &&
151 MRequirements.size() + MEvents.size() + MStreamStorage.size() == 0) {
157 std::vector<RT::PiEvent> RawEvents;
161 auto EnqueueKernel = [&]() {
163 pi_int32 Result = PI_ERROR_INVALID_VALUE;
165 if (MQueue->is_host()) {
166 MHostKernel->call(MNDRDesc, (NewEvent)
167 ? NewEvent->getHostProfilingInfo()
171 if (MQueue->getPlugin().getBackend() ==
172 backend::ext_intel_esimd_emulator) {
174 nullptr,
reinterpret_cast<pi_kernel>(MHostKernel->getPtr()),
175 MNDRDesc.Dims, &MNDRDesc.GlobalOffset[0],
176 &MNDRDesc.GlobalSize[0], &MNDRDesc.LocalSize[0], 0,
nullptr,
181 MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, MKernel,
182 MKernelName, MOSModuleHandle, RawEvents, OutEvent,
nullptr);
188 bool DiscardEvent =
false;
189 if (MQueue->has_discard_events_support()) {
191 bool KernelUsesAssert =
192 !(MKernel && MKernel->isInterop()) &&
193 detail::ProgramManager::getInstance().kernelUsesAssert(
194 MOSModuleHandle, MKernelName);
195 DiscardEvent = !KernelUsesAssert;
199 if (PI_SUCCESS != EnqueueKernel())
200 throw runtime_error(
"Enqueue process failed.",
201 PI_ERROR_INVALID_OPERATION);
203 NewEvent = std::make_shared<detail::event_impl>(MQueue);
204 NewEvent->setContextImpl(MQueue->getContextImplPtr());
205 NewEvent->setStateIncomplete();
206 OutEvent = &NewEvent->getHandleRef();
208 NewEvent->setSubmissionTime();
210 if (PI_SUCCESS != EnqueueKernel())
211 throw runtime_error(
"Enqueue process failed.",
212 PI_ERROR_INVALID_OPERATION);
213 else if (NewEvent->is_host() || NewEvent->getHandleRef() ==
nullptr)
214 NewEvent->setComplete();
216 MLastEvent = detail::createSyclObjFromImpl<event>(NewEvent);
222 std::unique_ptr<detail::CG> CommandGroup;
224 case detail::CG::Kernel:
225 case detail::CG::RunOnHostIntel: {
229 CommandGroup.reset(
new detail::CGExecKernel(
230 std::move(MNDRDesc), std::move(MHostKernel), std::move(MKernel),
231 std::move(MImpl->MKernelBundle), std::move(MArgsStorage),
232 std::move(MAccStorage), std::move(MSharedPtrStorage),
233 std::move(MRequirements), std::move(MEvents), std::move(MArgs),
234 MKernelName, MOSModuleHandle, std::move(MStreamStorage),
235 std::move(MImpl->MAuxiliaryResources), MCGType, MCodeLoc));
238 case detail::CG::CodeplayInteropTask:
239 CommandGroup.reset(
new detail::CGInteropTask(
240 std::move(MInteropTask), std::move(MArgsStorage),
241 std::move(MAccStorage), std::move(MSharedPtrStorage),
242 std::move(MRequirements), std::move(MEvents), MCGType, MCodeLoc));
244 case detail::CG::CopyAccToPtr:
245 case detail::CG::CopyPtrToAcc:
246 case detail::CG::CopyAccToAcc:
247 CommandGroup.reset(
new detail::CGCopy(
248 MCGType, MSrcPtr, MDstPtr, std::move(MArgsStorage),
249 std::move(MAccStorage), std::move(MSharedPtrStorage),
250 std::move(MRequirements), std::move(MEvents), MCodeLoc));
252 case detail::CG::Fill:
253 CommandGroup.reset(
new detail::CGFill(
254 std::move(MPattern), MDstPtr, std::move(MArgsStorage),
255 std::move(MAccStorage), std::move(MSharedPtrStorage),
256 std::move(MRequirements), std::move(MEvents), MCodeLoc));
258 case detail::CG::UpdateHost:
259 CommandGroup.reset(
new detail::CGUpdateHost(
260 MDstPtr, std::move(MArgsStorage), std::move(MAccStorage),
261 std::move(MSharedPtrStorage), std::move(MRequirements),
262 std::move(MEvents), MCodeLoc));
264 case detail::CG::CopyUSM:
265 CommandGroup.reset(
new detail::CGCopyUSM(
266 MSrcPtr, MDstPtr, MLength, std::move(MArgsStorage),
267 std::move(MAccStorage), std::move(MSharedPtrStorage),
268 std::move(MRequirements), std::move(MEvents), MCodeLoc));
270 case detail::CG::FillUSM:
271 CommandGroup.reset(
new detail::CGFillUSM(
272 std::move(MPattern), MDstPtr, MLength, std::move(MArgsStorage),
273 std::move(MAccStorage), std::move(MSharedPtrStorage),
274 std::move(MRequirements), std::move(MEvents), MCodeLoc));
276 case detail::CG::PrefetchUSM:
277 CommandGroup.reset(
new detail::CGPrefetchUSM(
278 MDstPtr, MLength, std::move(MArgsStorage), std::move(MAccStorage),
279 std::move(MSharedPtrStorage), std::move(MRequirements),
280 std::move(MEvents), MCodeLoc));
282 case detail::CG::AdviseUSM:
283 CommandGroup.reset(
new detail::CGAdviseUSM(
284 MDstPtr, MLength, MImpl->MAdvice, std::move(MArgsStorage),
285 std::move(MAccStorage), std::move(MSharedPtrStorage),
286 std::move(MRequirements), std::move(MEvents), MCGType, MCodeLoc));
288 case detail::CG::Copy2DUSM:
289 CommandGroup.reset(
new detail::CGCopy2DUSM(
290 MSrcPtr, MDstPtr, MImpl->MSrcPitch, MImpl->MDstPitch, MImpl->MWidth,
291 MImpl->MHeight, std::move(MArgsStorage), std::move(MAccStorage),
292 std::move(MSharedPtrStorage), std::move(MRequirements),
293 std::move(MEvents), MCodeLoc));
295 case detail::CG::Fill2DUSM:
296 CommandGroup.reset(
new detail::CGFill2DUSM(
297 std::move(MPattern), MDstPtr, MImpl->MDstPitch, MImpl->MWidth,
298 MImpl->MHeight, std::move(MArgsStorage), std::move(MAccStorage),
299 std::move(MSharedPtrStorage), std::move(MRequirements),
300 std::move(MEvents), MCodeLoc));
302 case detail::CG::Memset2DUSM:
303 CommandGroup.reset(
new detail::CGMemset2DUSM(
304 MPattern[0], MDstPtr, MImpl->MDstPitch, MImpl->MWidth, MImpl->MHeight,
305 std::move(MArgsStorage), std::move(MAccStorage),
306 std::move(MSharedPtrStorage), std::move(MRequirements),
307 std::move(MEvents), MCodeLoc));
309 case detail::CG::CodeplayHostTask:
310 CommandGroup.reset(
new detail::CGHostTask(
311 std::move(MHostTask), MQueue, MQueue->getContextImplPtr(),
312 std::move(MArgs), std::move(MArgsStorage), std::move(MAccStorage),
313 std::move(MSharedPtrStorage), std::move(MRequirements),
314 std::move(MEvents), MCGType, MCodeLoc));
316 case detail::CG::Barrier:
317 case detail::CG::BarrierWaitlist:
318 CommandGroup.reset(
new detail::CGBarrier(
319 std::move(MEventsWaitWithBarrier), std::move(MArgsStorage),
320 std::move(MAccStorage), std::move(MSharedPtrStorage),
321 std::move(MRequirements), std::move(MEvents), MCGType, MCodeLoc));
323 case detail::CG::None:
325 std::cout <<
"WARNING: An empty command group is submitted." << std::endl;
328 MLastEvent = detail::createSyclObjFromImpl<event>(Event);
333 throw sycl::runtime_error(
334 "Internal Error. Command group cannot be constructed.",
335 PI_ERROR_INVALID_OPERATION);
338 std::move(CommandGroup), std::move(MQueue));
340 MLastEvent = detail::createSyclObjFromImpl<event>(Event);
344 void handler::addReduction(
const std::shared_ptr<const void> &ReduObj) {
345 MImpl->MAuxiliaryResources.push_back(ReduObj);
353 MRequirements.push_back(Req);
355 MAccStorage.push_back(std::move(AccImpl));
358 MAssociatedAccesors.emplace_back(detail::kernel_param_kind_t::kind_accessor,
359 Req,
static_cast<int>(AccTarget),
364 size_t &IndexShift,
int Size,
365 bool IsKernelCreatedFromSource,
367 std::vector<detail::ArgDesc> &Args,
371 AccImpl->
resize(GlobalSize);
373 Args.emplace_back(kernel_param_kind_t::kind_accessor, AccImpl, Size,
379 if (!isESIMD && !IsKernelCreatedFromSource) {
382 const size_t SizeAccField =
383 sizeof(size_t) * (AccImpl->
MDims == 0 ? 1 : AccImpl->
MDims);
385 Args.emplace_back(kernel_param_kind_t::kind_std_layout,
389 Args.emplace_back(kernel_param_kind_t::kind_std_layout,
393 Args.emplace_back(kernel_param_kind_t::kind_std_layout,
394 &AccImpl->
MOffset[0], SizeAccField, Index + IndexShift);
399 const int Size,
const size_t Index,
size_t &IndexShift,
400 bool IsKernelCreatedFromSource,
bool IsESIMD) {
404 case kernel_param_kind_t::kind_std_layout:
405 case kernel_param_kind_t::kind_pointer: {
406 MArgs.emplace_back(Kind, Ptr, Size, Index + IndexShift);
409 case kernel_param_kind_t::kind_stream: {
411 stream *S =
static_cast<stream *
>(Ptr);
413 detail::AccessorBaseHost *GBufBase =
414 static_cast<detail::AccessorBaseHost *
>(&S->GlobalBuf);
418 IsKernelCreatedFromSource,
419 MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
421 detail::AccessorBaseHost *GOffsetBase =
422 static_cast<detail::AccessorBaseHost *
>(&S->GlobalOffset);
426 IsKernelCreatedFromSource,
427 MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
429 detail::AccessorBaseHost *GFlushBase =
430 static_cast<detail::AccessorBaseHost *
>(&S->GlobalFlushBuf);
434 size_t GlobalSize = MNDRDesc.GlobalSize.size();
440 if (GlobalSize == 0) {
442 GlobalSize = MNDRDesc.NumWorkGroups.size();
445 IsKernelCreatedFromSource, GlobalSize, MArgs,
448 MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
449 &S->FlushBufferSize,
sizeof(S->FlushBufferSize),
454 case kernel_param_kind_t::kind_accessor: {
459 case access::target::device:
460 case access::target::constant_buffer: {
463 IsKernelCreatedFromSource,
464 MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
467 case access::target::local: {
468 detail::LocalAccessorImplHost *LAcc =
469 static_cast<detail::LocalAccessorImplHost *
>(Ptr);
471 range<3> &Size = LAcc->MSize;
472 const int Dims = LAcc->MDims;
473 int SizeInBytes = LAcc->MElemSize;
474 for (
int I = 0; I < Dims; ++I)
475 SizeInBytes *= Size[I];
478 SizeInBytes = std::max(SizeInBytes, 1);
479 MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
nullptr,
480 SizeInBytes, Index + IndexShift);
481 if (!IsKernelCreatedFromSource) {
483 const size_t SizeAccField = Dims *
sizeof(Size[0]);
484 MArgs.emplace_back(kernel_param_kind_t::kind_std_layout, &Size,
485 SizeAccField, Index + IndexShift);
487 MArgs.emplace_back(kernel_param_kind_t::kind_std_layout, &Size,
488 SizeAccField, Index + IndexShift);
490 MArgs.emplace_back(kernel_param_kind_t::kind_std_layout, &Size,
491 SizeAccField, Index + IndexShift);
495 case access::target::image:
496 case access::target::image_array: {
498 MArgs.emplace_back(Kind, AccImpl, Size, Index + IndexShift);
499 if (!IsKernelCreatedFromSource) {
505 case access::target::host_image:
506 case access::target::host_buffer: {
507 throw sycl::invalid_parameter_error(
"Unsupported accessor target case.",
508 PI_ERROR_INVALID_OPERATION);
514 case kernel_param_kind_t::kind_sampler: {
515 MArgs.emplace_back(kernel_param_kind_t::kind_sampler, Ptr,
sizeof(sampler),
519 case kernel_param_kind_t::kind_specialization_constants_buffer: {
521 kernel_param_kind_t::kind_specialization_constants_buffer, Ptr, Size,
525 case kernel_param_kind_t::kind_invalid:
526 throw runtime_error(
"Invalid kernel param kind", PI_ERROR_INVALID_VALUE);
541 void handler::extractArgsAndReqs() {
542 assert(MKernel &&
"MKernel is not initialized");
543 std::vector<detail::ArgDesc> UnPreparedArgs = std::move(MArgs);
547 UnPreparedArgs.begin(), UnPreparedArgs.end(),
549 return (first.MIndex < second.MIndex);
552 const bool IsKernelCreatedFromSource = MKernel->isCreatedFromSource();
555 size_t IndexShift = 0;
556 for (
size_t I = 0; I < UnPreparedArgs.size(); ++I) {
557 void *Ptr = UnPreparedArgs[I].MPtr;
559 const int &Size = UnPreparedArgs[I].MSize;
560 const int Index = UnPreparedArgs[I].MIndex;
561 processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource,
566 void handler::extractArgsAndReqsFromLambda(
567 char *LambdaPtr,
size_t KernelArgsNum,
568 const detail::kernel_param_desc_t *KernelArgs,
bool IsESIMD) {
569 const bool IsKernelCreatedFromSource =
false;
570 size_t IndexShift = 0;
573 for (
size_t I = 0; I < KernelArgsNum; ++I) {
574 void *Ptr = LambdaPtr + KernelArgs[I].offset;
576 const int &Size = KernelArgs[I].info;
577 if (Kind == detail::kernel_param_kind_t::kind_accessor) {
582 if ((AccTarget == access::target::device ||
583 AccTarget == access::target::constant_buffer) ||
584 (AccTarget == access::target::image ||
585 AccTarget == access::target::image_array)) {
586 detail::AccessorBaseHost *AccBase =
587 static_cast<detail::AccessorBaseHost *
>(Ptr);
589 }
else if (AccTarget == access::target::local) {
590 detail::LocalAccessorBaseHost *LocalAccBase =
591 static_cast<detail::LocalAccessorBaseHost *
>(Ptr);
595 processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource,
603 std::string handler::getKernelName() {
604 return MKernel->get_info<info::kernel::function_name>();
607 void handler::verifyUsedKernelBundle(
const std::string &KernelName) {
608 auto UsedKernelBundleImplPtr =
609 getOrInsertHandlerKernelBundle(
false);
610 if (!UsedKernelBundleImplPtr)
614 if (!MImpl->isStateExplicitKernelBundle())
619 if (!UsedKernelBundleImplPtr->has_kernel(KernelID, Dev))
620 throw sycl::exception(
622 "The kernel bundle in use does not contain the kernel");
625 void handler::ext_oneapi_barrier(
const std::vector<event> &WaitList) {
626 throwIfActionIsCreated();
627 MCGType = detail::CG::BarrierWaitlist;
628 MEventsWaitWithBarrier.resize(WaitList.size());
630 WaitList.begin(), WaitList.end(), MEventsWaitWithBarrier.begin(),
631 [](
const event &Event) { return detail::getSyclObjImpl(Event); });
636 handler::ext_oneapi_barrier(WaitList);
639 using namespace sycl::detail;
640 bool handler::DisableRangeRounding() {
644 bool handler::RangeRoundingTrace() {
648 void handler::GetRangeRoundingSettings(
size_t &MinFactor,
size_t &GoodFactor,
650 SYCLConfig<SYCL_PARALLEL_FOR_RANGE_ROUNDING_PARAMS>::GetSettings(
651 MinFactor, GoodFactor, MinRange);
655 throwIfActionIsCreated();
656 MSrcPtr =
const_cast<void *
>(Src);
659 setType(detail::CG::CopyUSM);
662 void handler::memset(
void *Dest,
int Value,
size_t Count) {
663 throwIfActionIsCreated();
665 MPattern.push_back(
static_cast<char>(Value));
667 setType(detail::CG::FillUSM);
671 throwIfActionIsCreated();
672 MDstPtr =
const_cast<void *
>(Ptr);
674 setType(detail::CG::PrefetchUSM);
677 void handler::mem_advise(
const void *Ptr,
size_t Count,
int Advice) {
678 throwIfActionIsCreated();
679 MDstPtr =
const_cast<void *
>(Ptr);
682 setType(detail::CG::AdviseUSM);
685 void handler::ext_oneapi_memcpy2d_impl(
void *Dest,
size_t DestPitch,
686 const void *Src,
size_t SrcPitch,
687 size_t Width,
size_t Height) {
689 MSrcPtr =
const_cast<void *
>(Src);
691 MImpl->MSrcPitch = SrcPitch;
692 MImpl->MDstPitch = DestPitch;
693 MImpl->MWidth = Width;
694 MImpl->MHeight = Height;
695 setType(detail::CG::Copy2DUSM);
698 void handler::ext_oneapi_fill2d_impl(
void *Dest,
size_t DestPitch,
699 const void *Value,
size_t ValueSize,
700 size_t Width,
size_t Height) {
703 MPattern.resize(ValueSize);
705 MImpl->MDstPitch = DestPitch;
706 MImpl->MWidth = Width;
707 MImpl->MHeight = Height;
708 setType(detail::CG::Fill2DUSM);
711 void handler::ext_oneapi_memset2d_impl(
void *Dest,
size_t DestPitch,
int Value,
712 size_t Width,
size_t Height) {
715 MPattern.push_back(
static_cast<char>(Value));
716 MImpl->MDstPitch = DestPitch;
717 MImpl->MWidth = Width;
718 MImpl->MHeight = Height;
719 setType(detail::CG::Memset2DUSM);
722 void handler::use_kernel_bundle(
725 std::shared_ptr<detail::queue_impl> PrimaryQueue =
726 MImpl->MSubmissionPrimaryQueue;
727 if (PrimaryQueue->get_context() != ExecBundle.
get_context())
728 throw sycl::exception(
730 "Context associated with the primary queue is different from the "
731 "context associated with the kernel bundle");
733 std::shared_ptr<detail::queue_impl> SecondaryQueue =
734 MImpl->MSubmissionSecondaryQueue;
735 if (SecondaryQueue &&
736 SecondaryQueue->get_context() != ExecBundle.
get_context())
737 throw sycl::exception(
739 "Context associated with the secondary queue is different from the "
740 "context associated with the kernel bundle");
742 setStateExplicitKernelBundle();
746 void handler::depends_on(
event Event) {
748 if (EventImpl->isDiscarded()) {
750 "Queue operation cannot depend on discarded event.");
752 MEvents.push_back(EventImpl);
755 void handler::depends_on(
const std::vector<event> &Events) {
756 for (
const event &Event : Events) {
758 if (EventImpl->isDiscarded()) {
759 throw sycl::exception(
761 "Queue operation cannot depend on discarded event.");
763 MEvents.push_back(EventImpl);
770 auto &Plugin = ContextImpl->getPlugin();
774 &SupportsOp,
nullptr);
778 bool handler::supportsUSMMemcpy2D() {
779 for (
const std::shared_ptr<detail::queue_impl> &QueueImpl :
780 {MImpl->MSubmissionPrimaryQueue, MImpl->MSubmissionSecondaryQueue}) {
789 bool handler::supportsUSMFill2D() {
790 for (
const std::shared_ptr<detail::queue_impl> &QueueImpl :
791 {MImpl->MSubmissionPrimaryQueue, MImpl->MSubmissionSecondaryQueue}) {
800 bool handler::supportsUSMMemset2D() {
801 for (
const std::shared_ptr<detail::queue_impl> &QueueImpl :
802 {MImpl->MSubmissionPrimaryQueue, MImpl->MSubmissionSecondaryQueue}) {
811 id<2> handler::computeFallbackKernelBounds(
size_t Width,
size_t Height) {
812 device Dev = MQueue->get_device();
813 id<2> ItemLimit = Dev.get_info<info::device::max_work_item_sizes<2>>() *
814 Dev.get_info<info::device::max_compute_units>();
815 return id<2>{std::min(ItemLimit[0], Height), std::min(ItemLimit[1], Width)};
range< 3 > & MAccessRange
void resize(size_t GlobalSize)
range< 3 > & MMemoryRange
An event object can be used to synchronize memory transfers, enqueues of kernels and signaling barrie...
Command group handler class.
context get_context() const noexcept
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL2020_DEPRECATED(message)
__ESIMD_API void barrier()
Generic work-group barrier.
__SYCL_EXTERN_STREAM_ATTRS ostream cout
Linked to standard output.
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
bool trace(TraceLevel level)
::pi_context_info PiContextInfo
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
kernel_id get_kernel_id_impl(std::string KernelName)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
std::shared_ptr< event_impl > EventImplPtr
pi_int32 enqueueImpKernel(const QueueImplPtr &Queue, NDRDescT &NDRDesc, std::vector< ArgDesc > &Args, const std::shared_ptr< detail::kernel_bundle_impl > &KernelBundleImplPtr, const std::shared_ptr< detail::kernel_impl > &MSyclKernel, const std::string &KernelName, const detail::OSModuleHandle &OSModuleHandle, std::vector< RT::PiEvent > &RawEvents, RT::PiEvent *OutEvent, const std::function< void *(Requirement *Req)> &getMemAllocationFunc)
AccessorImplHost Requirement
void memcpy(void *Dst, const void *Src, std::size_t Size)
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
std::shared_ptr< AccessorImplHost > AccessorImplPtr
prefetch_impl< _B > prefetch
constexpr size_t MaxNumAdditionalArgs
kernel_bundle< bundle_state::executable > build(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
static bool checkContextSupports(const std::shared_ptr< detail::context_impl > &ContextImpl, detail::RT::PiContextInfo InfoQuery)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, size_t &IndexShift, int Size, bool IsKernelCreatedFromSource, size_t GlobalSize, std::vector< detail::ArgDesc > &Args, bool isESIMD)
---— Error handling, matching OpenCL plugin semantics.
pi_result piEnqueueKernelLaunch(pi_queue queue, pi_kernel kernel, pi_uint32 work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, pi_uint32 num_events_in_wait_list, const pi_event *event_wait_list, pi_event *event)
@ PI_EXT_ONEAPI_CONTEXT_INFO_USM_MEMSET2D_SUPPORT
@ PI_EXT_ONEAPI_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT
@ PI_EXT_ONEAPI_CONTEXT_INFO_USM_FILL2D_SUPPORT
pi_result piContextGetInfo(pi_context context, pi_context_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
C++ wrapper of extern "C" PI interfaces.
Implementation of a PI Kernel for CUDA.