From 823985f19ebc8888ecbe93d823f20cfbd062d0ed Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Fri, 13 Jan 2023 11:05:52 -0800
Subject: [PATCH 01/14] Implement host pipe unique name generation and mapping
 calls

---
 clang/include/clang/Basic/Attr.td             |  9 +++
 clang/include/clang/Basic/AttrDocs.td         | 20 ++++++
 clang/include/clang/Sema/Sema.h               | 11 ++++
 clang/lib/CodeGen/CodeGenModule.cpp           |  7 ++-
 clang/lib/Driver/ToolChains/Clang.cpp         |  3 +
 clang/lib/Sema/SemaSYCL.cpp                   | 61 ++++++++++++++++++-
 .../SYCLLowerIR/CompileTimePropertiesPass.cpp |  7 +++
 .../tools/sycl-post-link/host-pipes/basic.ll  | 23 +++++++
 llvm/tools/sycl-post-link/HostPipes.cpp       | 42 +++++++++++++
 llvm/tools/sycl-post-link/HostPipes.h         | 34 +++++++++++
 llvm/tools/sycl-post-link/sycl-post-link.cpp  |  8 ++-
 sycl/include/sycl/detail/host_pipe_map.hpp    | 21 +++++++
 12 files changed, 239 insertions(+), 7 deletions(-)
 create mode 100644 llvm/test/tools/sycl-post-link/host-pipes/basic.ll
 create mode 100644 llvm/tools/sycl-post-link/HostPipes.cpp
 create mode 100644 llvm/tools/sycl-post-link/HostPipes.h
 create mode 100644 sycl/include/sycl/detail/host_pipe_map.hpp

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 39dfa9bad949e..299ebefd4d1c2 100644
--- 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
index 657832824776e..12989451029dd 100644
--- 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
index 4dc888251959c..9a1a4b8850e0c 100644
--- 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
index 56c6e0a893b0d..e1c07d5be43e6 100644
--- 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
index ae467a22b500c..34ea69fdb4de8 100644
--- 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
index 326d3750542b3..66076697ced65 100644
--- 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,6 +5547,7 @@ 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.
@@ -5535,13 +5555,17 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
   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
index 1d2fa554ed8ee..a3063750a3d02 100644
--- 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
new file mode 100644
index 0000000000000..ca104c1fb24f9
--- /dev/null
+++ 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
new file mode 100644
index 0000000000000..591357e8984f0
--- /dev/null
+++ 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
new file mode 100644
index 0000000000000..cf3de06d5bac2
--- /dev/null
+++ 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
index f9110752bc331..b01e9f2cbdac7 100644
--- 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
new file mode 100644
index 0000000000000..5e2e222bd9fb5
--- /dev/null
+++ 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

From c905dcd7fc0bb0db9c0b40a66dfae82c3243f8ff Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Mon, 16 Jan 2023 10:41:28 -0800
Subject: [PATCH 02/14] Add front end lit tests for host pipes

---
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        | 28 +++++++++++
 clang/test/CodeGenSYCL/host_pipe.cpp          | 29 +++++++++++
 .../host_pipe_int_footer_header.cpp           | 48 +++++++++++++++++++
 3 files changed, 105 insertions(+)
 create mode 100644 clang/test/CodeGenSYCL/host_pipe.cpp
 create mode 100644 clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp

diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 0467659cd5492..3b653b4bfec12 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -151,6 +151,34 @@ 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__::add_ir_attributes_global_variable(
+    // "sycl-host-pipe", nullptr)]] [[__sycl_detail__::host_pipe]]
+  [[__sycl_detail__::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
new file mode 100644
index 0000000000000..c7c88ff2a5d62
--- /dev/null
+++ b/clang/test/CodeGenSYCL/host_pipe.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -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.
+
+// XFAIL:*
+
+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
new file mode 100644
index 0000000000000..cbf2c9f48dcff
--- /dev/null
+++ b/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp
@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -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
+

From e62716f2aeb03b3268510d0f7748d553a8465113 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Mon, 16 Jan 2023 10:42:35 -0800
Subject: [PATCH 03/14] Update driver lit tests to include sycl-post-link
 host-pipes argument

---
 clang/test/Driver/sycl-device-lib.cpp              |  2 +-
 clang/test/Driver/sycl-intelfpga-aoco.cpp          |  4 ++--
 clang/test/Driver/sycl-offload-intelfpga.cpp       | 14 +++++++-------
 ...pragma-attribute-supported-attributes-list.test |  1 +
 4 files changed, 11 insertions(+), 10 deletions(-)

diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp
index 4a2ad0fc8bcbc..e5012682a42ea 100644
--- a/clang/test/Driver/sycl-device-lib.cpp
+++ b/clang/test/Driver/sycl-device-lib.cpp
@@ -174,7 +174,7 @@
 // RUN:   | FileCheck %s -check-prefix=SYCL_LLVM_LINK_NO_DEVICE_LIB
 // SYCL_LLVM_LINK_NO_DEVICE_LIB: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device"
 // SYCL_LLVM_LINK_NO_DEVICE_LIB-NOT: llvm-link{{.*}}  "-only-needed"
-// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-emit-exported-symbols" "-split-esimd" "-lower-esimd" "-O2" "-spec-const=rt" "-device-globals" "-o" "{{.*}}.table" "{{.*}}.bc"
+// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-emit-exported-symbols" "-split-esimd" "-lower-esimd" "-O2" "-spec-const=rt" "-device-globals" "-host-pipes" "-o" "{{.*}}.table" "{{.*}}.bc"
 
 /// ###########################################################################
 /// test llvm-link behavior for special user input whose filename resembles SYCL device library
diff --git a/clang/test/Driver/sycl-intelfpga-aoco.cpp b/clang/test/Driver/sycl-intelfpga-aoco.cpp
index aefcd262edfb2..d1e82c835dcf1 100755
--- a/clang/test/Driver/sycl-intelfpga-aoco.cpp
+++ b/clang/test/Driver/sycl-intelfpga-aoco.cpp
@@ -52,7 +52,7 @@
 // RUN:  %clang_cl -fsycl -fno-sycl-instrument-device-code -fno-sycl-device-lib=all -fintelfpga -Xshardware %t_aoco_cl.a -### %s 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCO,CHK-FPGA-AOCO-WIN %s
 // CHK-FPGA-AOCO: llvm-link{{.*}} "-o" "[[LINKEDBC:.+\.bc]]"
-// CHK-FPGA-AOCO: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
+// CHK-FPGA-AOCO: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
 // CHK-FPGA-AOCO: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[SPLTABLE]]"
 // CHK-FPGA-AOCO: llvm-spirv{{.*}} "-o" "[[TARGSPV:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCO: clang-offload-bundler{{.*}} "-type=aoo" "-targets=sycl-fpga_aoco-intel-unknown" "-input=[[INPUTLIB:.+\.a]]" "-output=[[AOCOLIST:.+\.txt]]" "-unbundle"
@@ -103,7 +103,7 @@
 // CHK-FPGA-AOCO-EMU: clang-offload-bundler{{.*}} "-type=aoo" "-targets=sycl-spir64_fpga-unknown-unknown" "-input=[[INPUTLIB:.+\.a]]" "-output=[[OUTLIB:.+\.txt]]" "-unbundle"
 // CHK-FPGA-AOCO-EMU: llvm-foreach{{.*}} "--out-ext=txt" "--in-file-list=[[OUTLIB]]" "--in-replace=[[OUTLIB]]" "--out-file-list=[[DEVICELIST:.+\.txt]]" "--out-replace=[[DEVICELIST]]" "--" {{.*}}spirv-to-ir-wrapper{{.*}} "[[OUTLIB]]" "-o" "[[DEVICELIST]]"
 // CHK-FPGA-AOCO-EMU: llvm-link{{.*}} "@[[DEVICELIST]]" "-o" "[[LINKEDBC:.+\.bc]]"
-// CHK-FPGA-AOCO-EMU: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
+// CHK-FPGA-AOCO-EMU: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
 // CHK-FPGA-AOCO-EMU: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[SPLTABLE]]"
 // CHK-FPGA-AOCO-EMU: llvm-spirv{{.*}} "-o" "[[TARGSPV:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCO-EMU: opencl-aot{{.*}} "-device=fpga_fast_emu" "-spv=[[TARGSPV]]" "-ir=[[AOCXOUT:.+\.aocx]]"
diff --git a/clang/test/Driver/sycl-offload-intelfpga.cpp b/clang/test/Driver/sycl-offload-intelfpga.cpp
index ea015f9e66cd8..d9583f8fcc3eb 100644
--- a/clang/test/Driver/sycl-offload-intelfpga.cpp
+++ b/clang/test/Driver/sycl-offload-intelfpga.cpp
@@ -73,7 +73,7 @@
 // CHK-FPGA-LINK-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA-LINK: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]"
-// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
+// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
 // CHK-FPGA-LINK: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]"
 // CHK-FPGA-LINK: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-EARLY: aoc{{.*}} "-o" "[[OUTPUT4:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
@@ -107,7 +107,7 @@
 // CHK-FPGA-LINK-WIN-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK-WIN: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA-LINK-WIN: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]"
-// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
+// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
 // CHK-FPGA-LINK-WIN: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]"
 // CHK-FPGA-LINK-WIN: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-LINK-WIN: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
@@ -172,7 +172,7 @@
 // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=host-x86_64-unknown-linux-gnu,sycl-spir64_fpga-unknown-unknown" {{.*}} "-output=[[FINALLINK2:.+\.o]]" "-output=[[OUTPUT1:.+\.o]]" "-unbundle"
 // CHK-FPGA: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_BC:.+\.bc]]"
-// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]"
+// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]"
 // CHK-FPGA: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT3_TABLE]]"
 // CHK-FPGA: llvm-spirv{{.*}} "-o" "[[OUTPUT5:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-fpga_dep" {{.*}} "-output=[[DEPFILE:.+\.d]]" "-unbundle"
@@ -232,7 +232,7 @@
 // CHK-FPGA-AOCX-SRC: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.(o|obj)]]" "[[WRAPOUT]]"
 // CHK-FPGA-AOCX-SRC: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device" {{.*}} "-o" "[[DEVICEBC:.+\.bc]]"
 // CHK-FPGA-AOCX-SRC: llvm-link{{.*}} "[[DEVICEBC]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
+// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
 // CHK-FPGA-AOCX-SRC: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-SRC: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-SRC: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl"
@@ -256,7 +256,7 @@
 // CHK-FPGA-AOCX-OBJ: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-unbundle"
 // CHK-FPGA-AOCX-OBJ: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
+// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
 // CHK-FPGA-AOCX-OBJ: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-OBJ: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl"
@@ -275,7 +275,7 @@
 // CHK-FPGA-AOCX-OBJ2: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ2:.+\.(o|obj)]]" "-unbundle"
 // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=rt" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]"
+// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=rt" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-OBJ2: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" {{.*}} "-target=spir64" "-kind=sycl" "-batch"
@@ -286,7 +286,7 @@
 // CHK-FPGA-AOCX-OBJ2: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT2:.+\.(o|obj)]]" "[[WRAPOUT]]"
 // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ2]]" "-o" "[[IROUTPUT2:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT2]]" "-o" "[[LLVMLINKOUT2:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]"
+// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT2:.+\.txt]]" "[[POSTLINKOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT2:.+\.txt]]" {{.*}} "[[TABLEOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT2]]" "-sycl"
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index f71af5ded722b..1cecd9bb23367 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -171,6 +171,7 @@
 // CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLGlobalVariableAllowed (SubjectMatchRule_record)
+// CHECK-NEXT: SYCLHostPipe (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLIntelDisableLoopPipelining (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelInitiationInterval (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function)

From dc84cb934204c9d6e73bc716b48d8d284ea05e82 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Mon, 16 Jan 2023 11:24:38 -0800
Subject: [PATCH 04/14] Remove spurious XFAIL from host_pipe lit test

---
 clang/test/CodeGenSYCL/host_pipe.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/clang/test/CodeGenSYCL/host_pipe.cpp b/clang/test/CodeGenSYCL/host_pipe.cpp
index c7c88ff2a5d62..0ac50f6bcc5e1 100644
--- a/clang/test/CodeGenSYCL/host_pipe.cpp
+++ b/clang/test/CodeGenSYCL/host_pipe.cpp
@@ -5,8 +5,6 @@
 // global variable whose type is decorated with host_pipe attribute, and that a
 // unique string is generated.
 
-// XFAIL:*
-
 using namespace sycl::ext::intel::experimental;
 using namespace sycl;
 queue q;

From c2f11e01b8f3a48f1444d7c6682da4e8fde20017 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Wed, 18 Jan 2023 10:05:25 -0800
Subject: [PATCH 05/14] Use sycl_type to denote host pipes instead of new
 attribute

---
 clang/include/clang/Basic/Attr.td      | 13 ++---------
 clang/include/clang/Basic/AttrDocs.td  | 20 -----------------
 clang/lib/CodeGen/CodeGenModule.cpp    | 30 ++++++++++++++++++++++----
 clang/lib/Sema/SemaSYCL.cpp            |  8 +++----
 clang/test/CodeGenSYCL/Inputs/sycl.hpp |  4 +---
 5 files changed, 32 insertions(+), 43 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 299ebefd4d1c2..895218035a6d5 100644
--- 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];
 }
@@ -1554,15 +1554,6 @@ 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
index 12989451029dd..657832824776e 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3124,26 +3124,6 @@ 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/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index e1c07d5be43e6..dadadba12322e 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -110,6 +110,23 @@ static bool SYCLCUDAIsSYCLDevice(const clang::LangOptions &LangOpts) {
   return LangOpts.SYCLIsDevice && LangOpts.CUDA && !LangOpts.CUDAIsDevice;
 }
 
+static bool isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
+  const auto *RD = Ty->getAsCXXRecordDecl();
+  if (!RD)
+    return false;
+
+  if (const auto *Attr = RD->getAttr<SYCLTypeAttr>())
+    return Attr->getType() == TypeName;
+
+  if (const auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RD))
+    if (CXXRecordDecl *TemplateDecl =
+            CTSD->getSpecializedTemplate()->getTemplatedDecl())
+      if (const auto *Attr = TemplateDecl->getAttr<SYCLTypeAttr>())
+        return Attr->getType() == TypeName;
+
+  return false;
+}
+
 CodeGenModule::CodeGenModule(ASTContext &C,
                              IntrusiveRefCntPtr<llvm::vfs::FileSystem> FS,
                              const HeaderSearchOptions &HSO,
@@ -5512,10 +5529,15 @@ 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 or
-    // SYCL host_pipe attribute, emit IR attribute 'sycl-unique-id'.
-    if (RD && (RD->hasAttr<SYCLDeviceGlobalAttr>() || 
-               RD->hasAttr<SYCLHostPipeAttr>()))
+    // 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 VarDecl type is SYCLTypeAttr::host_pipe, emit the IR attribute 
+    // 'sycl-unique-id'.
+    auto Ty = D->getType();
+    if (isSyclType(Ty, SYCLTypeAttr::host_pipe))
       addSYCLUniqueID(GV, D, Context);
   }
 
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 66076697ced65..6bc12544c5058 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -5377,7 +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()) &&
+      !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
@@ -5563,9 +5563,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     // 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()) &&
-        !S.isTypeDecoratedWithDeclAttribute<SYCLHostPipeAttr>(
             VD->getType()))
       continue;
 
