[SYCL][FPGA] Add compute units library.

sycl/include/CL/sycl/INTEL/compute_units.hpp

@@ -0,0 +1,51 @@
+//==--------- compute_units.hpp - SYCL Compute Units -------*- C++ -*-------==//
+//
+// 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
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/detail/defines.hpp>
+#include <utility>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace INTEL {
+
+namespace {
+template <class Func, template <std::size_t> class Name, std::size_t index>
+class SubmitOneComputeUnit {
+public:
+  SubmitOneComputeUnit(Func &&f, sycl::queue &q) {
+
+    q.submit([&](sycl::handler &h) {
+      h.single_task<Name<index>>([=]() {
+        // verifies that f only takes a single argument
+        f(std::integral_constant<std::size_t, index>());
+      });
+    });
+  }
+};
+
+template <template <std::size_t> class Name, class Func, std::size_t... index>
+inline constexpr void ComputeUnitUnroller(sycl::queue &q, Func &&f,
+                                          std::index_sequence<index...>) {
+  (SubmitOneComputeUnit<Func, Name, index>(f, q), ...); // fold expression
+}
+} // namespace
+
+// N is the number of compute units
+// Name is the kernel's name
+template <std::size_t N, template <std::size_t ID> class Name, class Func>
+constexpr void submit_compute_units(sycl::queue &q, Func &&f) {
+  std::make_index_sequence<N> indices;
+  ComputeUnitUnroller<Name>(q, f, indices);
+}
+
+} // namespace INTEL
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/INTEL/fpga_extensions.hpp

@@ -7,7 +7,9 @@
 //===----------------------------------------------------------------------===//
 
 #pragma once
+#include <CL/sycl/INTEL/compute_units.hpp>
 #include <CL/sycl/INTEL/fpga_device_selector.hpp>
 #include <CL/sycl/INTEL/fpga_lsu.hpp>
 #include <CL/sycl/INTEL/fpga_reg.hpp>
+#include <CL/sycl/INTEL/pipe_array.hpp>
 #include <CL/sycl/INTEL/pipes.hpp>

sycl/include/CL/sycl/INTEL/pipe_array.hpp

@@ -0,0 +1,57 @@
+//==--------------- pipe_array.hpp - SYCL pipe array --------*- C++ -*------==//
+//
+// 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
+
+#include <CL/sycl/INTEL/pipes.hpp>
+#include <utility>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace INTEL {
+
+namespace {
+template <size_t dim1, size_t... dims> struct VerifierDimLayer {
+  template <size_t idx1, size_t... idxs> struct VerifierIdxLayer {
+    static constexpr bool IsValid() {
+      return idx1 < dim1 &&
+             (VerifierDimLayer<dims...>::template VerifierIdxLayer<
+                 idxs...>::IsValid());
+    }
+  };
+};
+template <size_t dim> struct VerifierDimLayer<dim> {
+  template <size_t idx> struct VerifierIdxLayer {
+    static constexpr bool IsValid() { return idx < dim; }
+  };
+};
+} // namespace
+
+template <class Id, typename BaseTy, size_t depth, size_t... dims>
+struct PipeArray {
+  PipeArray() = delete;
+
+  template <size_t... idxs> struct StructId;
+
+  template <size_t... idxs> struct VerifyIndices {
+    static_assert(sizeof...(idxs) == sizeof...(dims),
+                  "Indexing into a PipeArray requires as many indices as "
+                  "dimensions of the PipeArray.");
+    static_assert(VerifierDimLayer<dims...>::template VerifierIdxLayer<
+                      idxs...>::IsValid(),
+                  "Index out of bounds");
+    using VerifiedPipe = pipe<StructId<idxs...>, BaseTy, depth>;
+  };
+
+  template <size_t... idxs>
+  using PipeAt = typename VerifyIndices<idxs...>::VerifiedPipe;
+};
+
+} // namespace INTEL
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/fpga_tests/compute_units.cpp

@@ -0,0 +1,81 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==------------ compute_units.cpp - SYCL FPGA compute units 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/INTEL/fpga_extensions.hpp>
+
+using namespace sycl;
+
+constexpr float kTestData = 555;
+
+constexpr size_t kComputeUnits = 5;
+using Pipes = INTEL::PipeArray<class MyPipe, float, 4, kComputeUnits + 1>;
+
+class source_kernel;
+class sink_kernel;
+
+template <std::size_t ID> class chain_kernel;
+
+// Write the first piece of data to the pipeline
+void SourceKernel(queue &q, float data) {
+
+  q.submit([&](handler &h) {
+    h.single_task<source_kernel>([=]() { Pipes::PipeAt<0>::write(data); });
+  });
+}
+
+// Grab the data out of the pipeline and return to host in out_data array
+void SinkKernel(queue &q, std::array<float, 1> &out_data) {
+
+  buffer<float, 1> out_buf(out_data.data(), 1);
+
+  q.submit([&](handler &h) {
+    auto out_accessor = out_buf.get_access<access::mode::write>(h);
+    h.single_task<sink_kernel>(
+        [=]() { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });
+  });
+}
+
+template <int TestNumber> int test_compute_units(queue q) {
+
+  std::array<float, 1> out_data = {0};
+
+  SourceKernel(q, kTestData);
+
+  INTEL::submit_compute_units<kComputeUnits, chain_kernel>(q, [=](auto ID) {
+    // read from id, not id-1 because the index_sequence starts from 0
+    float f = Pipes::PipeAt<ID>::read();
+    Pipes::PipeAt<ID + 1>::write(f);
+  });
+
+  SinkKernel(q, out_data);
+
+  if (out_data[0] != kTestData) {
+    std::cout << "Test: " << TestNumber << "\nResult mismatches " << out_data[0]
+              << " Vs expected " << kTestData << std::endl;
+    return -1;
+  }
+
+  return 0;
+}
+
+int main() {
+  cl::sycl::queue Queue;
+
+  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;
+  }
+
+  int Result = test_compute_units</*test number*/ 1>(Queue);
+  return Result;
+}