Register new command group to execute host pipe read/write operation

Defines the flow of enqueue new host pipe operations (read/write),
User provide the queue to enqueue this event, and the runtime queries
the pipe address from registration using the given address and unique ID.
The runtime pass the pipe name, and host address into queue submit of new command group.
The enqueued command calls new opencl function, and provide the current program, queue,
event wait list, pipe name, host pointer of the data destination.

Spec: #5838

Author: sherry-yuan

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,
@@ -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<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; }
+};
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

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

sycl/include/sycl/ext/intel/experimental/host_pipes.hpp

@@ -0,0 +1,87 @@
+//==---------------- 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/intel/experimental/pipe_properties.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <type_traits>
+
+#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 intel {
+namespace experimental {
+
+template <class _name, class _dataT,
+          class _propertiesT = decltype(oneapi::experimental::properties{}),
+          class = void>
+class host_pipe {
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_list_v<_propertiesT>,
+      "Host pipe is available only through new property list");
+};
+
+using default_pipe_properties =
+    decltype(sycl::ext::oneapi::experimental::properties(min_capacity<0>));
+
+template <class _name, class _dataT, class _propertiesT>
+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<_name, _dataT, _propertiesT,
+              std::enable_if_t<sycl::ext::oneapi::experimental::
+                                   is_property_list_v<_propertiesT>>> {
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_list_v<_propertiesT>,
+      "Host pipe is available only through new property list");
+
+public:
+  using value_type = _dataT;
+  static constexpr int32_t min_cap =
+      _propertiesT::template has_property<min_capacity_key>()
+          ? _propertiesT::template get_property<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);
+
+#ifdef __SYCL_DEVICE_ONLY__
+  static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
+                                                           min_capacity};
+#endif // __SYCL_DEVICE_ONLY__
+};
+
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/CMakeLists.txt

@@ -142,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"

sycl/source/detail/host_pipe.cpp

@@ -0,0 +1,69 @@
+//==-------------------- 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 <sycl/ext/intel/experimental/host_pipes.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+
+template <class _name, class _dataT, typename _propertiesT>
+_dataT
+host_pipe<_name, _dataT, _propertiesT,
+          std::enable_if_t<sycl::ext::oneapi::experimental::is_property_list_v<
+              _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 std::string pipe_name = "pipename";
+  size_t size = 4;
+  event e = q.submit([=](handler &CGH) {
+    CGH.read_write_host_pipe(pipe_name, (void *)(&data), (size_t)size, false,
+                             true /* read */);
+  });
+  e.wait();
+  return data;
+}
+
+template <class _name, class _dataT, typename _propertiesT>
+void host_pipe<
+    _name, _dataT, _propertiesT,
+    std::enable_if_t<sycl::ext::oneapi::experimental::is_property_list_v<
+        _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 std::string pipe_name = "pipename";
+  const void *data_ptr = &data;
+  size_t size = 4;
+  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)

sycl/source/detail/scheduler/commands.cpp

@@ -2192,6 +2192,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<RT::PiEvent> &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::OSModuleHandle M =
+      detail::OSUtil::getOSModuleHandle("HostPipeReadWriteKernelName");
+  RT::PiProgram Program =
+      sycl::detail::ProgramManager::getInstance().getBuiltPIProgram(
+          M, Queue->getContextImplPtr(), Queue->getDeviceImplPtr(),
+          "HostPipeReadWriteKernelName");
+
+  // 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<sycl::detail::PiApiKind::piextEnqueueReadHostPipe>(
+            pi_q, Program, PipeName.c_str(), blocking, ptr, size,
+            RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0],
+            OutEvent);
+  } else {
+    Error =
+        Plugin.call_nocheck<sycl::detail::PiApiKind::piextEnqueueWriteHostPipe>(
+            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 +2591,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);
   }

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<RT::PiEvent> &RawEvents,
+                                 RT::PiEvent *OutEvent, bool read);
+
 cl_int enqueueImpKernel(
     const QueueImplPtr &Queue, NDRDescT &NDRDesc, std::vector<ArgDesc> &Args,
     const std::shared_ptr<detail::kernel_bundle_impl> &KernelBundleImplPtr,

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<event> &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)

sycl/test/abi/sycl_symbols_linux.dump

@@ -4006,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

sycl/test/abi/symbol_size_alignment.cpp

@@ -58,11 +58,11 @@ int main() {
   check<event, 16, 8>();
   check<gpu_selector, 8, 8>();
 #ifdef _MSC_VER
-  check<handler, 552, 8>();
+  check<handler, 608, 8>();
   check<detail::buffer_impl, 216, 8>();
   check<detail::image_impl<1>, 272, 8>();
 #else
-  check<handler, 560, 8>();
+  check<handler, 616, 8>();
   check<detail::buffer_impl, 184, 8>();
   check<detail::image_impl<1>, 240, 8>();
 #endif