diff --git a/buildbot/dependency.py b/buildbot/dependency.py index ec8dc9d423fe4..8ab0e08bb9c16 100644 --- a/buildbot/dependency.py +++ b/buildbot/dependency.py @@ -49,8 +49,8 @@ def do_dependency(args): # fetch OpenCL headers ocl_header_dir = os.path.join(args.obj_dir, "OpenCL-Headers") if not os.path.isdir(ocl_header_dir): - clone_cmd = ["git", "clone", "https://github.com/KhronosGroup/OpenCL-Headers", - "OpenCL-Headers", "-b", "main"] + clone_cmd = ["git", "clone", "https://github.com/sherry-yuan/OpenCL-Headers", + "OpenCL-Headers", "-b", "host_pipe"] # TODO: Remove change once upstream header changed subprocess.check_call(clone_cmd, cwd=args.obj_dir) else: fetch_cmd = ["git", "pull", "--ff", "--ff-only", "origin"] @@ -58,7 +58,7 @@ def do_dependency(args): # Checkout fixed version to avoid unexpected issues coming from upstream # Specific version can be uplifted as soon as such need arise - checkout_cmd = ["git", "checkout", "23710f1b99186065c1768fc3098ba681adc0f253"] + checkout_cmd = ["git", "checkout", "1f2cb76195fb77be7c0b4d811ecff244c864d2e2"] # TODO: Remove change once upstream header changed subprocess.check_call(checkout_cmd, cwd=ocl_header_dir) # fetch and build OpenCL ICD loader diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 310afd5f4cfa3..0a8852357cd36 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -193,6 +193,7 @@ class PropertySetRegistry { static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used"; static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols"; static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals"; + static constexpr char SYCL_HOST_PIPES[] = "SYCL/host pipes"; // Function for bulk addition of an entire property set under given category // (property set name). diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index 3541c11a8d7d2..4dc6a048e137c 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -202,6 +202,7 @@ constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; constexpr char PropertySetRegistry::SYCL_ASSERT_USED[]; constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[]; constexpr char PropertySetRegistry::SYCL_DEVICE_GLOBALS[]; +constexpr char PropertySetRegistry::SYCL_HOST_PIPES[]; } // namespace util } // namespace llvm diff --git a/opencl/CMakeLists.txt b/opencl/CMakeLists.txt index 4115fbb45e59a..00bc0c590de5b 100644 --- a/opencl/CMakeLists.txt +++ b/opencl/CMakeLists.txt @@ -14,14 +14,14 @@ endif() # Repo URLs set(OCL_HEADERS_REPO - "https://github.com/KhronosGroup/OpenCL-Headers.git") + "https://github.com/sherry-yuan/OpenCL-Headers.git") set(OCL_LOADER_REPO "https://github.com/KhronosGroup/OpenCL-ICD-Loader.git") # Repo tags/hashes -set(OCL_HEADERS_TAG dcd5bede6859d26833cd85f0d6bbcee7382dc9b3) -set(OCL_LOADER_TAG 5d9177ee79bfbcc75ee9a8cff6415eab2c3113f6) +set(OCL_HEADERS_TAG 1f2cb76195fb77be7c0b4d811ecff244c864d2e2) +set(OCL_LOADER_TAG 5f8249691ec8c25775789498951f8e9eb62c201d) # OpenCL Headers if(NOT OpenCL_HEADERS) diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp index 8ca7d28223cad..5c55b6c6738e9 100644 --- a/sycl/include/CL/sycl.hpp +++ b/sycl/include/CL/sycl.hpp @@ -71,3 +71,6 @@ #include #include #include + +#include +#include diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index ba70ddefcee2e..95b93356271e8 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, @@ -522,6 +523,36 @@ class CGBarrier : public CG { MEventsWaitWithBarrier(std::move(EventsWaitWithBarrier)) {} }; +/// "ReadWriteHostPipe" command group class. +class CGReadWriteHostPipe : public CG { + std::string PipeName; + bool Blocking; + void *HostPtr; + size_t TypeSize; + bool IsReadOp; + +public: + CGReadWriteHostPipe(const 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; } +}; + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/detail/host_pipe_map.hpp b/sycl/include/CL/sycl/detail/host_pipe_map.hpp new file mode 100755 index 0000000000000..9c82f11bf0a76 --- /dev/null +++ b/sycl/include/CL/sycl/detail/host_pipe_map.hpp @@ -0,0 +1,21 @@ +//==-------------------- host_pipe_map.hpp -----------------------------==// +// +// 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +namespace host_pipe_map { + +__SYCL_EXPORT void add(const void *HostPipePtr, const char *UniqueId); + +} // namespace host_pipe_map +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/detail/pi.def b/sycl/include/CL/sycl/detail/pi.def index c9a68c6cadec3..a0ee268af01e5 100644 --- a/sycl/include/CL/sycl/detail/pi.def +++ b/sycl/include/CL/sycl/detail/pi.def @@ -130,6 +130,9 @@ _PI_API(piextUSMEnqueueMemcpy) _PI_API(piextUSMEnqueuePrefetch) _PI_API(piextUSMEnqueueMemAdvise) _PI_API(piextUSMGetMemAllocInfo) +// Host pipes +_PI_API(piextEnqueueReadHostPipe) +_PI_API(piextEnqueueWriteHostPipe) _PI_API(piextKernelSetArgMemObj) _PI_API(piextKernelSetArgSampler) diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index ad2217368ddb1..b7f91300d0784 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -779,6 +779,8 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4; #define __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS "SYCL/exported symbols" /// PropertySetRegistry::SYCL_DEVICE_GLOBALS defined in PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS "SYCL/device globals" +/// PropertySetRegistry::SYCL_HOST_PIPES defined in PropertySetIO.h +#define __SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES "SYCL/host pipes" /// Program metadata tags recognized by the PI backends. For kernels the tag /// must appear after the kernel name. @@ -1782,6 +1784,56 @@ __SYCL_EXPORT pi_result piextUSMGetMemAllocInfo( pi_context context, const void *ptr, pi_mem_alloc_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 that will hold resulting data +/// from pipe +/// @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. +__SYCL_EXPORT 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 that holds data to be written +/// to host pipe. +/// @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. +__SYCL_EXPORT 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/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp index dbb7adabd2c26..c2e7a96eaf946 100644 --- a/sycl/include/CL/sycl/detail/pi.hpp +++ b/sycl/include/CL/sycl/detail/pi.hpp @@ -383,6 +383,13 @@ class DeviceBinaryImage { DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS); return DeviceGlobals; } + const PropertyRange getHostPipes() const { + // We can't have this variable as a class member, since it would break + // the ABI backwards compatibility. + DeviceBinaryImage::PropertyRange HostPipes; + HostPipes.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES); + return HostPipes; + } virtual ~DeviceBinaryImage() {} protected: diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 26dd0e718ba95..51a472b6d4f0e 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -2612,6 +2612,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 as identified by address of its const + /// expr __pipe member \param Size the size of data getting read back / to. + /// /// \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(const std::string &Name, void *Ptr, size_t Size, + bool Block, bool Read); + private: std::shared_ptr MQueue; /// The storage for the arguments passed. @@ -2660,6 +2669,16 @@ class __SYCL_EXPORT handler { /// The list of valid SYCL events that need to complete /// before barrier command can be executed std::vector MEventsWaitWithBarrier; + /// Pipe name that uniquely identifies a pipe. + std::string HostPipeName; + /// Pipe host pointer, the address of its constexpr __pipe member. + void *HostPipePtr = nullptr; + /// Host pipe read write operation is blocking. + bool HostPipeBlocking = false; + /// The size of returned type for each read. + size_t HostPipeTypeSize = 0; + /// If the pipe operation is read or write, 1 for read 0 for 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..e45d82f5896d2 --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/host_pipes.hpp @@ -0,0 +1,91 @@ +//==---------------- 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 + +#ifdef XPTI_ENABLE_INSTRUMENTATION +#include +#include +#endif + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +using default_pipe_properties = + decltype(sycl::ext::oneapi::experimental::properties(min_capacity<0>)); + +template +class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_attributes_global_variable("sycl-host-access", + "readwrite")]] +#endif + // TODO: change name to pipe, and merge into the existing pipe + // implementation + host_pipe { + + struct +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_global_variable_attributes( + "sycl-host-pipe", + nullptr)]] [[__sycl_detail__:: + host_pipe]] [[__sycl_detail__:: + global_variable_allowed]] // may + // not be + // needed +#endif + __pipeType { + const char __p; + }; + + static constexpr __pipeType __pipe = {0}; + +public: + using value_type = _dataT; + static constexpr int32_t min_cap = + _propertiesT::template has_property() + ? _propertiesT::template get_property().value + : 0; + + static const void *get_host_ptr() { return &__pipe; } + + // 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); + +#ifdef __SYCL_DEVICE_ONLY__ + static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment, + min_cap}; +#endif // __SYCL_DEVICE_ONLY__ +}; + +} // namespace experimental +} // namespace intel +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) \ No newline at end of file diff --git a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp new file mode 100644 index 0000000000000..e4e012ccd972d --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp @@ -0,0 +1,197 @@ +//==----- pipe_properties.hpp - SYCL properties associated with data flow pipe +//---==// +// +// 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +struct min_capacity_key { + template + using value_t = oneapi::experimental::property_value< + min_capacity_key, std::integral_constant>; +}; + +struct ready_latency_key { + template + using value_t = oneapi::experimental::property_value< + ready_latency_key, std::integral_constant>; +}; + +struct bits_per_symbol_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct uses_valid_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct uses_ready_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct in_csr_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct first_symbol_in_high_order_bits_key { + template + using value_t = oneapi::experimental::property_value< + first_symbol_in_high_order_bits_key, + sycl::detail::bool_constant>; +}; + +enum class protocol_name : std::uint16_t { AVALON, AXI }; +struct protocol_key { + template + using value_t = oneapi::experimental::property_value< + protocol_key, std::integral_constant>; +}; + +template +inline constexpr min_capacity_key::value_t min_capacity; + +template +inline constexpr ready_latency_key::value_t ready_latency; + +template +inline constexpr bits_per_symbol_key::value_t bits_per_symbol; + +template +inline constexpr uses_valid_key::value_t uses_valid; +inline constexpr uses_valid_key::value_t uses_valid_on; +inline constexpr uses_valid_key::value_t uses_valid_off; + +template +inline constexpr uses_ready_key::value_t uses_ready; +inline constexpr uses_ready_key::value_t uses_ready_on; +inline constexpr uses_ready_key::value_t uses_ready_off; + +template inline constexpr in_csr_key::value_t in_csr; +inline constexpr in_csr_key::value_t in_csr_on; +inline constexpr in_csr_key::value_t in_csr_off; + +template +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits; +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits_on; +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits_off; + +template +inline constexpr protocol_key::value_t protocol; +inline constexpr protocol_key::value_t protocol_avalon; +inline constexpr protocol_key::value_t protocol_axi; + +} // namespace experimental +} // namespace intel + +namespace oneapi { +namespace experimental { + +template <> +struct is_property_key : std::true_type { +}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; + +namespace detail { +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::MinCapacity; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::ReadyLatency; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::BitsPerSymbol; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::UsesValid; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::UsesReady; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::ImplementInCSR; +}; +template <> +struct PropertyToKind< + intel::experimental::first_symbol_in_high_order_bits_key> { + static constexpr PropKind Kind = PropKind::FirstSymbolInHigherOrderBit; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::PipeProtocol; +}; + +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty : std::true_type { +}; +template <> +struct IsCompileTimeProperty< + intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type { +}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; + +} // namespace detail +} // namespace experimental +} // namespace oneapi +} // 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..bde7ccd85c6b1 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -148,7 +148,14 @@ enum PropKind : uint32_t { HostAccess = 1, InitMode = 2, ImplementInCSR = 3, - PropKindSize = 4, + BitsPerSymbol = 4, + FirstSymbolInHigherOrderBit = 5, + MinCapacity = 6, + PipeProtocol = 7, + ReadyLatency = 8, + UsesReady = 9, + UsesValid = 10, + PropKindSize = 11, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 239d4be03faee..23c07830a471a 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -4949,6 +4949,43 @@ pi_result cuda_piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return result; } +/// Host Pipes +pi_result cuda_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) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + cl::sycl::detail::pi::die("cuda_piextEnqueueReadHostPipe not implemented"); + return {}; +} + +pi_result cuda_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) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + cl::sycl::detail::pi::die("cuda_piextEnqueueWriteHostPipe not implemented"); + return {}; +} + // This API is called by Sycl RT to notify the end of the plugin lifetime. // TODO: add a global variable lifetime management code here (see // pi_level_zero.cpp for reference) Currently this is just a NOOP. @@ -5091,6 +5128,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueueMemAdvise, cuda_piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, cuda_piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, cuda_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, cuda_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler) _PI_CL(piTearDown, cuda_piTearDown) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 88bab8fa142e8..64d486f00c0b2 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1887,6 +1887,25 @@ pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_alloc_info, DIE_NO_IMPLEMENTATION; } +/// Host Pipes +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 piKernelSetExecInfo(pi_kernel, pi_kernel_exec_info, size_t, const void *) { DIE_NO_IMPLEMENTATION; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index a715c7e064386..234f222b3da96 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -4840,6 +4840,45 @@ pi_result hip_piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return result; } +/// Host Pipes +pi_result hip_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) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + cl::sycl::detail::pi::die("hip_piextEnqueueReadHostPipe not implemented"); + return {}; +} + +pi_result hip_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) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + cl::sycl::detail::pi::die("hip_piextEnqueueWriteHostPipe not implemented"); + return {}; +} + // This API is called by Sycl RT to notify the end of the plugin lifetime. // TODO: add a global variable lifetime management code here (see // pi_level_zero.cpp for reference) Currently this is just a NOOP. @@ -4981,6 +5020,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueueMemAdvise, hip_piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, hip_piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, hip_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, hip_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, hip_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, hip_piextKernelSetArgSampler) _PI_CL(piTearDown, hip_piTearDown) diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 28e00dd73468c..86c767d4219f6 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -7888,6 +7888,72 @@ pi_result piextUSMGetMemAllocInfo(pi_context Context, const void *Ptr, return PI_SUCCESS; } +/// API for Read from host pipe. +/// +/// \param Queue is the queue +/// \param Program is the program containing the device variable +/// \param PipeSymbol is the unique identifier for the device variable +/// \param Blocking is true if the write should block +/// \param Ptr is a pointer to where the data will be copied to +/// \param Size is size of the data that is read/written from/to pipe +/// \param NumEventsInWaitList is a number of events in the wait list +/// \param EventWaitList is the wait list +/// \param Event is the resulting event +pi_result piextEnqueueReadHostPipe(pi_queue Queue, pi_program Program, + const char *PipeSymbol, pi_bool Blocking, + void *Ptr, size_t Size, + pi_uint32 NumEventsInWaitList, + const pi_event *EventsWaitList, + pi_event *Event) { + (void)Queue; + (void)Program; + (void)PipeSymbol; + (void)Blocking; + (void)Ptr; + (void)Size; + (void)NumEventsInWaitList; + (void)EventsWaitList; + (void)Event; + + PI_ASSERT(Queue, PI_INVALID_QUEUE); + + die("piextEnqueueReadHostPipe: not implemented"); + return {}; +} + +/// API for write to pipe of a given name. +/// +/// \param Queue is the queue +/// \param Program is the program containing the device variable +/// \param PipeSymbol is the unique identifier for the device variable +/// \param Blocking is true if the write should block +/// \param Ptr is a pointer to where the data must be copied from +/// \param Size is size of the data that is read/written from/to pipe +/// \param NumEventsInWaitList is a number of events in the wait list +/// \param EventWaitList is the wait list +/// \param Event is the resulting event +pi_result piextEnqueueWriteHostPipe(pi_queue Queue, pi_program Program, + const char *PipeSymbol, pi_bool Blocking, + void *Ptr, size_t Size, + pi_uint32 NumEventsInWaitList, + const pi_event *EventsWaitList, + pi_event *Event) { + (void)Queue; + (void)Program; + (void)PipeSymbol; + (void)Blocking; + (void)Ptr; + (void)Size; + (void)NumEventsInWaitList; + (void)EventsWaitList; + (void)Event; + + PI_ASSERT(Queue, PI_INVALID_QUEUE); + + die("piextEnqueueWriteHostPipe: not implemented"); + return {}; +} + 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 7adcc79fc8c65..40370ea6729cf 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -68,6 +68,9 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = "clSetProgramSpecializationConstant"; CONSTFIX char clGetDeviceFunctionPointerName[] = "clGetDeviceFunctionPointerINTEL"; +// Names of host pipe functions queried from OpenCL +CONSTFIX char clEnqueueReadHostPipeName[] = "clEnqueueReadHostPipeIntelFPGA"; +CONSTFIX char clEnqueueWriteHostPipeName[] = "clEnqueueWriteHostPipeIntelFPGA"; #undef CONSTFIX @@ -1308,6 +1311,64 @@ pi_result piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return RetVal; } +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) { + cl_context CLContext; + cl_int CLErr = + clGetCommandQueueInfo(cast(queue), CL_QUEUE_CONTEXT, + sizeof(cl_context), &CLContext, nullptr); + if (CLErr != CL_SUCCESS) { + return cast(CLErr); + } + + clEnqueueReadHostPipeIntelFPGA_fn FuncPtr = nullptr; + pi_result RetVal = getExtFuncFromContext( + cast(CLContext), &FuncPtr); + + if (FuncPtr) { + RetVal = cast(FuncPtr( + cast(queue), cast(program), pipe_symbol, + blocking, ptr, size, num_events_in_waitlist, + cast(events_waitlist), cast(event))); + } + + return RetVal; +} + +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) { + cl_context CLContext; + cl_int CLErr = + clGetCommandQueueInfo(cast(queue), CL_QUEUE_CONTEXT, + sizeof(cl_context), &CLContext, nullptr); + if (CLErr != CL_SUCCESS) { + return cast(CLErr); + } + + clEnqueueWriteHostPipeIntelFPGA_fn FuncPtr = nullptr; + pi_result RetVal = getExtFuncFromContext( + cast(CLContext), &FuncPtr); + + if (FuncPtr) { + RetVal = cast(FuncPtr( + cast(queue), cast(program), pipe_symbol, + 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 @@ -1540,6 +1601,9 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueuePrefetch, piextUSMEnqueuePrefetch) _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) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 5f23d77b1126a..804b37d31e491 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -132,6 +132,7 @@ set(SYCL_SOURCES "detail/context_impl.cpp" "detail/device_binary_image.cpp" "detail/device_filter.cpp" + "detail/host_pipe_map.cpp" "detail/device_global_map.cpp" "detail/device_impl.cpp" "detail/error_handling/enqueue_kernel.cpp" @@ -141,6 +142,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..5e2d0ab54be7b --- /dev/null +++ b/sycl/source/detail/host_pipe.cpp @@ -0,0 +1,70 @@ +//==-------------------- host_pipe.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 +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +template +_dataT host_pipe<_name, _dataT, _propertiesT>::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 the pipe registration + _dataT data; + const void *HostPipePtr = &__pipe; + detail::HostPipeMapEntry *hostPipeEntry = + detail::ProgramManager::getInstance().getHostPipeEntry(HostPipePtr); + const std::string pipe_name = hostPipeEntry->MUniqueId; + event e = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, (void *)(&data), sizeof(_dataT), false, + true /* read */); + }); + e.wait(); + return data; +} + +template +void host_pipe<_name, _dataT, _propertiesT>::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: get pipe name from the pipe registration + const void *HostPipePtr = &__pipe; + detail::HostPipeMapEntry *hostPipeEntry = + detail::ProgramManager::getInstance().getHostPipeEntry(HostPipePtr); + const std::string pipe_name = hostPipeEntry->MUniqueId; + const void *data_ptr = &data; + event e = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, (void *)data_ptr, sizeof(_dataT), false, + false /* write */); + }); + e.wait(); +} + +// TODO: implement non blocking version + +} // namespace experimental +} // namespace intel +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) \ No newline at end of file diff --git a/sycl/source/detail/host_pipe_map.cpp b/sycl/source/detail/host_pipe_map.cpp new file mode 100644 index 0000000000000..79f7f77f70716 --- /dev/null +++ b/sycl/source/detail/host_pipe_map.cpp @@ -0,0 +1,24 @@ +//==-------------------- 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 detail { +namespace host_pipe_map { + +__SYCL_EXPORT void add(const void *HostPipePtr, const char *UniqueId) { + detail::ProgramManager::getInstance().addOrInitHostPipeEntry(HostPipePtr, + UniqueId); +} + +} // namespace host_pipe_map +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/host_pipe_map_entry.hpp b/sycl/source/detail/host_pipe_map_entry.hpp new file mode 100644 index 0000000000000..a02653106e848 --- /dev/null +++ b/sycl/source/detail/host_pipe_map_entry.hpp @@ -0,0 +1,56 @@ +//==----------------- host_pipe_map_entry.hpp --------------------------==// +// +// 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { + +struct HostPipeMapEntry { + std::string MUniqueId; + // Pointer to the host_pipe on host. + const void *MHostPipePtr; + // Size of the underlying type in the host_pipe. + std::uint32_t MHostPipeTSize; + // The device image that pipe is associated with + const RTDeviceBinaryImage *mDeviceImage; + + // Constructor only initializes with the pointer and ID. + // Other members will be initialized later + HostPipeMapEntry(std::string UniqueId, const void *HostPipePtr) + : MUniqueId(UniqueId), MHostPipePtr(HostPipePtr), MHostPipeTSize(0) {} + + // Constructor only initializes with the size and ID. + // Other members will be initialized later + HostPipeMapEntry(std::string UniqueId, std::uint32_t HostPipeTSize) + : MUniqueId(UniqueId), MHostPipePtr(nullptr), + MHostPipeTSize(HostPipeTSize) {} + + void initialize(std::uint32_t HostPipeTSize) { + assert(HostPipeTSize != 0 && "Host pipe initialized with 0 size."); + assert(MHostPipeTSize == 0 && "Host pipe has already been initialized."); + MHostPipeTSize = HostPipeTSize; + } + + void initialize(const void *HostPipePtr) { + assert(!MHostPipePtr && "Host pipe pointer has already been initialized."); + MHostPipePtr = HostPipePtr; + } + void initialize(const RTDeviceBinaryImage *DeviceImage) { + mDeviceImage = DeviceImage; + } +}; + +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 8108b642c823b..3f6b08e2c06ec 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1229,6 +1229,42 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) { } } } + // ... and initialize associated host_pipe information + { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + + auto HostPipes = Img->getHostPipes(); + for (const pi_device_binary_property &HostPipe : HostPipes) { + pi::ByteArray HostPipeInfo = + pi::DeviceBinaryProperty(HostPipe).asByteArray(); + + // The supplied host_pipe info property is expected to contain: + // * 8 bytes - Size of the property. + // * 4 bytes - Size of the underlying type in the host_pipe. + // Note: Property may be padded. + constexpr unsigned int NumPropertySizeBytes = 8; + constexpr unsigned int NumTypeBytes = 4; + assert(HostPipeInfo.size() >= NumPropertySizeBytes + NumTypeBytes && + "Unexpected property size"); + auto TypeSize = *reinterpret_cast( + &HostPipeInfo[NumPropertySizeBytes]); + + auto ExistingHostPipe = m_HostPipes.find(HostPipe->Name); + if (ExistingHostPipe != m_HostPipes.end()) { + // If it has already been registered we update the information. + ExistingHostPipe->second->initialize(TypeSize); + ExistingHostPipe->second->initialize(Img.get()); + } else { + // If it has not already been registered we create a new entry. + // Note: Pointer to the host pipe is not available here, so it + // cannot be set until registration happens. + auto EntryUPtr = + std::make_unique(HostPipe->Name, TypeSize); + EntryUPtr->initialize(Img.get()); + m_HostPipes.emplace(HostPipe->Name, std::move(EntryUPtr)); + } + } + } m_DeviceImages[KSId].reset(new std::vector()); cacheKernelUsesAssertInfo(M, *Img); @@ -1497,6 +1533,39 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr, m_Ptr2DeviceGlobal.insert({DeviceGlobalPtr, NewEntry.first->second.get()}); } +void ProgramManager::addOrInitHostPipeEntry(const void *HostPipePtr, + const char *UniqueId) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + + assert(m_HostPipes.find(UniqueId) == m_HostPipes.end() && + "Host pipe has already been registered."); + auto ExistingHostPipe = m_HostPipes.find(UniqueId); + if (ExistingHostPipe != m_HostPipes.end()) { + ExistingHostPipe->second->initialize(HostPipePtr); + m_Ptr2HostPipe.insert({HostPipePtr, ExistingHostPipe->second.get()}); + return; + } + + auto EntryUPtr = std::make_unique(UniqueId, HostPipePtr); + auto NewEntry = m_HostPipes.emplace(UniqueId, std::move(EntryUPtr)); + m_Ptr2HostPipe.insert({HostPipePtr, NewEntry.first->second.get()}); +} + +HostPipeMapEntry * +ProgramManager::getHostPipeEntry(const std::string &UniqueId) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + auto Entry = m_HostPipes.find(UniqueId); + assert(Entry != m_HostPipes.end() && "Host pipe entry not found"); + return Entry->second.get(); +} + +HostPipeMapEntry *ProgramManager::getHostPipeEntry(const void *HostPipePtr) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + auto Entry = m_Ptr2HostPipe.find(HostPipePtr); + assert(Entry != m_Ptr2HostPipe.end() && "Host pipe entry not found"); + return Entry->second; +} + std::vector ProgramManager::getSYCLDeviceImagesWithCompatibleState( const context &Ctx, const std::vector &Devs, diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 504162e5deae8..ab85d86bb3c76 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -10,6 +10,7 @@ #include #include #include +#include #include #include #include @@ -17,6 +18,7 @@ #include #include #include +#include #include #include @@ -188,6 +190,18 @@ class ProgramManager { void addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr, const char *UniqueId); + // The function inserts or initializes a host_pipe entry into the + // host_pipe map. + void addOrInitHostPipeEntry(const void *HostPipePtr, const char *UniqueId); + + // The function gets a host_pipe entry identified by the unique ID from + // the host_pipe map. + HostPipeMapEntry *getHostPipeEntry(const std::string &UniqueId); + + // The function gets a host_pipe entry identified by the pointer to the + // host_pipe object from the host_pipe map. + HostPipeMapEntry *getHostPipeEntry(const void *HostPipePtr); + // The function returns a vector of SYCL device images that are compiled with // the required state and at least one device from the passed list of devices. std::vector getSYCLDeviceImagesWithCompatibleState( @@ -396,6 +410,14 @@ class ProgramManager { /// Protects m_DeviceGlobals and m_Ptr2DeviceGlobal. std::mutex m_DeviceGlobalsMutex; + + // Maps between host_pipe identifiers and associated information. + std::unordered_map> + m_HostPipes; + std::unordered_map m_Ptr2HostPipe; + + /// Protects m_HostPipes and m_Ptr2HostPipe. + std::mutex m_HostPipesMutex; }; } // namespace detail } // namespace sycl diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 6138b21723a57..e54ca1f872645 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -2192,6 +2193,43 @@ cl_int enqueueImpKernel( return PI_SUCCESS; } +cl_uint enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + const std::string &PipeName, bool blocking, + void *ptr, size_t size, + std::vector &RawEvents, + RT::PiEvent *OutEvent, bool read) { + // TODO: Few options of getting the kernel name / program object: + // 1. Encode this in the pipe registration + // 2. Initialize the pipe registration from first kernel launch, but then this + // will violate the spec + detail::HostPipeMapEntry *hostPipeEntry = + detail::ProgramManager::getInstance().getHostPipeEntry(PipeName); + RT::PiProgram Program = + sycl::detail::ProgramManager::getInstance().createPIProgram( + *(hostPipeEntry->mDeviceImage), Queue->get_context(), + Queue->get_device()); + + // 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(); @@ -2554,6 +2592,22 @@ 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, 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 75b2ae82ed3ef..b0e70038f621d 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -526,6 +526,12 @@ class MemCpyCommandHost : public Command { void **MDstPtr = nullptr; }; +cl_uint enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + 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..3358cc060a262 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; @@ -814,5 +821,16 @@ void handler::depends_on(const std::vector &Events) { } } +void handler::read_write_host_pipe(const 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); +} + } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/layout_handler.cpp b/sycl/test/abi/layout_handler.cpp index 72bdcf382fde8..67319577fdf3a 100644 --- a/sycl/test/abi/layout_handler.cpp +++ b/sycl/test/abi/layout_handler.cpp @@ -172,19 +172,32 @@ void foo() { // CHECK: 472 | std::_Vector_base, class std::allocator > >::pointer _M_start // CHECK-NEXT: 480 | std::_Vector_base, class std::allocator > >::pointer _M_finish // CHECK-NEXT: 488 | std::_Vector_base, class std::allocator > >::pointer _M_end_of_storage -// CHECK-NEXT: 496 | _Bool MIsHost -// CHECK-NEXT: 504 | struct sycl::detail::code_location MCodeLoc -// CHECK-NEXT: 504 | const char * MFileName -// CHECK-NEXT: 512 | const char * MFunctionName -// CHECK-NEXT: 520 | unsigned long MLineNo -// CHECK-NEXT: 528 | unsigned long MColumnNo -// CHECK-NEXT: 536 | _Bool MIsFinalized -// CHECK-NEXT: 544 | class sycl::event MLastEvent -// CHECK-NEXT: 544 | class std::shared_ptr impl -// CHECK-NEXT: 544 | class std::__shared_ptr (base) -// CHECK-NEXT: 544 | class std::__shared_ptr_access (base) (empty) -// CHECK-NEXT: 544 | std::__shared_ptr::element_type * _M_ptr -// CHECK-NEXT: 552 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount -// CHECK-NEXT: 552 | _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi -// CHECK-NEXT: | [sizeof=560, dsize=560, align=8, -// CHECK-NEXT: | nvsize=560, nvalign=8] +// CHECK-NEXT: 496 | class std::basic_string HostPipeName +// CHECK-NEXT: 496 | struct std::basic_string::_Alloc_hider _M_dataplus +// CHECK-NEXT: 496 | class std::allocator (base) (empty) +// CHECK-NEXT: 496 | class __gnu_cxx::new_allocator (base) (empty) +// CHECK-NEXT: 496 | std::basic_string::pointer _M_p +// CHECK-NEXT: 504 | std::basic_string::size_type _M_string_length +// CHECK-NEXT: 512 | union std::basic_string::(anonymous at /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/basic_string.h:175:7) +// CHECK-NEXT: 512 | char[16] _M_local_buf +// CHECK-NEXT: 512 | std::basic_string::size_type _M_allocated_capacity +// CHECK-NEXT: 528 | void * HostPipePtr +// CHECK-NEXT: 536 | _Bool HostPipeBlocking +// CHECK-NEXT: 544 | size_t HostPipeTypeSize +// CHECK-NEXT: 552 | _Bool HostPipeRead +// CHECK-NEXT: 553 | _Bool MIsHost +// CHECK-NEXT: 560 | struct sycl::detail::code_location MCodeLoc +// CHECK-NEXT: 560 | const char * MFileName +// CHECK-NEXT: 568 | const char * MFunctionName +// CHECK-NEXT: 576 | unsigned long MLineNo +// CHECK-NEXT: 584 | unsigned long MColumnNo +// CHECK-NEXT: 592 | _Bool MIsFinalized +// CHECK-NEXT: 600 | class sycl::event MLastEvent +// CHECK-NEXT: 600 | class std::shared_ptr impl +// CHECK-NEXT: 600 | class std::__shared_ptr (base) +// CHECK-NEXT: 600 | class std::__shared_ptr_access (base) (empty) +// CHECK-NEXT: 600 | std::__shared_ptr::element_type * _M_ptr +// CHECK-NEXT: 608 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount +// CHECK-NEXT: 608 | _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi +// CHECK-NEXT: | [sizeof=616, dsize=616, align=8, +// CHECK-NEXT: | nvsize=616, nvalign=8] diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index 491eb9d87f155..0443faf3b401e 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -86,6 +86,8 @@ piextContextSetExtendedDeleter piextDeviceCreateWithNativeHandle piextDeviceGetNativeHandle piextDeviceSelectBinary +piextEnqueueReadHostPipe +piextEnqueueWriteHostPipe piextEventCreateWithNativeHandle piextEventGetNativeHandle piextGetDeviceFunctionPointer diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 0e6aaa9dfa7a5..db890125035cb 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -33,6 +33,8 @@ piextContextGetNativeHandle piextDeviceCreateWithNativeHandle piextDeviceGetNativeHandle piextDeviceSelectBinary +piextEnqueueReadHostPipe +piextEnqueueWriteHostPipe piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer piextKernelCreateWithNativeHandle diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index dba6e1c3bf5db..3d368975dece9 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3883,6 +3883,7 @@ _ZN2cl4sycl6detail13MemoryManager8copy_usmEPKvSt10shared_ptrINS1_10queue_implEEm _ZN2cl4sycl6detail13MemoryManager8copy_usmEPKvSt10shared_ptrINS1_10queue_implEEmPvSt6vectorIP9_pi_eventSaISB_EERSB_ _ZN2cl4sycl6detail13MemoryManager8fill_usmEPvSt10shared_ptrINS1_10queue_implEEmiSt6vectorIP9_pi_eventSaIS9_EEPS9_ _ZN2cl4sycl6detail13MemoryManager8fill_usmEPvSt10shared_ptrINS1_10queue_implEEmiSt6vectorIP9_pi_eventSaIS9_EERS9_ +_ZN2cl4sycl6detail13host_pipe_map3addEPKvPKc _ZN2cl4sycl6detail13make_platformEmNS0_7backendE _ZN2cl4sycl6detail14getBorderColorENS0_19image_channel_orderE _ZN2cl4sycl6detail14host_half_impl4halfC1ERKf @@ -4005,6 +4006,7 @@ _ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler18extractArgsAndReqsEv _ZN2cl4sycl7handler20DisableRangeRoundingEv _ZN2cl4sycl7handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE +_ZN2cl4sycl7handler20read_write_host_pipeERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPvmbb _ZN2cl4sycl7handler20setStateSpecConstSetEv _ZN2cl4sycl7handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN2cl4sycl7handler22verifyUsedKernelBundleERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index df41bb5ddecc2..bdd8d7e8b6f75 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1094,6 +1094,7 @@ ?acospi@__host_std@cl@@YAMM@Z ?acospi@__host_std@cl@@YANN@Z ?add@device_global_map@detail@sycl@cl@@YAXPEBXPEBD@Z +?add@host_pipe_map@detail@sycl@cl@@YAXPEBXPEBD@Z ?addHostAccessorAndWait@detail@sycl@cl@@YAXPEAVAccessorImplHost@123@@Z ?addOrReplaceAccessorProperties@SYCLMemObjT@detail@sycl@cl@@QEAAXAEBVproperty_list@34@@Z ?addReduction@handler@sycl@cl@@AEAAXAEBV?$shared_ptr@$$CBX@std@@@Z diff --git a/sycl/test/abi/symbol_size_alignment.cpp b/sycl/test/abi/symbol_size_alignment.cpp index 68cd05e158745..b0138f275baca 100644 --- a/sycl/test/abi/symbol_size_alignment.cpp +++ b/sycl/test/abi/symbol_size_alignment.cpp @@ -58,11 +58,11 @@ int main() { check(); check(); #ifdef _MSC_VER - check(); + check(); check(); check, 272, 8>(); #else - check(); + check(); check(); check, 240, 8>(); #endif diff --git a/sycl/test/extensions/properties/properties_pipe.cpp b/sycl/test/extensions/properties/properties_pipe.cpp new file mode 100644 index 0000000000000..9179d9f66c003 --- /dev/null +++ b/sycl/test/extensions/properties/properties_pipe.cpp @@ -0,0 +1,168 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// expected-no-diagnostics + +#include + +#include + +using namespace sycl::ext; + +constexpr sycl::ext::intel::experimental::protocol_name TestProtocol = + sycl::ext::intel::experimental::protocol_name::AVALON; + +int main() { + // Check that is_property_key is correctly specialized. + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::min_capacity_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::ready_latency_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::bits_per_symbol_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::uses_valid_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::uses_ready_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::in_csr_key>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_key>:: + value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::protocol_key>::value); + + // Check that is_property_value is correctly specialized. + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::min_capacity<3>)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::ready_latency<3>)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::bits_per_symbol<3>)>::value); + + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid_on)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid_off)>::value); + + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_ready)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_ready_on)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_ready_off)>::value); + + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr_on)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr_off)>::value); + + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits_on)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits_off)>::value); + + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol)>:: + value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol_avalon)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol_axi)>::value); + + // Checks that fully specialized properties are the same as the templated + // variants. + static_assert(std::is_same_v< + decltype(sycl::ext::intel::experimental::uses_valid_on), + decltype(sycl::ext::intel::experimental::uses_valid)>); + static_assert(std::is_same_v< + decltype(sycl::ext::intel::experimental::uses_ready_off), + decltype(sycl::ext::intel::experimental::uses_ready)>); + static_assert( + std::is_same_v)>); + static_assert( + std::is_same_v)>); + static_assert( + std::is_same_v< + decltype(sycl::ext::intel::experimental::protocol_avalon), + decltype(sycl::ext::intel::experimental::protocol)>); + static_assert(std::is_same_v< + decltype(sycl::ext::intel::experimental::protocol_axi), + decltype(sycl::ext::intel::experimental::protocol< + sycl::ext::intel::experimental::protocol_name::AXI>)>); + + // Check that property lists will accept the new properties. + using P = decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::intel::experimental::min_capacity<0>, + sycl::ext::intel::experimental::ready_latency<1>, + sycl::ext::intel::experimental::bits_per_symbol<2>, + sycl::ext::intel::experimental::uses_valid, + sycl::ext::intel::experimental::uses_ready, + sycl::ext::intel::experimental::in_csr, + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off, + sycl::ext::intel::experimental::protocol_avalon)); + static_assert(sycl::ext::oneapi::experimental::is_property_list_v

); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert(P::has_property()); + static_assert(P::has_property()); + static_assert( + P::has_property()); + + static_assert( + P::get_property() == + sycl::ext::intel::experimental::min_capacity<0>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::ready_latency<1>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::bits_per_symbol<2>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::uses_valid); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::uses_ready); + static_assert(P::get_property() == + sycl::ext::intel::experimental::in_csr); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::protocol_avalon); +} diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 39b0921629336..43ebbccf2eac5 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -40,6 +40,7 @@ add_subdirectory(scheduler) add_subdirectory(stream) add_subdirectory(SYCL2020) add_subdirectory(thread_safety) +add_subdirectory(pipes) add_subdirectory(program_manager) add_subdirectory(assert) add_subdirectory(Extensions) diff --git a/sycl/unittests/pipes/CMakeLists.txt b/sycl/unittests/pipes/CMakeLists.txt new file mode 100644 index 0000000000000..9bec8a94609c1 --- /dev/null +++ b/sycl/unittests/pipes/CMakeLists.txt @@ -0,0 +1,8 @@ +set(CMAKE_CXX_EXTENSIONS OFF) + +add_sycl_unittest(PipeTests OBJECT + host_pipe_registration.cpp +) + +add_dependencies(PipeTests sycl) +target_include_directories(PipeTests PRIVATE SYSTEM ${sycl_inc_dir}) \ No newline at end of file diff --git a/sycl/unittests/pipes/host_pipe_registration.cpp b/sycl/unittests/pipes/host_pipe_registration.cpp new file mode 100644 index 0000000000000..39fbe4931df4c --- /dev/null +++ b/sycl/unittests/pipes/host_pipe_registration.cpp @@ -0,0 +1,154 @@ +//==-------------- host_pipe_registration.cpp - Host pipe tests------------==// +// +// 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 +#include + +#include +#include +#include +#include + +namespace { +using namespace cl::sycl; +using pipe_prop = decltype(ext::oneapi::experimental::properties( + ext::intel::experimental::min_capacity<5>)); + +template struct pipe_id { + static constexpr unsigned id = ID; +}; + +class test_data_type { +public: + int num; +}; + +using test_host_pipe = + ext::intel::experimental::host_pipe, test_data_type, pipe_prop>; + +pi_device_binary_struct generate_device_binary() { + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + unittest::PiArray Entries = + unittest::makeEmptyKernels({"TestKernel"}); + unittest::PiPropertySet PropSet; + pi_device_binary_struct MBinaryDesc = pi_device_binary_struct{ + PI_DEVICE_BINARY_VERSION, + PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL, + PI_DEVICE_BINARY_TYPE_SPIRV, + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, + "", + "", + nullptr, + nullptr, + &*Bin.begin(), + (&*Bin.begin()) + Bin.size(), + Entries.begin(), + Entries.end(), + PropSet.begin(), + PropSet.end(), + }; + return MBinaryDesc; +} +pi_event READ = reinterpret_cast(0); +pi_event WRITE = reinterpret_cast(1); +static constexpr test_data_type PipeReadVal = {8}; +static test_data_type PipeWriteVal = {0}; +pi_result redefinedEnqueueReadHostPipe(pi_queue, pi_program, const char *, + pi_bool, void *ptr, size_t, pi_uint32, + const pi_event *, pi_event *event) { + *(((test_data_type *)ptr)) = PipeReadVal; + *event = READ; + return PI_SUCCESS; +} +pi_result redefinedEnqueueWriteHostPipe(pi_queue, pi_program, const char *, + pi_bool, void *ptr, size_t, pi_uint32, + const pi_event *, pi_event *event) { + test_data_type tmp = {9}; + PipeWriteVal = tmp; + *event = WRITE; + return PI_SUCCESS; +} + +bool preparePiMock(platform &Plt) { + if (Plt.is_host()) { + std::cout << "Not run on host - no PI events created in that case" + << std::endl; + return false; + } + + unittest::PiMock Mock{Plt}; + Mock.redefine( + redefinedEnqueueReadHostPipe); + Mock.redefine( + redefinedEnqueueWriteHostPipe); + return true; +} + +class PipeTest : public ::testing::Test { +protected: + void SetUp() override { + platform Plt{default_selector()}; + if (!preparePiMock(Plt)) + return; + context Ctx{Plt.get_devices()[0]}; + queue Q{Ctx, default_selector()}; + plat = Plt; + ctx = Ctx; + q = Q; + + // Fake registration of host pipes + sycl::detail::host_pipe_map::add(test_host_pipe::get_host_ptr(), + "test_host_pipe_unique_id"); + // Fake registration of device image + static constexpr size_t NumberOfImages = 1; + pi_device_binary_struct MNativeImages[NumberOfImages]; + MNativeImages[0] = generate_device_binary(); + MAllBinaries = pi_device_binaries_struct{ + PI_DEVICE_BINARIES_VERSION, + NumberOfImages, + MNativeImages, + nullptr, // not used, put here for compatibility with OpenMP + nullptr, // not used, put here for compatibility with OpenMP + }; + __sycl_register_lib(&MAllBinaries); + } + + void TearDown() override { __sycl_unregister_lib(&MAllBinaries); } + + platform plat; + context ctx; + queue q; + pi_device_binaries_struct MAllBinaries; +}; + +TEST_F(PipeTest, Basic) { + const void *HostPipePtr = test_host_pipe::get_host_ptr(); + detail::HostPipeMapEntry *hostPipeEntry = + detail::ProgramManager::getInstance().getHostPipeEntry(HostPipePtr); + const std::string pipe_name = hostPipeEntry->MUniqueId; + test_data_type host_pipe_read_data = {}; + void *data_ptr = &host_pipe_read_data; + event e = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, data_ptr, sizeof(test_data_type), false, + true /* read */); + }); + e.wait(); + // auto host_pipe_read_data = test_host_pipe::read(q); + assert(host_pipe_read_data.num == PipeReadVal.num); + test_data_type tmp = {9}; + data_ptr = &tmp; + event e_write = q.submit([=](handler &CGH) { + CGH.read_write_host_pipe(pipe_name, data_ptr, sizeof(test_data_type), false, + false /* write */); + }); + e_write.wait(); + // test_host_pipe::write(q, tmp); + assert(PipeWriteVal.num == 9); +} + +} // namespace