[SYCL] Implement host pipe unique name generation and mapping calls (#8009)

rho180 · AlexeySachkov · web-flow · commit 752e4d384e52 · 2023-02-06T18:21:04.000-08:00
Implementation of host pipes outlined in the design document in this PR: #5850 1. Generation a unique pipe id for GVs marked with the new "sycl-host-pipe" attribute. Id generation utilizes the same method as used for name generation for device global. 2. Added a host pipe map to map the addresses of marked GVs with the unique id. This host pipe map is generated by a constructor and method calls added to the header and footer. 3. Modified the sycl-post-link tool to generate compile time properties metadata for these GVs. This metadata contains the unique id generated for the GV to be consumed by the device backend compiler. PR for accompanying runtime changes: #7468 --------- Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
@@ -1331,12 +1331,12 @@ def SYCLType: InheritableAttr {
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
-                            "stream", "sampler"],
+                            "stream", "sampler", "host_pipe"],
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
-                            "stream", "sampler"]>];
+                            "stream", "sampler", "host_pipe"]>];
   // Only used internally by SYCL implementation
   let Documentation = [InternalOnly];
 }
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
@@ -5508,14 +5508,22 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
 
   if (getLangOpts().SYCLIsDevice) {
     const RecordDecl *RD = D->getType()->getAsRecordDecl();
-    // Add IR attributes if add_ir_attribute_global_variable is attached to
-    // 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>())
-      addSYCLUniqueID(GV, D, Context);
+
+    if (RD) {
+      // Add IR attributes if add_ir_attribute_global_variable is attached to
+      // type.
+      if (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->hasAttr<SYCLDeviceGlobalAttr>())
+        addSYCLUniqueID(GV, D, Context);
+      // If VarDecl type is SYCLTypeAttr::host_pipe, emit the IR attribute 
+      // 'sycl-unique-id'.
+      if (const auto *Attr = RD->getAttr<SYCLTypeAttr>())
+        if (Attr->getType() == SYCLTypeAttr::SYCLType::host_pipe)
+          addSYCLUniqueID(GV, D, Context);
+    }
   }
 
   if (D->getType().isRestrictQualified()) {
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) &&
+      !isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) &&
       !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
           VD->getType())) {
     // Handle the case where this could be a deduced type, such as a deduction
@@ -5528,19 +5547,23 @@ 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) &&
+        !isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && 
         !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
             VD->getType()))
       continue;
@@ -5551,7 +5574,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 +5594,20 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
                                                           VD);
       DeviceGlobOS << "\");\n";
+    } else if (isSyclType(VD->getType(), SYCLTypeAttr::host_pipe)) {
+      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 +5651,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/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -151,6 +151,32 @@ class [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allo
 } // namespace oneapi
 } // namespace ext
 