@@ -5595,8 +5594,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
                                                           VD);
       DeviceGlobOS << "\");\n";
-    } else if (S.isTypeDecoratedWithDeclAttribute<SYCLHostPipeAttr>(
-            VD->getType())) {
+    } else if (isSyclType(VD->getType(), SYCLTypeAttr::host_pipe)) {
       HostPipesEmitted = true;
       HostPipesOS << "host_pipe_map::add(";
       HostPipesOS << "(void *)&";
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 3b653b4bfec12..1b97f531bde2a 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -163,9 +163,7 @@ host_pipe {
 public:
   struct
 #ifdef __SYCL_DEVICE_ONLY__
-  // [[ __sycl_detail__::add_ir_attributes_global_variable(
-    // "sycl-host-pipe", nullptr)]] [[__sycl_detail__::host_pipe]]
-  [[__sycl_detail__::host_pipe]]
+  [[__sycl_detail__::sycl_type(host_pipe)]]
 #endif
   __pipeType { const char __p; };
 

From 20e3d807a6950cc3496290368a4733e9c5db014f Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Wed, 18 Jan 2023 10:45:22 -0800
Subject: [PATCH 06/14] Remove check for HostPipe attribute in attribute list
 test

---
 clang/test/Misc/pragma-attribute-supported-attributes-list.test | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 1cecd9bb23367..f71af5ded722b 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -171,7 +171,6 @@
 // CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLGlobalVariableAllowed (SubjectMatchRule_record)
-// CHECK-NEXT: SYCLHostPipe (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLIntelDisableLoopPipelining (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelInitiationInterval (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function)

From e8845cc53eb35c48ecb5e5595cd559ba8d1b07c3 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Tue, 24 Jan 2023 11:16:45 -0800
Subject: [PATCH 07/14] Simplify host_pipe attribute checking

---
 clang/lib/CodeGen/CodeGenModule.cpp | 47 ++++++++++-------------------
 1 file changed, 16 insertions(+), 31 deletions(-)

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index dadadba12322e..e46fae2c45d6a 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -110,23 +110,6 @@ static bool SYCLCUDAIsSYCLDevice(const clang::LangOptions &LangOpts) {
   return LangOpts.SYCLIsDevice && LangOpts.CUDA && !LangOpts.CUDAIsDevice;
 }
 
-static bool isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
-  const auto *RD = Ty->getAsCXXRecordDecl();
-  if (!RD)
-    return false;
-
-  if (const auto *Attr = RD->getAttr<SYCLTypeAttr>())
-    return Attr->getType() == TypeName;
-
-  if (const auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RD))
-    if (CXXRecordDecl *TemplateDecl =
-            CTSD->getSpecializedTemplate()->getTemplatedDecl())
-      if (const auto *Attr = TemplateDecl->getAttr<SYCLTypeAttr>())
-        return Attr->getType() == TypeName;
-
-  return false;
-}
-
 CodeGenModule::CodeGenModule(ASTContext &C,
                              IntrusiveRefCntPtr<llvm::vfs::FileSystem> FS,
                              const HeaderSearchOptions &HSO,
@@ -5525,20 +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 VarDecl type is SYCLTypeAttr::host_pipe, emit the IR attribute 
-    // 'sycl-unique-id'.
-    auto Ty = D->getType();
-    if (isSyclType(Ty, SYCLTypeAttr::host_pipe))
-      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()) {

From d0fe353947180dd29edd9b4685a618aaa8539f3e Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Wed, 1 Feb 2023 10:05:46 -0800
Subject: [PATCH 08/14] Remove unnecessary lit test args for host_pipe test

---
 clang/test/CodeGenSYCL/host_pipe.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CodeGenSYCL/host_pipe.cpp b/clang/test/CodeGenSYCL/host_pipe.cpp
index 0ac50f6bcc5e1..8cb6bd1977074 100644
--- a/clang/test/CodeGenSYCL/host_pipe.cpp
+++ b/clang/test/CodeGenSYCL/host_pipe.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -opaque-pointers -emit-llvm %s -o - | FileCheck %s
+// 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

From 7b828d77261f6cbc9ba7e13e2c9fb7dc1592ed66 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Wed, 1 Feb 2023 10:57:26 -0800
Subject: [PATCH 09/14] Remove another unnecessary arg for host_pipe test

---
 clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp b/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp
index cbf2c9f48dcff..31899db188470 100644
--- a/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -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: %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"

From e34ac9e23b47da9e157f0effee092354cd342f60 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Thu, 2 Feb 2023 13:07:46 -0800
Subject: [PATCH 10/14] Move host pipe compile time properties processing to
 opt

---
 clang/lib/Driver/ToolChains/Clang.cpp                         | 3 ---
 .../sycl-post-link => include/llvm/SYCLLowerIR}/HostPipes.h   | 0
 llvm/lib/SYCLLowerIR/CMakeLists.txt                           | 1 +
 llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp            | 1 +
 llvm/{tools/sycl-post-link => lib/SYCLLowerIR}/HostPipes.cpp  | 4 ++--
 .../CompileTimePropertiesPass}/host-pipes/basic.ll            | 4 +---
 6 files changed, 5 insertions(+), 8 deletions(-)
 rename llvm/{tools/sycl-post-link => include/llvm/SYCLLowerIR}/HostPipes.h (100%)
 rename llvm/{tools/sycl-post-link => lib/SYCLLowerIR}/HostPipes.cpp (92%)
 rename llvm/test/{tools/sycl-post-link => SYCLLowerIR/CompileTimePropertiesPass}/host-pipes/basic.ll (86%)

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 34ea69fdb4de8..ae467a22b500c 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -9702,9 +9702,6 @@ 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/llvm/tools/sycl-post-link/HostPipes.h b/llvm/include/llvm/SYCLLowerIR/HostPipes.h
similarity index 100%
rename from llvm/tools/sycl-post-link/HostPipes.h
rename to llvm/include/llvm/SYCLLowerIR/HostPipes.h
diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt
index b694a82aa0641..b95e09160d07a 100644
--- a/llvm/lib/SYCLLowerIR/CMakeLists.txt
+++ b/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
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index a3063750a3d02..0bc5de91ae636 100644
--- 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/HostPipes.h"
 
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/StringMap.h"
diff --git a/llvm/tools/sycl-post-link/HostPipes.cpp b/llvm/lib/SYCLLowerIR/HostPipes.cpp
similarity index 92%
rename from llvm/tools/sycl-post-link/HostPipes.cpp
rename to llvm/lib/SYCLLowerIR/HostPipes.cpp
index 591357e8984f0..5445fa4a01dd8 100644
--- a/llvm/tools/sycl-post-link/HostPipes.cpp
+++ b/llvm/lib/SYCLLowerIR/HostPipes.cpp
@@ -8,8 +8,8 @@
 // See comments in the header.
 //===----------------------------------------------------------------------===//
 
