[SYCL] Set compiled program for hostpipe (#9060)

HostPipe implementation should pass the compiled program to backend when
invoke the read / write operation.

There are two scenarios for HostPipe read / write:
   * Invoke device kernel before HostPipe read / write
   * Invoke HostPipe read / write before device kernel.

For the 1st scenario, the program should be compiled before HostPipe
read / write. We get the program from Cache for HostPipe entry.

For the 2nd scenario, the program should NOT be compiled. We need to
compile the program first for HostPipe entry. The compiled program will
be cached for future device kernel.

Author: cdai2

sycl/source/detail/context_impl.cpp

@@ -394,26 +394,29 @@ void context_impl::DeviceGlobalInitializer::ClearEvents(const plugin &Plugin) {
   MDeviceGlobalInitEvents.clear();
 }
 
-std::optional<RT::PiProgram> context_impl::getProgramForDeviceGlobal(
-    const device &Device, DeviceGlobalMapEntry *DeviceGlobalEntry) {
+std::optional<RT::PiProgram> context_impl::getProgramForDevImgs(
+    const device &Device, const std::set<std::uintptr_t> &ImgIdentifiers,
+    const std::string &ObjectTypeName) {
+
   KernelProgramCache::ProgramWithBuildStateT *BuildRes = nullptr;
   {
     auto LockedCache = MKernelProgramCache.acquireCachedPrograms();
     auto &KeyMap = LockedCache.get().KeyMap;
     auto &Cache = LockedCache.get().Cache;
     RT::PiDevice &DevHandle = getSyclObjImpl(Device)->getHandleRef();
-    for (std::uintptr_t ImageIDs : DeviceGlobalEntry->MImageIdentifiers) {
+    for (std::uintptr_t ImageIDs : ImgIdentifiers) {
       auto OuterKey = std::make_pair(ImageIDs, DevHandle);
       size_t NProgs = KeyMap.count(OuterKey);
       if (NProgs == 0)
         continue;
       // If the cache has multiple programs for the identifiers or if we have
-      // already found a program in the cache with the device_global, we cannot
-      // proceed.
+      // already found a program in the cache with the device_global or host
+      // pipe we cannot proceed.
       if (NProgs > 1 || (BuildRes && NProgs == 1))
-        throw sycl::exception(
-            make_error_code(errc::invalid),
-            "More than one image exists with the device_global.");
+        throw sycl::exception(make_error_code(errc::invalid),
+                              "More than one image exists with the " +
+                                  ObjectTypeName + ".");
+
       auto KeyMappingsIt = KeyMap.find(OuterKey);
       assert(KeyMappingsIt != KeyMap.end());
       auto CachedProgIt = Cache.find(KeyMappingsIt->second);
@@ -426,6 +429,21 @@ std::optional<RT::PiProgram> context_impl::getProgramForDeviceGlobal(
   return *MKernelProgramCache.waitUntilBuilt<compile_program_error>(BuildRes);
 }
 
+std::optional<RT::PiProgram> context_impl::getProgramForDeviceGlobal(
+    const device &Device, DeviceGlobalMapEntry *DeviceGlobalEntry) {
+  return getProgramForDevImgs(Device, DeviceGlobalEntry->MImageIdentifiers,
+                              "device_global");
+}
+/// Gets a program associated with a HostPipe Entry from the cache.
+std::optional<RT::PiProgram>
+context_impl::getProgramForHostPipe(const device &Device,
+                                    HostPipeMapEntry *HostPipeEntry) {
+  // One HostPipe entry belongs to one Img
+  std::set<std::uintptr_t> ImgIdentifiers;
+  ImgIdentifiers.insert(HostPipeEntry->getDevBinImage()->getImageID());
+  return getProgramForDevImgs(Device, ImgIdentifiers, "host_pipe");
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/detail/context_impl.hpp

@@ -215,6 +215,15 @@ class context_impl {
   std::optional<RT::PiProgram>
   getProgramForDeviceGlobal(const device &Device,
                             DeviceGlobalMapEntry *DeviceGlobalEntry);
+  /// Gets a program associated with a HostPipe Entry from the cache.
+  std::optional<RT::PiProgram>
+  getProgramForHostPipe(const device &Device, HostPipeMapEntry *HostPipeEntry);
+
+  /// Gets a program associated with Dev / Images pairs.
+  std::optional<RT::PiProgram>
+  getProgramForDevImgs(const device &Device,
+                       const std::set<std::uintptr_t> &ImgIdentifiers,
+                       const std::string &ObjectTypeName);
 
   enum PropertySupport { NotSupported = 0, Supported = 1, NotChecked = 2 };
 

sycl/source/detail/host_pipe_map_entry.hpp

@@ -50,6 +50,10 @@ struct HostPipeMapEntry {
   void initialize(const RTDeviceBinaryImage *DeviceImage) {
     mDeviceImage = DeviceImage;
   }
+
+  RTDeviceBinaryImage *getDevBinImage() {
+    return const_cast<RTDeviceBinaryImage *>(mDeviceImage);
+  }
 };
 
 } // namespace detail

sycl/source/detail/scheduler/commands.cpp

@@ -2343,9 +2343,24 @@ pi_int32 enqueueReadWriteHostPipe(const QueueImplPtr &Queue,
   detail::HostPipeMapEntry *hostPipeEntry =
       ProgramManager::getInstance().getHostPipeEntry(PipeName);
 
-  RT::PiProgram Program = ProgramManager::getInstance().createPIProgram(
-      *(hostPipeEntry->mDeviceImage), Queue->get_context(),
-      Queue->get_device());
+  RT::PiProgram Program = nullptr;
+  device Device = Queue->get_device();
+  ContextImplPtr ContextImpl = Queue->getContextImplPtr();
+  std::optional<RT::PiProgram> CachedProgram =
+      ContextImpl->getProgramForHostPipe(Device, hostPipeEntry);
+  if (CachedProgram)
+    Program = *CachedProgram;
+  else {
+    // If there was no cached program, build one.
+    device_image_plain devImgPlain =
+        ProgramManager::getInstance().getDeviceImageFromBinaryImage(
+            hostPipeEntry->getDevBinImage(), Queue->get_context(),
+            Queue->get_device());
+    device_image_plain BuiltImage =
+        ProgramManager::getInstance().build(devImgPlain, {Device}, {});
+    Program = getSyclObjImpl(BuiltImage)->get_program_ref();
+  }
+  assert(Program && "Program for this hostpipe is not compiled.");
 
   // Get plugin for calling opencl functions
   const detail::plugin &Plugin = Queue->getPlugin();