[SYCL][FPGA] Implement IO pipes interface

MrSidims · bader · commit c900248f5d4c · 2020-02-27T10:13:10.000+03:00
kernel_readable_io_pipe and kernel_writeable_io_pipe classes
were added.

Signed-off-by: Dmitry Sidorov &lt;dmitry.sidorov@intel.com&gt;
diff --git a/sycl/doc/extensions/fpga_io_pipes_design.rst b/sycl/doc/extensions/fpga_io_pipes_design.rst
@@ -0,0 +1,81 @@
+IO pipes design
+===============
+
+Requirements
+------------
+ - Device shall be able to distinguish kernel-to-kernel pipes and I/O pipes;
+ - No changes shall be in SYCL pipe specification;
+ - I/O pipe namings/IDs are provided by a vendor in a separated header, like:
+.. code:: cpp
+  namespace intelfpga {
+    template <unsigned ID>
+    struct ethernet_pipe_id {
+      static constexpr unsigned id = ID;
+    };
+    using ethernet_read_pipe =
+        cl::sycl::intel::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
+    using ethernet_write_pipe =
+        cl::sycl::intel::kernel_writeable_io_pipe<ethernet_pipe_id<1>, int, 0>;
+  }
+
+  Thus, the user interacts only with vendor-defined pipe objects.
+
+Links
+-----
+.. _Spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DataFlowPipes/data_flow_pipes.asciidoc
+.. _Interesting comment from Ronan: https://github.com/intel/llvm/pull/635#discussion_r325851766
+
+Summary
+-------
+This document describes a design of I/O pipes implementation in SYCL compiler.
+It includes changes in SYCL library, clang and SPIR-V to LLVM IR translator.
+It adds extra attribute '__attribute__((io_pipe_id(ID)))' that generates
+metadata, attached to a pipe storage declaration. By this metadata a backend can
+recognize I/O pipe and distinguish different I/O pipes from each other.
+
+There is another notable solution which was proposed by Ronan (see the link
+above): don't make any compiler/library changes - just make backend recognizing
+I/O pipe by demangling it's name. This proposal wasn't picked, because it will
+make backend support of the feature more difficult. So far we already have
+two devices' backends that support SYCL_INTEL_data_flow_pipes extension
+(Intel FPGA HW and Intel FPGA emulator) and in the future this number may
+increase. So efforts put in compiler implementation shall be payed off even
+more.
+
+clang
+-----
+Need to implement additional attribute, that possible to attach to pipe storage
+declaration. The attribute shall accept a compile-time known integer argument
+(the pipe ID). With the attribute applied, clang generates a metadata attached
+the to pipe storage declaration, that contains the I/O pipe ID (argument).
+
+llvm-spirv translator
+---------------------
+Need to implement additional decoration, that saves the I/O pipe ID information,
+that can be collected from a metadata attached to the pipe storage object.
+
+SYCL implementation in headers
+------------------------------
+Following the spec, we need to add two more classes for pipes:
+ - 'kernel_readable_io_pipe'
+ - 'kernel_writeable_io_pipe'
+
+ with the same member functions and fields as it is already done for pipe class.
+
+The attribute should be attached to a pipe storage declaration in the headers
+and it would be looking like:
+.. code:: cpp
+  static constexpr int32_t ID = name::id;
+  static constexpr struct ConstantPipeStorage
+  m_Storage __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
+
+
+ where 'name' is some class used in pipe type construction, like:
+.. code:: cpp
+ using pipe_type = pipe<name, dataT, min_capacity>;
+
+When specific io_pipe_def.h is included in a user's code and 'name' is mapped to
+the pipe name defined in this header (for example 'ethernet_pipe_id' structure
+defined above) 'name::id' returns the actual I/O pipe ID (compile-time known
+integer constant) that is passed as the attribute's argument and used to
+indentify the I/O pipe in some backend.
diff --git a/sycl/include/CL/sycl/intel/pipes.hpp b/sycl/include/CL/sycl/intel/pipes.hpp
@@ -84,6 +84,112 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
 #endif // __SYCL_DEVICE_ONLY__
 };
 