-#include "HostPipes.h"
-#include "CompileTimePropertiesPass.h"
+#include "llvm/SYCLLowerIR/HostPipes.h"
+#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/StringRef.h"
diff --git a/llvm/test/tools/sycl-post-link/host-pipes/basic.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll
similarity index 86%
rename from llvm/test/tools/sycl-post-link/host-pipes/basic.ll
rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll
index ca104c1fb24f9..f6009ad0b4583 100644
--- a/llvm/test/tools/sycl-post-link/host-pipes/basic.ll
+++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll
@@ -1,6 +1,4 @@
-; 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
+; 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

From 9e28e1690e47cc216198ebd4dc78569028ec32e9 Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Thu, 2 Feb 2023 13:19:23 -0800
Subject: [PATCH 11/14] Remove -host-pipes arg from clang driver tests

---
 clang/test/Driver/sycl-device-lib.cpp        |  2 +-
 clang/test/Driver/sycl-intelfpga-aoco.cpp    |  4 ++--
 clang/test/Driver/sycl-offload-intelfpga.cpp | 14 +++++++-------
 3 files changed, 10 insertions(+), 10 deletions(-)

diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp
index e5012682a42ea..4a2ad0fc8bcbc 100644
--- a/clang/test/Driver/sycl-device-lib.cpp
+++ b/clang/test/Driver/sycl-device-lib.cpp
@@ -174,7 +174,7 @@
 // RUN:   | FileCheck %s -check-prefix=SYCL_LLVM_LINK_NO_DEVICE_LIB
 // SYCL_LLVM_LINK_NO_DEVICE_LIB: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device"
 // SYCL_LLVM_LINK_NO_DEVICE_LIB-NOT: llvm-link{{.*}}  "-only-needed"
-// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-emit-exported-symbols" "-split-esimd" "-lower-esimd" "-O2" "-spec-const=rt" "-device-globals" "-host-pipes" "-o" "{{.*}}.table" "{{.*}}.bc"
+// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-emit-exported-symbols" "-split-esimd" "-lower-esimd" "-O2" "-spec-const=rt" "-device-globals" "-o" "{{.*}}.table" "{{.*}}.bc"
 
 /// ###########################################################################
 /// test llvm-link behavior for special user input whose filename resembles SYCL device library
diff --git a/clang/test/Driver/sycl-intelfpga-aoco.cpp b/clang/test/Driver/sycl-intelfpga-aoco.cpp
index d1e82c835dcf1..aefcd262edfb2 100755
--- a/clang/test/Driver/sycl-intelfpga-aoco.cpp
+++ b/clang/test/Driver/sycl-intelfpga-aoco.cpp
@@ -52,7 +52,7 @@
 // RUN:  %clang_cl -fsycl -fno-sycl-instrument-device-code -fno-sycl-device-lib=all -fintelfpga -Xshardware %t_aoco_cl.a -### %s 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCO,CHK-FPGA-AOCO-WIN %s
 // CHK-FPGA-AOCO: llvm-link{{.*}} "-o" "[[LINKEDBC:.+\.bc]]"
