Implement host pipe unique name generation and mapping calls

rho180 · rho180 · commit 823985f19ebc · 2023-02-02T11:16:44.000-08:00
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
@@ -1554,6 +1554,15 @@ def SYCLDeviceGlobal: InheritableAttr {
   let SimpleHandler = 1;
 }
 
+def SYCLHostPipe: InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "host_pipe">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  // Only used internally by SYCL implementation
+  let Documentation = [SYCLHostPipeAttrDocs];
+  let SimpleHandler = 1;
+}
+
 def SYCLGlobalVariableAllowed : InheritableAttr {
   let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">];
   let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
@@ -3124,6 +3124,26 @@ so we have this attribute in sycl_detail namespace.
   }];
 }
 
+def SYCLHostPipeAttrDocs : Documentation {
+  let Category = DocCatType;
+  let Heading = "__sycl_detail__::host_pipe";
+  let Content = [{
+This attribute is part of support for SYCL host_pipe feature.
+Global or static variables of type decorated with this attribute have
+`sycl-unique-id`, an LLVM IR attribute, added to the definition of each such
+variable, which provides a unique string identifier using
+__builtin_sycl_unique_stable_id.
+We do not intend to support this as a general attribute that user code can use,
+so we have this attribute in sycl_detail namespace.
+
+.. code-block:: c++
+  struct
+  [[__sycl_detail__::host_pipe]] __pipeType {}
+
+  __pipeType __pipe;
+  }];
+}
+
 def SYCLGlobalVariableAllowedAttrDocs : Documentation {
   let Category = DocCatType;
   let Heading = "__sycl_detail__::global_variable_allowed";
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
@@ -371,6 +371,13 @@ class SYCLIntegrationHeader {
     NeedToEmitDeviceGlobalRegistration = true;
   }
 
+  /// Signals that emission of __sycl_host_pipe_registration type and
+  /// declaration of variable __sycl_host_pipe_registrar of this type in
+  /// integration header is required.
+  void addHostPipeRegistration() {
+    NeedToEmitHostPipeRegistration = true;
+  }
+
 private:
   // Kernel actual parameter descriptor.
   struct KernelParamDesc {
@@ -454,6 +461,10 @@ class SYCLIntegrationHeader {
   /// Keeps track of whether declaration of __sycl_device_global_registration
   /// type and __sycl_device_global_registrar variable are required to emit.
   bool NeedToEmitDeviceGlobalRegistration = false;
+
+  /// Keeps track of whether declaration of __sycl_host_pipe_registration
+  /// type and __sycl_host_pipe_registrar variable are required to emit.
+  bool NeedToEmitHostPipeRegistration = false;
 };
 
 class SYCLIntegrationFooter {
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5512,9 +5512,10 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
     // type.
     if (RD && RD->hasAttr<SYCLAddIRAttributesGlobalVariableAttr>())
       AddGlobalSYCLIRAttributes(GV, RD);
-    // If VarDecl has a type decorated with SYCL device_global attribute, emit
-    // IR attribute 'sycl-unique-id'.
-    if (RD && RD->hasAttr<SYCLDeviceGlobalAttr>())
+    // If VarDecl has a type decorated with SYCL device_global attribute or
+    // SYCL host_pipe attribute, emit IR attribute 'sycl-unique-id'.
+    if (RD && (RD->hasAttr<SYCLDeviceGlobalAttr>() || 
+               RD->hasAttr<SYCLHostPipeAttr>()))
       addSYCLUniqueID(GV, D, Context);
   }
 
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -9702,6 +9702,9 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
   // Process device-globals.
   addArgs(CmdArgs, TCArgs, {"-device-globals"});
 
+  // Process host pipes.
+  addArgs(CmdArgs, TCArgs, {"-host-pipes"});
+
   // Make ESIMD accessors use stateless memory accesses.
   if (TCArgs.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem,
                      options::OPT_fno_sycl_esimd_force_stateless_mem, false))
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
@@ -5169,6 +5169,24 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "\n";
   }
 
