[SYCL] Cache host pipe to device image mapping

To build program before first kernel launch, the program manager requires knowing the device image, device and context associated with a host pipe.
It then will take from cache or build the program from device image.
This work around the difficulty of program being built on first kernel launch, and allows the host pipe read/write to happen at any point.

Author: sherry-yuan

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

@@ -52,16 +52,19 @@ class
 
   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
+      [[__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};  
+      __pipeType {
+    const char __p;
+  };
+
+  static constexpr __pipeType __pipe = {0};
 
 public:
   using value_type = _dataT;

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;
       }

sycl/source/detail/host_pipe.cpp

@@ -6,9 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <sycl/ext/intel/experimental/host_pipes.hpp>
 #include <detail/host_pipe_map_entry.hpp>
 #include <detail/program_manager/program_manager.hpp>
+#include <sycl/ext/intel/experimental/host_pipes.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -29,9 +29,10 @@ host_pipe<_name, _dataT, _propertiesT,
   }
   // 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;
+  const void *HostPipePtr = &__pipe;
+  detail::HostPipeMapEntry *hostPipeEntry =
+      detail::ProgramManager::getInstance().getHostPipeEntry(HostPipePtr);
+  const std::string pipe_name = hostPipeEntry->MUniqueId;
   size_t size = 4;
   event e = q.submit([=](handler &CGH) {
     CGH.read_write_host_pipe(pipe_name, (void *)(&data), (size_t)size, false,
@@ -54,9 +55,10 @@ void host_pipe<
     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 *HostPipePtr = &__pipe;
+  detail::HostPipeMapEntry *hostPipeEntry =
+      detail::ProgramManager::getInstance().getHostPipeEntry(HostPipePtr);
+  const std::string pipe_name = hostPipeEntry->MUniqueId;
   const void *data_ptr = &data;
   size_t size = 4;
   event e = q.submit([=](handler &CGH) {

sycl/source/detail/host_pipe_map_entry.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/detail/device_binary_image.hpp>
 #include <cstdint>
 #include <unordered_map>
 
@@ -21,6 +22,8 @@ struct HostPipeMapEntry {
   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
@@ -43,6 +46,9 @@ struct HostPipeMapEntry {
     assert(!MHostPipePtr && "Host pipe pointer has already been initialized.");
     MHostPipePtr = HostPipePtr;
   }
+  void initialize(const RTDeviceBinaryImage *DeviceImage) {
+    mDeviceImage = DeviceImage;
+  }
 };
 
 } // namespace detail

sycl/source/detail/program_manager/program_manager.cpp

@@ -1253,12 +1253,14 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) {
           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<HostPipeMapEntry>(HostPipe->Name, TypeSize);
+            EntryUPtr->initialize(Img.get());
             m_HostPipes.emplace(HostPipe->Name, std::move(EntryUPtr));
           }
         }

sycl/source/detail/scheduler/commands.cpp

@@ -18,6 +18,7 @@
 #include <CL/sycl/sampler.hpp>
 #include <detail/context_impl.hpp>
 #include <detail/event_impl.hpp>
+#include <detail/host_pipe_map_entry.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
 #include <detail/kernel_info.hpp>
@@ -2201,12 +2202,12 @@ cl_uint enqueueReadWriteHostPipe(const QueueImplPtr &Queue,
   // 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");
+  detail::HostPipeMapEntry *hostPipeEntry =
+      detail::ProgramManager::getInstance().getHostPipeEntry(PipeName);
   RT::PiProgram Program =
-      sycl::detail::ProgramManager::getInstance().getBuiltPIProgram(
-          M, Queue->getContextImplPtr(), Queue->getDeviceImplPtr(),
-          "HostPipeReadWriteKernelName");
+      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();

sycl/test/abi/layout_handler.cpp

@@ -1,3 +1,4 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s
 // RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s | FileCheck %s
 // REQUIRES: linux
 // UNSUPPORTED: libcxx