+namespace ext {
+namespace intel {
+namespace experimental {
+
+// host_pipe class decorated with attribute
+template <class _name, class _dataT>
+class
+host_pipe {
+
+public:
+  struct
+#ifdef __SYCL_DEVICE_ONLY__
+  [[__sycl_detail__::sycl_type(host_pipe)]]
+#endif
+  __pipeType { const char __p; };
+
+  static constexpr __pipeType __pipe = {0};
+  static _dataT read() {
+    (void)__pipe;
+  }
+};
+
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+
 template <int dim>
 struct id {
   template <typename... T>
diff --git a/clang/test/CodeGenSYCL/host_pipe.cpp b/clang/test/CodeGenSYCL/host_pipe.cpp
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-unique-prefix=THE_PREFIX -opaque-pointers -emit-llvm %s -o - | FileCheck %s
+#include "sycl.hpp"
+
+// Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the
+// global variable whose type is decorated with host_pipe attribute, and that a
+// unique string is generated.
+
+using namespace sycl::ext::intel::experimental;
+using namespace sycl;
+queue q;
+
+// check that "sycl-unique-id" attribute is created for host pipes
+// CHECK: @_ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE5HPIntiE6__pipeE = internal addrspace(1) constant %"struct.sycl::_V1::ext::intel::experimental::host_pipe<HPInt, int>::__pipeType" zeroinitializer, align 1 #[[HPINT_ATTRS:[0-9]+]]
+// CHECK: @_ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE7HPFloatiE6__pipeE = internal addrspace(1) constant %"struct.sycl::_V1::ext::intel::experimental::host_pipe<HPFloat, int>::__pipeType" zeroinitializer, align 1 #[[HPFLOAT_ATTRS:[0-9]+]]
+
+void foo() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name_1>([=]() {
+      host_pipe<class HPInt, int>::read();
+      host_pipe<class HPFloat, int>::read();
+    });
+  });
+}
+
+// CHECK: attributes #[[HPINT_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE5HPIntiE6__pipeE" }
+// CHECK: attributes #[[HPFLOAT_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE7HPFloatiE6__pipeE" 
+
diff --git a/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp b/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp
@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h -fsycl-unique-prefix=THE_PREFIX %s -emit-llvm -o %t.ll
+// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
+// RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
+#include "sycl.hpp"
+
+// Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the
+// global variable whose type is decorated with host_pipe attribute, and that a
+// unique string is generated.
+
+using namespace sycl::ext::intel::experimental;
+using namespace sycl;
+queue q;
+
+void foo() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name_1>([=]() {
+      host_pipe<class HPInt, int>::read();
+      host_pipe<class HPFloat, int>::read();
+    });
+  });
+}
+
+// CHECK-HEADER: namespace sycl {
+// CHECK-HEADER-NEXT: __SYCL_INLINE_VER_NAMESPACE(_V1) {
+// CHECK-HEADER-NEXT: namespace detail {
+// CHECK-HEADER-NEXT: namespace {
+// CHECK-HEADER-NEXT: class __sycl_host_pipe_registration {
+// CHECK-HEADER-NEXT: public:
+// CHECK-HEADER-NEXT: __sycl_host_pipe_registration() noexcept;
+// CHECK-HEADER-NEXT: };
+// CHECK-HEADER-NEXT: __sycl_host_pipe_registration __sycl_host_pipe_registrar;
+// CHECK-HEADER-NEXT: } // namespace
+// CHECK-HEADER: } // namespace detail
+// CHECK-HEADER: } // __SYCL_INLINE_VER_NAMESPACE(_V1)
+// CHECK-HEADER: } // namespace sycl
+
+// CHECK-FOOTER: #include <sycl/detail/defines_elementary.hpp>
+// CHECK-FOOTER: #include <sycl/detail/host_pipe_map.hpp>
+// CHECK-FOOTER-NEXT: namespace sycl::detail {
+// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER-NEXT: __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept {
+
+// CHECK-FOOTER: host_pipe_map::add((void *)&::sycl::ext::intel::experimental::host_pipe<HPInt, int>::__pipe, "THE_PREFIX____ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE5HPIntiE6__pipeE");
+// CHECK-FOOTER: host_pipe_map::add((void *)&::sycl::ext::intel::experimental::host_pipe<HPFloat, int>::__pipe, "THE_PREFIX____ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE7HPFloatiE6__pipeE");
+
+// CHECK-FOOTER: } // namespace (unnamed)
+// CHECK-FOOTER: } // namespace sycl::detail
+
diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
@@ -12,6 +12,7 @@
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/IR/Function.h"
+#include "llvm/IR/GlobalVariable.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Operator.h"
 
@@ -21,6 +22,7 @@ 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 =
@@ -115,6 +117,10 @@ 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
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -10,6 +10,7 @@
 
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/StringMap.h"
@@ -266,6 +267,13 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
                                               HostAccessDecorValue, VarName));
     }
 
+    if (sycl::utils::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/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll
@@ -0,0 +1,21 @@
+; RUN: opt -passes=compile-time-properties %s -S | 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/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
