Add pi extension API for host pipes

Setup lower runtime extension functions for host pipes.

See also
#5766
#5851

Host pipe sycl spec: #5838

Author: sherry-yuan

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

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

@@ -1784,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/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)

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;

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)

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -4123,12 +4123,12 @@ pi_result piProgramGetInfo(pi_program Program, pi_program_info ParamName,
         uint32_t Count = 0;
         ZE_CALL(zeModuleGetKernelNames, (Program->ZeModule, &Count, nullptr));
         std::unique_ptr<const char *[]> PNames(new const char *[Count]);
-        ZE_CALL(zeModuleGetKernelNames,
-                (Program->ZeModule, &Count, PNames.get()));
-        for (uint32_t I = 0; I < Count; ++I) {
-          PINames += (I > 0 ? ";" : "");
-          PINames += PNames[I];
-        }
+            ZE_CALL(zeModuleGetKernelNames,
+                    (Program->ZeModule, &Count, PNames.get()));
+            for (uint32_t I = 0; I < Count; ++I) {
+              PINames += (I > 0 ? ";" : "");
+              PINames += PNames[I];
+            }
       } else {
         return PI_INVALID_PROGRAM;
       }
@@ -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;