diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index ba70ddefcee2e..75bd4d80b28b6 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -170,6 +170,7 @@ class CG { CodeplayInteropTask = 13, CodeplayHostTask = 14, AdviseUSM = 15, + ReadWriteHostPipe = 16, }; CG(CGTYPE Type, std::vector> ArgsStorage, @@ -387,6 +388,35 @@ class CGCopyUSM : public CG { size_t getLength() { return MLength; } }; +/// "ReadWriteHostPipe" command group class. +class CGReadWriteHostPipe : public CG { + std::string PipeName; + bool Blocking; + void *HostPtr; + size_t TypeSize; + bool IsReadOp; + +public: + CGReadWriteHostPipe(std::string Name, bool Block, void *Ptr, size_t Size, + bool Read, std::vector> ArgsStorage, + std::vector AccStorage, + std::vector> SharedPtrStorage, + std::vector Requirements, + std::vector Events, + detail::code_location loc = {}) + : CG(ReadWriteHostPipe, std::move(ArgsStorage), std::move(AccStorage), + std::move(SharedPtrStorage), std::move(Requirements), + std::move(Events), std::move(loc)), + PipeName(Name), Blocking(Block), HostPtr(Ptr), TypeSize(Size), + IsReadOp(Read) {} + + std::string getPipeName() { return PipeName; } + void *getHostPtr() { return HostPtr; } + size_t getTypeSize() { return TypeSize; } + bool isBlocking() { return Blocking; } + bool isReadHostPipe() { return IsReadOp; } +}; + /// "Fill USM" command group class. class CGFillUSM : public CG { std::vector MPattern; diff --git a/sycl/include/CL/sycl/detail/pi.def b/sycl/include/CL/sycl/detail/pi.def index c9a68c6cadec3..57bd906541698 100644 --- a/sycl/include/CL/sycl/detail/pi.def +++ b/sycl/include/CL/sycl/detail/pi.def @@ -120,6 +120,9 @@ _PI_API(piEnqueueMemImageCopy) _PI_API(piEnqueueMemImageFill) _PI_API(piEnqueueMemBufferMap) _PI_API(piEnqueueMemUnmap) +// Host pipes +_PI_API(piextEnqueueReadHostPipe) +_PI_API(piextEnqueueWriteHostPipe) // USM _PI_API(piextUSMHostAlloc) _PI_API(piextUSMDeviceAlloc) @@ -138,4 +141,5 @@ _PI_API(piextPluginGetOpaqueData) _PI_API(piTearDown) + #undef _PI_API diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 37fa88a0ddf69..076e20d7a81ff 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -1760,6 +1760,60 @@ __SYCL_EXPORT pi_result piextUSMGetMemAllocInfo( pi_context context, const void *ptr, pi_mem_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret); +/// +// Host Pipes +/// + +/// Read from pipe of a given name +/// +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event); + +/// Write to pipe of a given name +/// +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event); + /// API to get Plugin internal data, opaque to SYCL RT. Some devices whose /// device code is compiled by the host compiler (e.g. CPU emulators) may use it /// to access some device code functionality implemented in/behind the plugin. diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 3004f438e32b9..f7964601da01c 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -2615,6 +2615,15 @@ class __SYCL_EXPORT handler { /// \param Advice is a device-defined advice for the specified allocation. void mem_advise(const void *Ptr, size_t Length, int Advice); + /// Read from or write to host pipes given a host address and + /// \param Name name of the host pipe to be passed into lower level runtime + /// \param Ptr host pointer of host pipe + /// \param Size the size of data getting read back / to. + /// \param Blocking if read/write opeartion is blocking + /// \param Read 1 for read, 0 for write + void read_write_host_pipe(std::string Name, void *Ptr, size_t Size, + bool Block, bool Read); + private: std::shared_ptr MQueue; /// The storage for the arguments passed. @@ -2663,6 +2672,16 @@ class __SYCL_EXPORT handler { /// The list of valid SYCL events that need to complete /// before barrier command can be executed std::vector MEventsWaitWithBarrier; + /// Host pipe name + std::string HostPipeName; + /// Host pipe host pointer + void *HostPipePtr = nullptr; + /// Host pipe read write operation is blocking + bool HostPipeBlocking = false; + /// Host pipe pointer type size + size_t HostPipeTypeSize = 0; + /// if the operation is read or write + bool HostPipeRead = true; bool MIsHost = false; diff --git a/sycl/include/sycl/ext/intel/experimental/host_pipes.hpp b/sycl/include/sycl/ext/intel/experimental/host_pipes.hpp new file mode 100644 index 0000000000000..b7fd1fdf57e9d --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/host_pipes.hpp @@ -0,0 +1,115 @@ +//==---------------- pipes.hpp - SYCL pipes ------------*- C++ -*-----------==// +// +// 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 +// +// ===--------------------------------------------------------------------=== // + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#ifdef XPTI_ENABLE_INSTRUMENTATION +#include +#include +#endif + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace oneapi { +namespace experimental { + +// min_capacity property has one integer non-type parameter. +struct min_capacity_key { + template + using value_t = + property_value>; +}; +// min_capacity is an object of a property value type of min_capacity. +template +inline constexpr min_capacity_key::value_t min_capacity; + +template <> struct is_property_key : std::true_type {}; + +namespace detail { + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::MinCapacity; +}; + +template <> struct IsCompileTimeProperty : std::true_type {}; + +} // namespace detail + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +using default_pipe_properties = + decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::oneapi::experimental::min_capacity<0>)); + +template +class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_attributes_global_variable("sycl-host-access", + "readwrite")]] +#endif + host_pipe { // TODO change name to pipe, and merge into the existing pipe + // implementation + static_assert( + sycl::ext::oneapi::experimental::is_property_list_v, + "Host pipe is available only through new property list"); + +public: + using value_type = _dataT; + static constexpr int32_t min_cap = + PropertyList::template has_property< + sycl::ext::oneapi::experimental::min_capacity_key>() + ? PropertyList::template get_property< + sycl::ext::oneapi::experimental::min_capacity_key>() + .value + : 0; + + // Blocking pipes + static _dataT read(queue & q, memory_order order = memory_order::seq_cst); + static void write(queue & q, const _dataT &data, + memory_order order = memory_order::seq_cst); + // Non-blocking pipes + static _dataT read(queue & q, bool &success_code, + memory_order order = memory_order::seq_cst); + static void write(queue & q, const _dataT &data, bool &success_code, + memory_order order = memory_order::seq_cst); + +private: + static constexpr int32_t m_Size = sizeof(_dataT); + static constexpr int32_t m_Alignment = alignof(_dataT); + static constexpr int32_t ID = _name::id; +#ifdef __SYCL_DEVICE_ONLY__ + static constexpr struct ConstantPipeStorage m_Storage + __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity}; +#endif // __SYCL_DEVICE_ONLY__ +}; + +} // namespace experimental +} // namespace intel +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 658fa6721ca20..b9c33628ca018 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -148,7 +148,8 @@ enum PropKind : uint32_t { HostAccess = 1, InitMode = 2, ImplementInCSR = 3, - PropKindSize = 4, + MinCapacity = 4, + PropKindSize = 5, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 03f0061f8a8f0..37f41351ec9d9 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1763,6 +1763,25 @@ pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, DIE_NO_IMPLEMENTATION; } +/// Host Pips +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + DIE_NO_IMPLEMENTATION; +} + +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + DIE_NO_IMPLEMENTATION; +} + pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_info, size_t, void *, size_t *) { DIE_NO_IMPLEMENTATION; diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 9071c6bdca92c..1ec1035c15ad6 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -7577,6 +7577,62 @@ pi_result piextUSMGetMemAllocInfo(pi_context Context, const void *Ptr, return PI_SUCCESS; } +/// Host Pips + +/// API to read host pipe +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + + return PI_SUCCESS; +} + +/// API to write host pipe +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + + return PI_SUCCESS; +} + pi_result piKernelSetExecInfo(pi_kernel Kernel, pi_kernel_exec_info ParamName, size_t ParamValueSize, const void *ParamValue) { (void)ParamValueSize; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 8075f169347e4..ebc2817db7bba 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -69,6 +69,10 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = CONSTFIX char clGetDeviceFunctionPointerName[] = "clGetDeviceFunctionPointerINTEL"; +// Names of host pipe functions queried from OpenCL +CONSTFIX char clEnqueueReadHostPipeName[] = "clEnqueueReadHostPipeIntelFPGA"; +CONSTFIX char clEnqueueWriteHostPipeName[] = "clEnqueueWriteHostPipeIntelFPGA"; + #undef CONSTFIX // USM helper function to get an extension function pointer @@ -1305,6 +1309,94 @@ pi_result piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return RetVal; } +/// API to read host pipe +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + + return cast(0); + // TODO: Uncomment below once upstream khronos header include opencl host pipe + // API clEnqueueReadHostPipeIntelFPGA_fn FuncPtr = nullptr; pi_result RetVal = + // getExtFuncFromContext( + // queue, program, pipe_symbol, blocking, ptr, size, + // num_events_in_waitlist, events_waitlist, event); + + // if (FuncPtr) { + // RetVal = cast( + // FuncPtr(cast(queue), cast(program), + // blocking, ptr, size, + // num_events_in_waitlist, cast(events_waitlist), cast(event))); + // } + + // return RetVal; +} + +/// API to write host pipe +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory where data is to be read into +/// or written from. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + + return cast(0); + + // TODO: Uncomment below once upstream khronos header include opencl host pipe + // API clEnqueueWriteHostPipeIntelFPGA_fn FuncPtr = nullptr; pi_result RetVal + // = + // getExtFuncFromContext( + // queue, program, pipe_symbol, blocking, ptr, size, + // num_events_in_waitlist, events_waitlist, event); + + // if (FuncPtr) { + // RetVal = cast( + // FuncPtr(cast(queue), cast(program), + // blocking, ptr, size, + // num_events_in_waitlist, cast(events_waitlist), cast(event))); + // } + + // return RetVal; +} + /// API to set attributes controlling kernel execution /// /// \param kernel is the pi kernel to execute @@ -1538,6 +1630,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueueMemAdvise, piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, piextKernelSetArgSampler) _PI_CL(piTearDown, piTearDown) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index b89835c5f4864..25113de20524f 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -143,6 +143,7 @@ set(SYCL_SOURCES "detail/global_handler.cpp" "detail/helpers.cpp" "detail/handler_proxy.cpp" + "detail/host_pipe.cpp" "detail/image_accessor_util.cpp" "detail/image_impl.cpp" "detail/kernel_impl.cpp" diff --git a/sycl/source/detail/host_pipe.cpp b/sycl/source/detail/host_pipe.cpp new file mode 100644 index 0000000000000..2fc2aa72fc643 --- /dev/null +++ b/sycl/source/detail/host_pipe.cpp @@ -0,0 +1,67 @@ +//==-------------------- host_pipe_map.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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +template +_dataT host_pipe<_name, _dataT, PropertyList>::read(queue &q, + memory_order order) { + const device Dev = q.get_device(); + bool IsReadPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsReadPipeSupported) { + return &_dataT(); + } + // TODO: Get pipe name from template, get host pointer by quering the host + // pipe registration / host pipe mapping + _dataT data; + const std::string pipe_name = std::to_string(ID); + size_t size = 4; + // event e = q.read_host_pipe(name, (void*)pipe_ptr, (size_t)size, false); + event e = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, (void *)(&data), (size_t)size, false, + false /* write */); + }); + e.wait(); + return data; +} + +template +void host_pipe<_name, _dataT, PropertyList>::write(queue &q, const _dataT &data, + memory_order order) { + const device Dev = q.get_device(); + bool IsReadPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsReadPipeSupported) { + return; + } + // TODO: is this the correct logical name of the pipe? + const std::string pipe_name = std::to_string(ID); + const void *data_ptr = &data; + size_t size = 4; + // event e = q.write_host_pipe(name, (void*)pipe_ptr, (size_t)size, false); + event e = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, (void *)data_ptr, (size_t)size, false, + false /* write */); + }); + e.wait(); +} + +// TODO: implement non blocking version + +} // namespace experimental +} // namespace intel +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index f9b36419f6294..58edb08a0ff2d 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2118,6 +2118,39 @@ cl_int enqueueImpKernel( return PI_SUCCESS; } +cl_uint enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + const std::string &KernelName, + const std::string &PipeName, bool blocking, + void *ptr, size_t size, + std::vector &RawEvents, + RT::PiEvent *OutEvent, bool read) { + detail::OSModuleHandle M = + detail::OSUtil::getOSModuleHandle("HostPipeReadWriteKernelName"); + RT::PiProgram Program = + sycl::detail::ProgramManager::getInstance().getBuiltPIProgram( + M, Queue->getContextImplPtr(), Queue->getDeviceImplPtr(), KernelName); + + // Get plugin for calling opencl functions + const detail::plugin &Plugin = Queue->getPlugin(); + + pi_queue pi_q = Queue->getHandleRef(); + pi_result Error; + if (read) { + Error = + Plugin.call_nocheck( + pi_q, Program, PipeName.c_str(), blocking, ptr, size, + RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0], + OutEvent); + } else { + Error = + Plugin.call_nocheck( + pi_q, Program, PipeName.c_str(), blocking, ptr, size, + RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0], + OutEvent); + } + return Error; +} + cl_int ExecCGCommand::enqueueImp() { if (getCG().getType() != CG::CGTYPE::CodeplayHostTask) waitForPreparedHostEvents(); @@ -2479,6 +2512,23 @@ cl_int ExecCGCommand::enqueueImp() { return PI_SUCCESS; } + case CG::CGTYPE::ReadWriteHostPipe: { + CGReadWriteHostPipe *ExecReadWriteHostPipe = + (CGReadWriteHostPipe *)MCommandGroup.get(); + std::string pipeName = ExecReadWriteHostPipe->getPipeName(); + void *hostPtr = ExecReadWriteHostPipe->getHostPtr(); + size_t typeSize = ExecReadWriteHostPipe->getTypeSize(); + bool blocking = ExecReadWriteHostPipe->isBlocking(); + bool read = ExecReadWriteHostPipe->isReadHostPipe(); + + if (!Event) { + Event = &MEvent->getHandleRef(); + } + + return enqueueReadWriteHostPipe(MQueue, "ReadWriteHostPipeKernelName", + pipeName, blocking, hostPtr, typeSize, + RawEvents, Event, read); + } case CG::CGTYPE::None: throw runtime_error("CG type not implemented.", PI_INVALID_OPERATION); } diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 0a4a55579c078..0f7eb82e9d733 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -526,6 +526,13 @@ class MemCpyCommandHost : public Command { void **MDstPtr = nullptr; }; +cl_uint enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + const std::string &KernelName, + const std::string &PipeName, bool blocking, + void *ptr, size_t size, + std::vector &RawEvents, + RT::PiEvent *OutEvent, bool read); + cl_int enqueueImpKernel( const QueueImplPtr &Queue, NDRDescT &NDRDesc, std::vector &Args, const std::shared_ptr &KernelBundleImplPtr, diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index aa349c90a6f33..a080e2f766a97 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -381,6 +381,13 @@ event handler::finalize() { std::move(MAccStorage), std::move(MSharedPtrStorage), std::move(MRequirements), std::move(MEvents), MCGType, MCodeLoc)); break; + case detail::CG::ReadWriteHostPipe: + CommandGroup.reset(new detail::CGReadWriteHostPipe( + HostPipeName, HostPipeBlocking, HostPipePtr, HostPipeTypeSize, + HostPipeRead, std::move(MArgsStorage), std::move(MAccStorage), + std::move(MSharedPtrStorage), std::move(MRequirements), + std::move(MEvents), MCodeLoc)); + break; case detail::CG::None: if (detail::pi::trace(detail::pi::TraceLevel::PI_TRACE_ALL)) { std::cout << "WARNING: An empty command group is submitted." << std::endl; @@ -733,6 +740,17 @@ void handler::memcpy(void *Dest, const void *Src, size_t Count) { setType(detail::CG::CopyUSM); } +void handler::read_write_host_pipe(std::string Name, void *Ptr, size_t Size, + bool Block, bool Read) { + throwIfActionIsCreated(); + HostPipeName = Name; + HostPipePtr = Ptr; + HostPipeTypeSize = Size; + HostPipeBlocking = Block; + HostPipeRead = Read; + setType(detail::CG::ReadWriteHostPipe); +} + void handler::memset(void *Dest, int Value, size_t Count) { throwIfActionIsCreated(); MDstPtr = Dest;