+// IO pipes that provide interface to connect with hardware peripheral.
+// Their name aliases are defined in vendor-provided header, below you can see
+// an example of this header. There are defined aliases to ethernet_read_pipe
+// and ethernet_write_pipe that users can use in their code to connect with
+// HW peripheral.
+/* namespace intelfpga {
+template <int32_t ID>
+struct ethernet_pipe_id {
+  static constexpr int32_t id = ID;
+};
+
+template <class dataT, size_t min_capacity>
+using ethernet_read_pipe =
+  kernel_readable_io_pipe<ethernet_pipe_id<0>, dataT, min_capacity>;
+
+template <class dataT, size_t min_capacity>
+using ethernet_write_pipe =
+  kernel_writeable_io_pipe<ethernet_pipe_id<1>, dataT, min_capacity>;
+} // namespace intelfpga */
+
+template <class name, class dataT, size_t min_capacity = 0>
+class kernel_readable_io_pipe {
+public:
+  // Non-blocking pipes
+  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
+  // friendly LLVM IR.
+  static dataT read(bool &Success) {
+#ifdef __SYCL_DEVICE_ONLY__
+    RPipeTy<dataT> RPipe =
+        __spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
+    dataT TempData;
+    Success = !static_cast<bool>(
+        __spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
+    return TempData;
+#else
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  // Blocking pipes
+  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
+  // friendly LLVM IR.
+  static dataT read() {
+#ifdef __SYCL_DEVICE_ONLY__
+    RPipeTy<dataT> RPipe =
+        __spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
+    dataT TempData;
+    __spirv_ReadPipeBlockingINTEL(RPipe, &TempData, m_Size, m_Alignment);
+    return TempData;
+#else
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+private:
+  static constexpr int32_t m_Size = sizeof(dataT);
+  static constexpr int32_t m_Alignment = alignof(dataT);
+  static constexpr int32_t m_Capacity = min_capacity;
+  static constexpr int32_t ID = name::id;
+#ifdef __SYCL_DEVICE_ONLY__
+  static constexpr struct ConstantPipeStorage m_Storage
+      __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
+#endif // __SYCL_DEVICE_ONLY__
+};
+
+template <class name, class dataT, size_t min_capacity = 0>
+class kernel_writeable_io_pipe {
+public:
+  // Non-blocking pipes
+  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
+  // friendly LLVM IR.
+  static void write(const dataT &Data, bool &Success) {
+#ifdef __SYCL_DEVICE_ONLY__
+    WPipeTy<dataT> WPipe =
+        __spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
+    Success = !static_cast<bool>(
+        __spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
+#else
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  // Blocking pipes
+  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
+  // friendly LLVM IR.
+  static void write(const dataT &Data) {
+#ifdef __SYCL_DEVICE_ONLY__
+    WPipeTy<dataT> WPipe =
+        __spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
+    __spirv_WritePipeBlockingINTEL(WPipe, &Data, m_Size, m_Alignment);
+#else
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+private:
+  static constexpr int32_t m_Size = sizeof(dataT);
+  static constexpr int32_t m_Alignment = alignof(dataT);
+  static constexpr int32_t m_Capacity = min_capacity;
+  static constexpr int32_t ID = name::id;
+#ifdef __SYCL_DEVICE_ONLY__
+  static constexpr struct ConstantPipeStorage m_Storage
+      __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
+#endif // __SYCL_DEVICE_ONLY__
+};
+
 } // namespace intel
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/test/fpga_tests/fpga_io_pipes.cpp b/sycl/test/fpga_tests/fpga_io_pipes.cpp
@@ -0,0 +1,132 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+//==------------ fpga_io_pipes.cpp - SYCL FPGA pipes test ------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <CL/sycl/intel/fpga_extensions.hpp>
+#include <fstream>
+#include <iostream>
+
+#include "io_pipe_def.h"
+
+// TODO: run is disabled, since no support added in FPGA backend yet. Check
+// implementation correctness from CXX and SYCL languages perspective.
+
+// This test is supposed to be run only on Intel FPGA emulator. Change it when
+// we have more experience with IO pipe feature in SYCL.
+// The emulator creates files (one for I pipe, another for O pipe) with the
+// appropriate naming, where a data flowing through a pipe can be stored.
+// So in the test we need to create these files and use them appropriately.
+// The name is taken as IO pipe ID.
+const size_t InputData = 42;
+const std::string InputFileName = "0.txt";
+const std::string OutputFileName = "1.txt";
+
+void createInputFile(const std::string &filename) {
+  std::ofstream Input(filename);
+  if (Input.is_open()) {
+    Input << InputData;
+    Input.close();
+  }
+}
+
+int validateOutputFile(const std::string &filename) {
+  std::ifstream Output(filename);
+  std::string Line;
+  std::vector<size_t> Result;
+  if (Output.is_open()) {
+    // In the test we write only one number into the pipe, but a backend might
+    // have a bug of incorrect interpretetion of capacity of the pipe. In this
+    // case let's read all the lines of the output file to catch this.
+    while (std::getline(Output, Line))
+      Result.push_back(stoi(Line));
+    Output.close();
+  }
+  if (Result.size() != 1 || Result[0] != InputData) {
+    std::cout << "Result mismatches " << Result[0] << " Vs expected "
+              << InputData << std::endl;
+    return -1;
+  }
+
+  return 0;
+}
+
+// Test for simple non-blocking pipes
+int test_io_nb_pipe(cl::sycl::queue Queue) {
+  createInputFile(InputFileName);
+
+  cl::sycl::buffer<int, 1> writeBuf(1);
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    auto write_acc = writeBuf.get_access<cl::sycl::access::mode::write>(cgh);
+
+    cgh.single_task<class nb_io_transfer>([=]() {
+      bool SuccessCodeI = false;
+      do {
+        write_acc[0] = intelfpga::ethernet_read_pipe::read(SuccessCodeI);
+      } while (!SuccessCodeI);
+      bool SuccessCodeO = false;
+      do {
+        intelfpga::ethernet_write_pipe::write(write_acc[0], SuccessCodeO);
+      } while (!SuccessCodeO);
+    });
+  });
+
+  auto readHostBuffer = writeBuf.get_access<cl::sycl::access::mode::read>();
+  if (readHostBuffer[0] != InputData) {
+    std::cout << "Read from a file mismatches " << readHostBuffer[0]
+              << " Vs expected " << InputData << std::endl;
+
+    return -1;
+  }
+
+  return validateOutputFile(OutputFileName);
+}
+
+// Test for simple blocking pipes
+int test_io_bl_pipe(cl::sycl::queue Queue) {
+  createInputFile(InputFileName);
+
+  cl::sycl::buffer<int, 1> writeBuf(1);
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    auto write_acc = writeBuf.get_access<cl::sycl::access::mode::write>(cgh);
+
+    cgh.single_task<class bl_io_transfer>([=]() {
+      write_acc[0] = intelfpga::ethernet_read_pipe::read();
+      intelfpga::ethernet_write_pipe::write(write_acc[0]);
+    });
+  });
+
+  auto readHostBuffer = writeBuf.get_access<cl::sycl::access::mode::read>();
+  if (readHostBuffer[0] != InputData) {
+    std::cout << "Read from a file mismatches " << readHostBuffer[0]
+              << " Vs expected " << InputData << std::endl;
+
+    return -1;
+  }
+
+  return validateOutputFile(OutputFileName);
+}
+
+int main() {
+  cl::sycl::queue Queue{cl::sycl::intel::fpga_emulator_selector{}};
+
+  if (!Queue.get_device()
+           .get_info<cl::sycl::info::device::kernel_kernel_pipe_support>()) {
+    std::cout << "SYCL_INTEL_data_flow_pipes not supported, skipping"
+              << std::endl;
+    return 0;
+  }
+
+  // Non-blocking pipes
+  int Result = test_io_nb_pipe(Queue);
+
+  // Blocking pipes
+  Result &= test_io_bl_pipe(Queue);
+
+  return Result;
+}
diff --git a/sycl/test/fpga_tests/io_pipe_def.h b/sycl/test/fpga_tests/io_pipe_def.h
@@ -0,0 +1,12 @@
+#include <CL/sycl/intel/fpga_extensions.hpp>
+
+namespace intelfpga {
+template <unsigned ID> struct ethernet_pipe_id {
+  static constexpr unsigned id = ID;
+};
+
+using ethernet_read_pipe =
+    sycl::intel::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
+using ethernet_write_pipe =
+    sycl::intel::kernel_writeable_io_pipe<ethernet_pipe_id<1>, int, 0>;
+} // namespace intelfpga
