[SYCL] Host pipe runtime implementation

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)