[WIP][SYCL][FPGA] Runtime implementation for host_pipes

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<std::vector<char>> 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<std::vector<char>> ArgsStorage,
+                      std::vector<detail::AccessorImplPtr> AccStorage,
+                      std::vector<std::shared_ptr<const void>> SharedPtrStorage,
+                      std::vector<Requirement *> Requirements,
+                      std::vector<detail::EventImplPtr> 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<char> MPattern;

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

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.

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<detail::queue_impl> 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<detail::EventImplPtr> 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;
 

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 <CL/sycl/context.hpp>
+#include <CL/sycl/device.hpp>
+#include <CL/sycl/queue.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/ext/oneapi/properties/property_utils.hpp>
+#include <sycl/ext/oneapi/properties/property_value.hpp>
+
+#ifdef XPTI_ENABLE_INSTRUMENTATION
+#include <xpti/xpti_data_types.h>
+#include <xpti/xpti_trace_framework.hpp>
+#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 <int capacity>
+  using value_t =
+      property_value<min_capacity_key, std::integral_constant<int, capacity>>;
+};
+// min_capacity is an object of a property value type of min_capacity.
+template <int capacity>
+inline constexpr min_capacity_key::value_t<capacity> min_capacity;
+
+template <> struct is_property_key<min_capacity_key> : std::true_type {};
+
+namespace detail {
+
+template <> struct PropertyToKind<min_capacity_key> {
+  static constexpr PropKind Kind = PropKind::MinCapacity;
+};
+
+template <> struct IsCompileTimeProperty<min_capacity_key> : 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 _name, typename _dataT,
+          typename PropertyList = default_pipe_properties>
+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<PropertyList>,
+      "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)

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

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;

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;