-// CHK-FPGA-AOCO: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
+// CHK-FPGA-AOCO: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
 // CHK-FPGA-AOCO: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[SPLTABLE]]"
 // CHK-FPGA-AOCO: llvm-spirv{{.*}} "-o" "[[TARGSPV:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCO: clang-offload-bundler{{.*}} "-type=aoo" "-targets=sycl-fpga_aoco-intel-unknown" "-input=[[INPUTLIB:.+\.a]]" "-output=[[AOCOLIST:.+\.txt]]" "-unbundle"
@@ -103,7 +103,7 @@
 // CHK-FPGA-AOCO-EMU: clang-offload-bundler{{.*}} "-type=aoo" "-targets=sycl-spir64_fpga-unknown-unknown" "-input=[[INPUTLIB:.+\.a]]" "-output=[[OUTLIB:.+\.txt]]" "-unbundle"
 // CHK-FPGA-AOCO-EMU: llvm-foreach{{.*}} "--out-ext=txt" "--in-file-list=[[OUTLIB]]" "--in-replace=[[OUTLIB]]" "--out-file-list=[[DEVICELIST:.+\.txt]]" "--out-replace=[[DEVICELIST]]" "--" {{.*}}spirv-to-ir-wrapper{{.*}} "[[OUTLIB]]" "-o" "[[DEVICELIST]]"
 // CHK-FPGA-AOCO-EMU: llvm-link{{.*}} "@[[DEVICELIST]]" "-o" "[[LINKEDBC:.+\.bc]]"
-// CHK-FPGA-AOCO-EMU: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
+// CHK-FPGA-AOCO-EMU: sycl-post-link{{.*}} "-split-esimd"{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[SPLTABLE:.+\.table]]" "[[LINKEDBC]]"
 // CHK-FPGA-AOCO-EMU: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[SPLTABLE]]"
 // CHK-FPGA-AOCO-EMU: llvm-spirv{{.*}} "-o" "[[TARGSPV:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCO-EMU: opencl-aot{{.*}} "-device=fpga_fast_emu" "-spv=[[TARGSPV]]" "-ir=[[AOCXOUT:.+\.aocx]]"