+  // Generate declaration of variable of type __sycl_host_pipe_registration
+  // whose sole purpose is to run its constructor before the application's
+  // main() function.
+  if (NeedToEmitHostPipeRegistration) {
+    O << "namespace {\n";
+
+    O << "class __sycl_host_pipe_registration {\n";
+    O << "public:\n";
+    O << "  __sycl_host_pipe_registration() noexcept;\n";
+    O << "};\n";
+    O << "__sycl_host_pipe_registration __sycl_host_pipe_registrar;\n";
+
+    O << "} // namespace\n";
+
+    O << "\n";
+  }
+
+
   O << "// names of all kernels defined in the corresponding source\n";
   O << "static constexpr\n";
   O << "const char* const kernel_names[] = {\n";
@@ -5359,6 +5377,7 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
     return;
   // Step 1: ensure that this is of the correct type template specialization.
   if (!isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) &&
+      !S.isTypeDecoratedWithDeclAttribute<SYCLHostPipeAttr>(VD->getType()) &&
       !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
           VD->getType())) {
     // Handle the case where this could be a deduced type, such as a deduction
@@ -5528,20 +5547,25 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   llvm::SmallSet<const VarDecl *, 8> Visited;
   bool EmittedFirstSpecConstant = false;
   bool DeviceGlobalsEmitted = false;
+  bool HostPipesEmitted = false;
 
   // Used to uniquely name the 'shim's as we generate the names in each
   // anonymous namespace.
   unsigned ShimCounter = 0;
 
   std::string DeviceGlobalsBuf;
   llvm::raw_string_ostream DeviceGlobOS(DeviceGlobalsBuf);
+  std::string HostPipesBuf;
+  llvm::raw_string_ostream HostPipesOS(HostPipesBuf);
   for (const VarDecl *VD : GlobalVars) {
     VD = VD->getCanonicalDecl();
 
-    // Skip if this isn't a SpecIdType or DeviceGlobal.  This can happen if it
-    // was a deduced type.
+    // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe.  This 
+    // can happen if it was a deduced type.
     if (!isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) &&
         !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            VD->getType()) &&
+        !S.isTypeDecoratedWithDeclAttribute<SYCLHostPipeAttr>(
             VD->getType()))
       continue;
 
@@ -5551,7 +5575,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
     // We only want to emit the #includes if we have a variable that needs
     // them, so emit this one on the first time through the loop.
-    if (!EmittedFirstSpecConstant && !DeviceGlobalsEmitted)
+    if (!EmittedFirstSpecConstant && !DeviceGlobalsEmitted && !HostPipesEmitted)
       OS << "#include <sycl/detail/defines_elementary.hpp>\n";
 
     Visited.insert(VD);
@@ -5571,6 +5595,21 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
                                                           VD);
       DeviceGlobOS << "\");\n";
+    } else if (S.isTypeDecoratedWithDeclAttribute<SYCLHostPipeAttr>(
+            VD->getType())) {
+      HostPipesEmitted = true;
+      HostPipesOS << "host_pipe_map::add(";
+      HostPipesOS << "(void *)&";
+      if (VD->isInAnonymousNamespace()) {
+        HostPipesOS << TopShim;
+      } else {
+        HostPipesOS << "::";
+        VD->getNameForDiagnostic(HostPipesOS, Policy, true);
+      }
+      HostPipesOS << ", \"";
+      HostPipesOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
+                                                         VD);
+      HostPipesOS << "\");\n";
     } else {
       EmittedFirstSpecConstant = true;
       OS << "namespace sycl {\n";
@@ -5614,5 +5653,21 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
     S.getSyclIntegrationHeader().addDeviceGlobalRegistration();
   }
+
+  if (HostPipesEmitted) {
+    OS << "#include <sycl/detail/host_pipe_map.hpp>\n";
+    HostPipesOS.flush();
+    OS << "namespace sycl::detail {\n";
+    OS << "namespace {\n";
+    OS << "__sycl_host_pipe_registration::__sycl_host_pipe_"
+          "registration() noexcept {\n";
+    OS << HostPipesBuf;
+    OS << "}\n";
+    OS << "} // namespace (unnamed)\n";
+    OS << "} // namespace sycl::detail\n";
+
+    S.getSyclIntegrationHeader().addHostPipeRegistration();
+  }
+
   return true;
 }
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -266,6 +266,13 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
                                               HostAccessDecorValue, VarName));
     }
 
