Merge branch 'sycl/unified/master' into MergeIntel

README.md

@@ -56,7 +56,7 @@ including for example an Intel FPGA and an AMD/Xilinx FPGA.
 
 - AMD/Xilinx FPGA Tests Documentation
   - [Tests.md](sycl/doc/Tests.md) covers a few details about the the
-    additional [xocc_tests](sycl/test/xocc_tests) directory we added
+    additional [vitis](sycl/test/vitis) directory we added
     to the [sycl/test](sycl/test) directory among some other small
     details.
 

clang/lib/Sema/SemaSYCL.cpp

@@ -1909,7 +1909,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   // attribute.
   void handleXilinxProperty(ParmVarDecl *Param, QualType PropTy,
                             SourceLocation Loc) {
-    if (!isSyclXilinxType(PropTy))
+    if (!SemaRef.getASTContext().getTargetInfo().getTriple().isXilinxFPGA() ||
+        !isSyclXilinxType(PropTy))
       return;
     ASTContext &Ctx = SemaRef.getASTContext();
     const CXXRecordDecl *RD = PropTy->getAsCXXRecordDecl();

llvm/include/llvm/SYCL/KernelProperties.h

@@ -66,7 +66,9 @@ class KernelProperties {
   SmallVector<MAXIBundle, 8> Bundles;
 
 public:
-  static bool isArgBuffer(Argument* Arg, bool SyclHLSFlow);
+  /// Return true iff Arg represents a buffer in the OpenCL sense equivalent to
+  /// a SYCL accessor's pointer on the data
+  static bool isArgBuffer(Argument *Arg, bool SyclHLSFlow);
   KernelProperties(Function &F, bool SyclHlsFlow);
   KernelProperties(KernelProperties &) = delete;
 

llvm/lib/SYCL/InSPIRation.cpp

@@ -235,22 +235,6 @@ struct InSPIRationState {
     return false;
   }
 
-  /// This function gives llvm::function arguments with no name
-  /// a default name e.g. arg_0, arg_1..
-  ///
-  /// This is because if your arguments have no name v++ will commit seppuku
-  /// when generating XML. Perhaps it's possible to move this to the Clang
-  /// Frontend by generating the name from the accessor/capture the arguments
-  /// come from, but I believe it requires a special compiler invocation option
-  /// to keep arg names from the frontend in the LLVM bitcode.
-  void giveNameToArguments(Function &F) {
-    int Counter = 0;
-    for (auto &Arg : F.args()) {
-      if (!Arg.hasName())
-        Arg.setName("arg_" + Twine{Counter++});
-    }
-  }
-
   // Hopeful list/probably impractical asks for v++:
   // 1) Make XML generator/reader a little kinder towards arguments with no
   //   names if possible
@@ -274,10 +258,10 @@ struct InSPIRationState {
 
     std::vector<Function *> Declarations;
     for (auto &F : M.functions()) {
-      if (isKernelFunc(&F)) {
+      if (sycl::isKernelFunc(&F)) {
         kernelSPIRify(F);
         applyKernelProperties(F);
-        giveNameToArguments(F);
+        sycl::giveNameToArguments(F);
 
         /// \todo Possible: We don't modify declarations right now as this
         /// will destroy the names of SPIR/CL intrinsics as they aren't
@@ -313,7 +297,7 @@ struct InSPIRationState {
         //
         // It doesn't require application to the SPIR intrinsics as we're
         // linking against the HLS SPIR library, which is already conformant.
-        giveNameToArguments(F);
+        sycl::giveNameToArguments(F);
       } else if (isTransitiveNonIntrinsicFunc(F) && F.isDeclaration()) {
         // push back intrinsics to make sure we handle naming after changing the
         // name of all functions to sycl_func.

llvm/lib/SYCL/KernelPropGen.cpp

@@ -118,18 +118,76 @@ struct KernelPropGenState {
     return {};
   }
 
+  struct PipeEndpoint {
+    /// Name of the kernel function in IR.
+    StringRef Kernel;
+    /// Name of the pipe function argument in IR.
+    StringRef Arg;
+  };
+  struct PipeProp {
+    /// Depth it defaults to -1 to indicate it is unset
+    int Depth = -1;
+    PipeEndpoint write;
+    PipeEndpoint read;
+  };
+
+  /// Pipes are matched in read and write pairs by their ID. Their ID is a
+  /// string matching the name of the global variable Intel uses to identify its
+  /// pipes
+  StringMap<PipeProp> PipeConnections;
+
+  void collectPipeConnections(Module &M) {
+    for (auto &F : M.functions())
+      if (sycl::isKernelFunc(&F))
+        for (auto &Arg : F.args())
+          if (sycl::isPipe(&Arg)) {
+            /// Build the endpoint for this pipe
+            PipeEndpoint endPoint{F.getName(), Arg.getName()};
+            PipeProp &Prop = PipeConnections[sycl::getPipeID(&Arg)];
+
+            /// Figure out the correct endpoint to write to
+            PipeEndpoint &mapEndPoint =
+                sycl::isReadPipe(&Arg) ? Prop.read : Prop.write;
+            assert(mapEndPoint.Arg.empty() && mapEndPoint.Kernel.empty() &&
+                   "multiple reader or writers");
+            mapEndPoint = endPoint;
+
+            /// If the Depth is unset, set it
+            if (Prop.Depth == -1)
+              Prop.Depth = sycl::getPipeDepth(&Arg);
+
+            assert(sycl::getPipeDepth(&Arg) == Prop.Depth &&
+                   "read and write depth not matching");
+          }
+  }
+
   /// Print in O the property file for all kernels of M
   void generateProperties(Module &M, llvm::raw_fd_ostream &O) {
     json::OStream J(O, 2);
     llvm::json::Array Kernels{};
     bool SyclHlsFlow = Triple(M.getTargetTriple()).isXilinxHLS();
     bool VitisHlsFlow = Triple(M.getTargetTriple()).getArch() == llvm::Triple::vitis_ip;
 
+    collectPipeConnections(M);
+
     J.objectBegin();
+    J.attributeBegin("pipe_connections");
+    J.arrayBegin();
+    for (auto& Elem : PipeConnections) {
+      J.objectBegin();
+      J.attribute("writer_kernel", Elem.second.write.Kernel);
+      J.attribute("writer_arg", Elem.second.write.Arg);
+      J.attribute("reader_kernel", Elem.second.read.Kernel);
+      J.attribute("reader_arg", Elem.second.read.Arg);
+      J.attribute("depth", Elem.second.Depth);
+      J.objectEnd();
+    }
+    J.arrayEnd();
+    J.attributeEnd();
     J.attributeBegin("kernels");
     J.arrayBegin();
     for (auto &F : M.functions()) {
-      if (isKernelFunc(&F)) {
+      if (sycl::isKernelFunc(&F)) {
         KernelProperties KProp(F, SyclHlsFlow);
         J.objectBegin();
         J.attribute("name", F.getName());
@@ -162,8 +220,14 @@ struct KernelPropGenState {
         J.attributeBegin("arg_bundle_mapping");
         J.arrayBegin();
         for (auto &Arg : F.args()) {
-          if (!VitisHlsFlow &&
-              KernelProperties::isArgBuffer(&Arg, SyclHlsFlow)) {
+          if (VitisHlsFlow)
+            continue;
+          /// Vitis's clang doesn't support string attributes on arguments which
+          /// we use to annotate a pipe, so we remove it here. But we could
+          /// remove them in the downgrader instead too.
+          if (sycl::isPipe(&Arg))
+            sycl::removePipeAnnotation(&Arg);
+          else if (KernelProperties::isArgBuffer(&Arg, SyclHlsFlow)) {
             // This currently forces a default assignment of DDR banks to 0
             // as some platforms have different Default DDR banks and buffers
             // default to DDR Bank 0. Perhaps it is possible to query the

llvm/lib/SYCL/KernelProperties.cpp

@@ -28,6 +28,8 @@
 #include <string>
 #include <utility>
 
+#include "SYCLUtils.h"
+
 using namespace llvm;
 namespace {
 enum SPIRAddressSpace {
@@ -103,6 +105,10 @@ Optional<KernelProperties::MemBankSpec> getUserSpecifiedBank(
 
 namespace llvm {
 bool KernelProperties::isArgBuffer(Argument *Arg, bool SyclHLSFlow) {
+  /// We consider that pointer arguments that are not byval or pipes are
+  /// buffers.
+  if (sycl::isPipe(Arg))
+    return false;
   if (Arg->getType()->isPointerTy() &&
       (SyclHLSFlow ||
        Arg->getType()->getPointerAddressSpace() == SPIRAS_Global ||