diff --git a/clang/test/Driver/sycl-offload-intelfpga.cpp b/clang/test/Driver/sycl-offload-intelfpga.cpp
index d9583f8fcc3eb..ea015f9e66cd8 100644
--- a/clang/test/Driver/sycl-offload-intelfpga.cpp
+++ b/clang/test/Driver/sycl-offload-intelfpga.cpp
@@ -73,7 +73,7 @@
 // CHK-FPGA-LINK-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA-LINK: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]"
-// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
+// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
 // CHK-FPGA-LINK: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]"
 // CHK-FPGA-LINK: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-EARLY: aoc{{.*}} "-o" "[[OUTPUT4:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
@@ -107,7 +107,7 @@
 // CHK-FPGA-LINK-WIN-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK-WIN: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA-LINK-WIN: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]"
-// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
+// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]"
 // CHK-FPGA-LINK-WIN: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]"
 // CHK-FPGA-LINK-WIN: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-LINK-WIN: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
@@ -172,7 +172,7 @@
 // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=host-x86_64-unknown-linux-gnu,sycl-spir64_fpga-unknown-unknown" {{.*}} "-output=[[FINALLINK2:.+\.o]]" "-output=[[OUTPUT1:.+\.o]]" "-unbundle"
 // CHK-FPGA: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]"
 // CHK-FPGA: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_BC:.+\.bc]]"
-// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]"
+// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]"
 // CHK-FPGA: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT3_TABLE]]"
 // CHK-FPGA: llvm-spirv{{.*}} "-o" "[[OUTPUT5:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]"
 // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-fpga_dep" {{.*}} "-output=[[DEPFILE:.+\.d]]" "-unbundle"
@@ -232,7 +232,7 @@
 // CHK-FPGA-AOCX-SRC: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.(o|obj)]]" "[[WRAPOUT]]"
 // CHK-FPGA-AOCX-SRC: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device" {{.*}} "-o" "[[DEVICEBC:.+\.bc]]"
 // CHK-FPGA-AOCX-SRC: llvm-link{{.*}} "[[DEVICEBC]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
+// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
 // CHK-FPGA-AOCX-SRC: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-SRC: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-SRC: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl"
@@ -256,7 +256,7 @@
 // CHK-FPGA-AOCX-OBJ: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-unbundle"
 // CHK-FPGA-AOCX-OBJ: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
+// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]
 // CHK-FPGA-AOCX-OBJ: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-OBJ: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl"
@@ -275,7 +275,7 @@
 // CHK-FPGA-AOCX-OBJ2: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ2:.+\.(o|obj)]]" "-unbundle"
 // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=rt" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]"
+// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=rt" "-device-globals" "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]"
 // CHK-FPGA-AOCX-OBJ2: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" {{.*}} "-target=spir64" "-kind=sycl" "-batch"
@@ -286,7 +286,7 @@
 // CHK-FPGA-AOCX-OBJ2: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT2:.+\.(o|obj)]]" "[[WRAPOUT]]"
 // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ2]]" "-o" "[[IROUTPUT2:.+\.bc]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT2]]" "-o" "[[LLVMLINKOUT2:.+\.bc]]" "--suppress-warnings"
-// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-host-pipes" "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]"
+// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-spec-const=default" "-device-globals" "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT2:.+\.txt]]" "[[POSTLINKOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT2:.+\.txt]]" {{.*}} "[[TABLEOUT2]]"
 // CHK-FPGA-AOCX-OBJ2: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT2]]" "-sycl"

From aad8b8c4561b12e47fd794ad1e53dd576a51660b Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Thu, 2 Feb 2023 15:15:56 -0800
Subject: [PATCH 12/14] Cleanup unnecessary include and class declarations in
 HostPipes header

---
 llvm/include/llvm/SYCLLowerIR/HostPipes.h | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/llvm/include/llvm/SYCLLowerIR/HostPipes.h b/llvm/include/llvm/SYCLLowerIR/HostPipes.h
index cf3de06d5bac2..6f360fc560fa6 100644
--- a/llvm/include/llvm/SYCLLowerIR/HostPipes.h
+++ b/llvm/include/llvm/SYCLLowerIR/HostPipes.h
@@ -15,13 +15,10 @@
 #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.
 ///

From dacaed38fdb6b821affa29ff6dcf1bb53b5397af Mon Sep 17 00:00:00 2001
From: "Ho, Robert" <robert.ho@intel.com>
Date: Mon, 6 Feb 2023 10:07:15 -0800
Subject: [PATCH 13/14] Move HostPipes check function to SYCLUtils; remove
 extraneous sycl-post-link flag

