[SYCL] Host pipe runtime implementation (#7468)

Disclaimer:

This work is a continuation of a previous approved Sherry's PR:
#5766 and her draft work
#5894.

We are implementing Hostpipes based on the Spec
[here](https://github.com/rho180/llvm/blob/rho180-pipe-design/sycl/doc/design/HostPipes.md).
OpenCL spec is
[here](https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc).

The following is the outline of the design.

1. The host pipe properties need to be added to the device image
(probably similar to the previous change here
a9ad3af)

2.The frontend calls the [registration of device image, maybe similar to
this
code](https://github.com/intel/llvm/blob/af858c74e4be95b2163ce6ba545ce588ff0aca4c/sycl/source/detail/program_manager/program_manager.cpp#L2014):
which is where the host pipe information is available. This is where the
mapping from host pipe name to host pipe pointer is extracted.

3. The frontend also calls `host_pipe_map::add` to register/initialize
the host pipes. This is the new function. We are not sure about the
ordering of registration of device image / registration of host pipe.
for which ever one that comes later, it need to initialize the remaining
attribute of host pipe (such as its properties).

4. The opencl runtime will need to get a cl_program object, which is
typically not available until the first kernel launch. To get a program
object early on, the host pipe name/pointer to device image mapping is
cached during registration. And when the specific host pipe is needed,
build the program and get its ocl runtime representation. This is done
in the first couple commits.

5. Since a host pipe read/write need to depend on other write operation
finishing before it (including the inter kernel write). This means the
pipe needs to know the dependency of kernel execution. For this reason,
the host pipe read and write ocl function cannot be called with no dep
event. therefore, it is implemented with handler , which is aware of the
event that it is supposed to wait upon. This is done in the "Register
new command group .." commit.

6. Unit test
    - mock a fake device image.
    - register the fake device image
    - register fake pipe with some name you specified
- fake the opencl functionality, this can be done with
unittest::PiMock::redefine

---------

Co-authored-by: Sherry Yuan <[email protected]>
Co-authored-by: Ho, Robert <[email protected]>

Author: 3 people

buildbot/dependency.py

@@ -58,7 +58,7 @@ def do_dependency(args):
 
     # 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", "9ddb236e6eb3cf844f9e2f81677e1045f9bf838e"]
     subprocess.check_call(checkout_cmd, cwd=ocl_header_dir)
 
     # fetch and build OpenCL ICD loader

llvm/include/llvm/SYCLLowerIR/HostPipes.h

@@ -0,0 +1,59 @@
+//===------- HostPipes.h - get required info about FPGA Host Pipes --------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// The file contains a number of functions to extract corresponding attributes
+// of the host pipe global variables and save them as a property set for the
+// runtime.
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/ADT/MapVector.h"
+
+#include <cstdint>
+#include <vector>
+
+namespace llvm {
+
+class GlobalVariable;
+class Module;
+class StringRef;
+
+// Represents a host pipe variable - at SYCL RT level host pipe
+// variables are being represented as a byte-array.
+struct HostPipeProperty {
+  HostPipeProperty(uint32_t Size) : Size(Size) {}
+
+  // Encodes size of the underlying type T of the host pipe variable.
+  uint32_t Size;
+};
+
+using HostPipePropertyMapTy =
+    MapVector<StringRef, std::vector<HostPipeProperty>>;
+
+/// Return \c true if the variable @GV is a host pipe variable.
+///
+/// The function checks whether the variable has the LLVM IR attribute \c
+/// sycl-host-pipe
+/// @param GV [in] A variable to test.
+///
+/// @return \c true if the variable is a host pipe variable, \c false
+/// otherwise.
+bool isHostPipeVariable(const GlobalVariable &GV);
+
+/// Searches given module for occurrences of host pipe variable-specific
+/// metadata and builds "host pipe variable name" ->
+/// vector<"variable properties"> map.
+///
+/// @param M [in] LLVM Module.
+///
+/// @returns the "host pipe variable name" -> vector<"variable properties">
+/// map.
+HostPipePropertyMapTy collectHostPipeProperties(const Module &M);
+
+} // end namespace llvm

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

@@ -22,7 +22,6 @@ namespace llvm {
 namespace sycl {
 namespace utils {
 constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id";
-constexpr StringRef SYCL_HOST_PIPE_ATTR = "sycl-host-pipe";
 
 using CallGraphNodeAction = ::std::function<void(Function *)>;
 using CallGraphFunctionFilter =
@@ -117,10 +116,6 @@ inline bool isSYCLExternalFunction(const Function *F) {
   return F->hasFnAttribute(ATTR_SYCL_MODULE_ID);
 }
 
-inline bool isHostPipeVariable(const GlobalVariable &GV) {
-  return GV.hasAttribute(SYCL_HOST_PIPE_ATTR);
-}
-
 } // namespace utils
 } // namespace sycl
 } // namespace llvm

llvm/include/llvm/Support/PropertySetIO.h

@@ -197,6 +197,7 @@ class PropertySetRegistry {
   static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols";
   static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals";
   static constexpr char SYCL_DEVICE_REQUIREMENTS[] = "SYCL/device requirements";
+  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/SYCLLowerIR/CMakeLists.txt

@@ -57,6 +57,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/LowerESIMDVecArg.cpp
   ESIMD/LowerESIMDVLoadVStore.cpp
   ESIMD/LowerESIMDSlmReservation.cpp
+  HostPipes.cpp
   LowerInvokeSimd.cpp
   LowerKernelProps.cpp
   LowerWGLocalMemory.cpp

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -10,7 +10,7 @@
 
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
-#include "llvm/SYCLLowerIR/SYCLUtils.h"
+#include "llvm/SYCLLowerIR/HostPipes.h"
 
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/StringMap.h"
@@ -343,7 +343,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
                                               HostAccessDecorValue, VarName));
     }
 
-    if (sycl::utils::isHostPipeVariable(GV)) {
+    if (isHostPipeVariable(GV)) {
       auto VarName = getGlobalVariableUniqueId(GV);
       MDOps.push_back(buildSpirvDecorMetadata(Ctx, SPIRV_HOST_ACCESS_DECOR,
                                               SPIRV_HOST_ACCESS_DEFAULT_VALUE, 

llvm/lib/SYCLLowerIR/DeviceGlobals.cpp

@@ -72,19 +72,19 @@ bool hasDeviceImageScopeProperty(const GlobalVariable &GV) {
   return hasProperty(GV, SYCL_DEVICE_IMAGE_SCOPE_ATTR);
 }
 
-/// Returns the unique id for the device global variable.
+/// Returns the unique id for the device global or host pipe variable.
 ///
 /// The function gets this value from the LLVM IR attribute \c
 /// sycl-unique-id.
 ///
-/// @param GV [in] Device Global variable.
+/// @param GV [in] Device Global or Hostpipe variable.
 ///
-/// @returns the unique id of the device global variable represented
-/// in the LLVM IR by \c GV.
+/// @returns the unique id of the device global or hostpipe variable
+/// represented in the LLVM IR by \c GV.
 StringRef getGlobalVariableUniqueId(const GlobalVariable &GV) {
   assert(GV.hasAttribute(SYCL_UNIQUE_ID_ATTR) &&
          "a 'sycl-unique-id' string must be associated with every device "
-         "global variable");
+         "global or hostpipe variable");
   return GV.getAttribute(SYCL_UNIQUE_ID_ATTR).getValueAsString();
 }
 

llvm/lib/SYCLLowerIR/HostPipes.cpp

@@ -0,0 +1,80 @@
+//===------------- HostPipes.cpp - SYCL Host Pipes Pass -------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// See comments in the header.
+//===----------------------------------------------------------------------===//
+
+#include "llvm/SYCLLowerIR/HostPipes.h"
+#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
+#include "llvm/SYCLLowerIR/DeviceGlobals.h"
+
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/IR/Module.h"
+
+#include <cassert>
+
+using namespace llvm;
+
+namespace {
+
+constexpr StringRef SYCL_HOST_PIPE_ATTR = "sycl-host-pipe";
+constexpr StringRef SYCL_HOST_PIPE_SIZE_ATTR = "sycl-host-pipe-size";
+
+/// Returns the size (in bytes) of the type \c T of the host
+/// pipe variable.
+///
+/// The function gets this value from the LLVM IR attribute \c
+/// sycl-host-pipe-size.
+///
+/// @param GV [in] Host Pipe variable.
+///
+/// @returns the size (int bytes) of the underlying type \c T of the
+/// host pipe variable represented in the LLVM IR by  @GV.
+uint32_t getHostPipeTypeSize(const GlobalVariable &GV) {
+  assert(GV.hasAttribute(SYCL_HOST_PIPE_SIZE_ATTR) &&
+         "The host pipe variable must have the 'sycl-host-pipe-size' "
+         "attribute that must contain a number representing the size of the "
+         "underlying type T of the host pipe variable");
+  return getAttributeAsInteger<uint32_t>(GV, SYCL_HOST_PIPE_SIZE_ATTR);
+}
+
+} // anonymous namespace
+
+namespace llvm {
+
+/// Return \c true if the variable @GV is a host pipe variable.
+///
+/// The function checks whether the variable has the LLVM IR attribute \c
+/// sycl-host-pipe.
+/// @param GV [in] A variable to test.
+///
+/// @return \c true if the variable is a host pipe variable, \c false
+/// otherwise.
+bool isHostPipeVariable(const GlobalVariable &GV) {
+  return GV.hasAttribute(SYCL_HOST_PIPE_ATTR);
+}
+
+HostPipePropertyMapTy collectHostPipeProperties(const Module &M) {
+  HostPipePropertyMapTy HPM;
+  auto HostPipeNum = count_if(M.globals(), isHostPipeVariable);
+  if (HostPipeNum == 0)
+    return HPM;
+
+  HPM.reserve(HostPipeNum);
+
+  for (auto &GV : M.globals()) {
+    if (!isHostPipeVariable(GV))
+      continue;
+
+    HPM[getGlobalVariableUniqueId(GV)] = {getHostPipeTypeSize(GV)};
+  }
+
+  return HPM;
+}
+
+} // namespace llvm

llvm/lib/Support/PropertySetIO.cpp

@@ -203,6 +203,7 @@ constexpr char PropertySetRegistry::SYCL_ASSERT_USED[];
 constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[];
 constexpr char PropertySetRegistry::SYCL_DEVICE_GLOBALS[];
 constexpr char PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS[];
+constexpr char PropertySetRegistry::SYCL_HOST_PIPES[];
 
 } // namespace util
 } // namespace llvm

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll

@@ -14,7 +14,7 @@ $_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experime
 @_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1 #0
 ; CHECK-IR: @_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1, !spirv.Decorations ![[#MN0:]]
 
-attributes #0 = { "sycl-host-pipe" "sycl-unique-id"="_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE" }
+attributes #0 = { "sycl-host-pipe" "sycl-host-pipe-size"="4" "sycl-unique-id"="_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE" }
 
 ; Ensure that the generated metadata nodes are correct
 ; CHECK-IR-DAG: ![[#MN0]] = !{![[#MN1:]]}

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -37,6 +37,7 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
+#include "llvm/SYCLLowerIR/HostPipes.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
 #include "llvm/SYCLLowerIR/LowerKernelProps.h"
 #include "llvm/Support/CommandLine.h"
@@ -466,6 +467,11 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
       PropSet.add(PropSetRegTy::SYCL_DEVICE_GLOBALS, DevGlobalPropertyMap);
   }
 
+  auto HostPipePropertyMap = collectHostPipeProperties(M);
+  if (!HostPipePropertyMap.empty()) {
+    PropSet.add(PropSetRegTy::SYCL_HOST_PIPES, HostPipePropertyMap);
+  }
+
   std::error_code EC;
   std::string SCFile = makeResultFileName(".prop", I, Suff);
   raw_fd_ostream SCOut(SCFile, EC);

opencl/CMakeLists.txt

@@ -20,7 +20,7 @@ set(OCL_LOADER_REPO
 
 # Repo tags/hashes
 
-set(OCL_HEADERS_TAG dcd5bede6859d26833cd85f0d6bbcee7382dc9b3)
+set(OCL_HEADERS_TAG 9ddb236e6eb3cf844f9e2f81677e1045f9bf838e)
 set(OCL_LOADER_TAG 9a3e962f16f5097d2054233ad8b6dad51b6f41b7)
 
 # OpenCL Headers

sycl/include/sycl/detail/cg.hpp

@@ -74,6 +74,7 @@ class CG {
     Memset2DUSM = 18,
     CopyToDeviceGlobal = 19,
     CopyFromDeviceGlobal = 20,
+    ReadWriteHostPipe = 21,
   };
 
   CG(CGTYPE Type, std::vector<std::vector<char>> ArgsStorage,
@@ -495,6 +496,36 @@ class CGMemset2DUSM : public CG {
   char getValue() const { return MValue; }
 };
 
+/// "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<AccessorImplHost *> 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; }
+};
+
 /// "Copy to device_global" command group class.
 class CGCopyToDeviceGlobal : public CG {
   void *MSrc;

sycl/include/sycl/detail/pi.def

@@ -131,6 +131,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)