[WIP][SYCL] Host pipe runtime implementation

buildbot/dependency.py

@@ -49,16 +49,16 @@ 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"]
         subprocess.check_call(fetch_cmd, cwd=ocl_header_dir)
 
     # 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

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).

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

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)

sycl/include/CL/sycl.hpp

@@ -71,3 +71,6 @@
 #include <sycl/ext/oneapi/reduction.hpp>
 #include <sycl/ext/oneapi/sub_group.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp>
+
+#include <sycl/ext/intel/experimental/host_pipes.hpp>
+#include <sycl/ext/intel/experimental/pipe_properties.hpp>

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/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)

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)

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.

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:

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,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 <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 {
+
+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 {
+
+  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<min_capacity_key>()
+          ? _propertiesT::template get_property<min_capacity_key>().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)