+    if (isHostPipeVariable(GV)) {
+      auto VarName = getGlobalVariableUniqueId(GV);
+      MDOps.push_back(buildSpirvDecorMetadata(Ctx, SPIRV_HOST_ACCESS_DECOR,
+                                              SPIRV_HOST_ACCESS_DEFAULT_VALUE, 
+                                              VarName));
+    }
+
     // Add the generated metadata to the variable
     if (!MDOps.empty()) {
       GV.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps));
diff --git a/llvm/test/tools/sycl-post-link/host-pipes/basic.ll b/llvm/test/tools/sycl-post-link/host-pipes/basic.ll
@@ -0,0 +1,23 @@
+; RUN: sycl-post-link --host-pipes -S %s -o %t.files.table
+; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
+; RUN: sycl-post-link --host-pipes --ir-output-only %s -S -o - | FileCheck %s --check-prefix CHECK-IR
+
+; This test is intended to check that CompileTimePropertiesPass adds all the required
+; metadata nodes to host pipe vars decorated with the "sycl-host-pipe" attribute
+
+source_filename = "basic.cpp"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64_fpga-unknown-unknown"
+
+%struct.BasicKernel = type { i8 }
+
+$_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = comdat any
+
+@_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" }
+
+; Ensure that the generated metadata nodes are correct
+; CHECK-IR-DAG: ![[#MN0]] = !{![[#MN1:]]}
+; CHECK-IR-DAG: ![[#MN1]] = !{i32 6147, i32 2, !"_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE"}
diff --git a/llvm/tools/sycl-post-link/HostPipes.cpp b/llvm/tools/sycl-post-link/HostPipes.cpp
@@ -0,0 +1,42 @@
+//===------------- 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 "HostPipes.h"
+#include "CompileTimePropertiesPass.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";
+
+} // anonymous namespace
+
+namespace llvm {
+
+/// Return \c true if the variable @GV is a device global 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);
+}
+
+} // namespace llvm
diff --git a/llvm/tools/sycl-post-link/HostPipes.h b/llvm/tools/sycl-post-link/HostPipes.h
@@ -0,0 +1,34 @@
+//===------- HostPipes.h - get required into about SYCL 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 host pipe 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;
+
+/// Return \c true if the variable @GV is a host pipe variable.
+///
+/// @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);
+
+} // end namespace llvm
diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -207,6 +207,11 @@ cl::opt<bool> DeviceGlobals{
     cl::desc("Lower and generate information about device global variables"),
     cl::cat(PostLinkCat)};
 
+cl::opt<bool> HostPipes{
+    "host-pipes",
+    cl::desc("Lower and generate information about host pipe variables"),
+    cl::cat(PostLinkCat)};
+
 struct GlobalBinImageProps {
   bool EmitKernelParamInfo;
   bool EmitProgramMetadata;
@@ -970,10 +975,11 @@ int main(int argc, char **argv) {
   bool DoProgMetadata = EmitProgramMetadata.getNumOccurrences() > 0;
   bool DoExportedSyms = EmitExportedSymbols.getNumOccurrences() > 0;
   bool DoDeviceGlobals = DeviceGlobals.getNumOccurrences() > 0;
+  bool DoHostPipes = HostPipes.getNumOccurrences() > 0;
 
   if (!DoSplit && !DoSpecConst && !DoSymGen && !DoParamInfo &&
       !DoProgMetadata && !DoSplitEsimd && !DoExportedSyms && !DoDeviceGlobals &&
-      !DoLowerEsimd) {
+      !DoLowerEsimd && !DoHostPipes) {
     errs() << "no actions specified; try --help for usage info\n";
     return 1;
   }
diff --git a/sycl/include/sycl/detail/host_pipe_map.hpp b/sycl/include/sycl/detail/host_pipe_map.hpp
@@ -0,0 +1,21 @@
+//==-------------------- host_pipe_map.hpp -----------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+namespace host_pipe_map {
+
+__SYCL_EXPORT void add(const void *HostPipePtr, const char *UniqueId);
+
+} // namespace host_pipe_map
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