---
 llvm/include/llvm/SYCLLowerIR/HostPipes.h     | 31 --------------
 llvm/include/llvm/SYCLLowerIR/SYCLUtils.h     |  7 ++++
 llvm/lib/SYCLLowerIR/CMakeLists.txt           |  1 -
 .../SYCLLowerIR/CompileTimePropertiesPass.cpp |  4 +-
 llvm/lib/SYCLLowerIR/HostPipes.cpp            | 42 -------------------
 llvm/tools/sycl-post-link/sycl-post-link.cpp  |  8 +---
 6 files changed, 10 insertions(+), 83 deletions(-)
 delete mode 100644 llvm/include/llvm/SYCLLowerIR/HostPipes.h
 delete mode 100644 llvm/lib/SYCLLowerIR/HostPipes.cpp

diff --git a/llvm/include/llvm/SYCLLowerIR/HostPipes.h b/llvm/include/llvm/SYCLLowerIR/HostPipes.h
deleted file mode 100644
index 6f360fc560fa6..0000000000000
--- a/llvm/include/llvm/SYCLLowerIR/HostPipes.h
+++ /dev/null
@@ -1,31 +0,0 @@
-//===------- 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>
-
-namespace llvm {
-
-class GlobalVariable;
-
-/// 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/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
index 84d37370f4994..65bedad4a6d11 100644
--- 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,11 +22,13 @@ 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 =
     std::function<bool(const Instruction *, const Function *)>;
 
+
 // Traverses call graph starting from given function up the call chain applying
 // given action to each function met on the way. If \c ErrorOnNonCallUse
 // parameter is true, then no functions' uses are allowed except calls.
@@ -115,6 +118,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/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt
index b95e09160d07a..b694a82aa0641 100644
--- a/llvm/lib/SYCLLowerIR/CMakeLists.txt
+++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt
@@ -57,7 +57,6 @@ add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/LowerESIMDVecArg.cpp
   ESIMD/LowerESIMDVLoadVStore.cpp
   ESIMD/LowerESIMDSlmReservation.cpp
-  HostPipes.cpp
   LowerInvokeSimd.cpp
   LowerKernelProps.cpp
   LowerWGLocalMemory.cpp
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 0bc5de91ae636..43e5d94b1df3d 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -10,7 +10,7 @@
 
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
-#include "llvm/SYCLLowerIR/HostPipes.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/StringMap.h"
@@ -267,7 +267,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
                                               HostAccessDecorValue, VarName));
     }
 
-    if (isHostPipeVariable(GV)) {
+    if (sycl::utils::isHostPipeVariable(GV)) {
       auto VarName = getGlobalVariableUniqueId(GV);
       MDOps.push_back(buildSpirvDecorMetadata(Ctx, SPIRV_HOST_ACCESS_DECOR,
                                               SPIRV_HOST_ACCESS_DEFAULT_VALUE, 
diff --git a/llvm/lib/SYCLLowerIR/HostPipes.cpp b/llvm/lib/SYCLLowerIR/HostPipes.cpp
deleted file mode 100644
index 5445fa4a01dd8..0000000000000
--- a/llvm/lib/SYCLLowerIR/HostPipes.cpp
+++ /dev/null
@@ -1,42 +0,0 @@
-//===------------- 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/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/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp
index b01e9f2cbdac7..f9110752bc331 100644
--- a/llvm/tools/sycl-post-link/sycl-post-link.cpp
+++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -207,11 +207,6 @@ 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;
@@ -975,11 +970,10 @@ 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 && !DoHostPipes) {
+      !DoLowerEsimd) {
     errs() << "no actions specified; try --help for usage info\n";
     return 1;
   }

From 82a1efa750a54826a50abbe017e065b10a5e2112 Mon Sep 17 00:00:00 2001
From: Robert Ho <84344325+rho180@users.noreply.github.com>
Date: Mon, 6 Feb 2023 14:54:04 -0500
Subject: [PATCH 14/14] Update llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
---
 llvm/include/llvm/SYCLLowerIR/SYCLUtils.h | 1 -
 1 file changed, 1 deletion(-)

diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
index 65bedad4a6d11..2d82e651b35fb 100644
--- a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
+++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
@@ -28,7 +28,6 @@ using CallGraphNodeAction = ::std::function<void(Function *)>;
 using CallGraphFunctionFilter =
     std::function<bool(const Instruction *, const Function *)>;
 
-
 // Traverses call graph starting from given function up the call chain applying
 // given action to each function met on the way. If \c ErrorOnNonCallUse
 // parameter is true, then no functions' uses are allowed except calls.