FPGA: Remove DataBundle from convolution2d code sample, update gasket IP (#2310)

Simplify convolution2d code sample by removing the proprietary DataBundle header and replace it with std::array for simplicity and portability.

Author: whitepau

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/convolution2d/README.md

@@ -118,7 +118,7 @@ struct Convolution2d {
         myLineBuffer(rows, cols);
 
     bool keep_going = true;
-    bool bypass = true;
+    bool bypass = false;
 
     [[intel::initiation_interval(1)]]  // NO-FORMAT: Attribute
     while (keep_going) {
@@ -196,9 +196,9 @@ conv2d::PixelType ConvolutionFunction(
       // handle the case where the center of the window is at the image edge.
       // In this design, simply 'reflect' pixels that are already in the
       // window.
-      SaturateWindowCoordinates(w_row, w_col,  //
-                                row, col,      //
-                                rows, cols,    //
+      SaturateWindowCoordinates(w_row, w_col,  // NO-FORMAT: Alignment
+                                row, col,      // NO-FORMAT: Alignment
+                                rows, cols,    // NO-FORMAT: Alignment
                                 r_select, c_select);
       conv2d::PixelType pixel =
           buffer[c_select + r_select * conv2d::kWindowSize];
@@ -245,7 +245,7 @@ For convenience, you may use the header file included in `quartus_project_files/
 
 ### Test bench utilities
 
-In this design, pipes are used to transfer data between kernels, and between the design and the testbench (host code). An aggregate type (`fpga_tools::DataBundle`) is used to allow multiple pixels to transfer in one clock cycle. To help with this, this reference design uses the `WriteFrameToPipe()` and `ReadFrameFromPipe()` functions, which are defined in `include/vvp_stream_adapters.hpp`. 
+In this design, pipes are used to transfer data between kernels, and between the design and the testbench (host code). An aggregate type (`std::array`) is used to allow multiple pixels to transfer in one clock cycle. To help with this, this reference design uses the `WriteFrameToPipe()` and `ReadFrameFromPipe()` functions, which are defined in `include/vvp_stream_adapters.hpp`. 
 
 `WriteFrameToPipe()` writes the contents of an array of pixels *into* a SYCL pipe that can be consumed by a oneAPI kernel. It detects the parameterization of the aggregate type used by the pipe, and groups pixels together accordingly. It also generates start-of-packet and end-of-packet sideband signals like a VVP FPGA IP would, so you can test that your IP can interface with other IPs that use the VVP standard. 
 
@@ -260,7 +260,11 @@ The following diagram illustrates how these functions adapt image data to pipes,
 The following code snippet demonstrates how you can use these functions to populate a pipe with image data before invoking a kernel, and how you can parse the output.
 
 ```c++
-bool TestTinyFrameOnStencil(sycl::queue q) {
+bool TestTinyFrameOnStencil(sycl::queue q, bool print_debug_info) {
+  std::cout << "\n**********************************\n"
+            << "Check Tiny frame... "
+            << "\n**********************************\n"
+            << std::endl;
   constexpr int rows_small = 3;
   constexpr int cols_small = 8;
 
@@ -274,28 +278,22 @@ bool TestTinyFrameOnStencil(sycl::queue q) {
   vvp_stream_adapters::WriteFrameToPipe<InputImageStreamGrey>(
       q, rows_small, cols_small, grey_pixels_in);
 
-  // extra pixels to flush out the FIFO
+  // add extra pixels to flush out the FIFO after all image frames
+  // have been added
   int dummy_pixels = cols_small * conv2d::kWindowSize;
   vvp_stream_adapters::WriteDummyPixelsToPipe<InputImageStreamGrey>(
       q, dummy_pixels, (uint16_t)15);
 
-  // disable bypass, since it's on by default
-  BypassCSR::write(q, false);
-
-  // Make sure that there is no 'true' still sitting in the 'stop' register from
-  // the last time the kernel was stopped
-  StopCSR::write(q, false);
-
   sycl::event e = q.single_task<ID_Convolution2d>(
       Convolution2d<InputImageStreamGrey, OutputImageStreamGrey>{
           (int)rows_small, (int)cols_small, identity_coeffs});
 
   conv2d::PixelType grey_pixels_out[pixels_count];
   bool sidebands_ok;
-  int defective_frames;
+  int parsed_frames;
   vvp_stream_adapters::ReadFrameFromPipe<OutputImageStreamGrey>(
-      q, rows_small, cols_small, grey_pixels_out, sidebands_ok,
-      defective_frames);
+      q, rows_small, cols_small, grey_pixels_out, sidebands_ok, parsed_frames,
+      print_debug_info);
 
   bool pixels_match = true;
   for (int i = 0; i < pixels_count; i++) {

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/convolution2d/assets/testbench_architecture.svg

@@ -5,8 +5,9 @@
 
 #pragma once
 
+#include <array>
+
 #include "comparisons.hpp"
-#include "data_bundle.hpp"
 #include "shift_reg.hpp"
 
 namespace line_buffer_2d {
@@ -29,10 +30,8 @@ template <typename PixelTypeIn, typename PixelTypeOut, short kStencilSize,
 class LineBuffer2d {
  public:
   // types used by LineBuffer2d
-  using LineBufferDataBundleIn =
-      fpga_tools::DataBundle<PixelTypeIn, kParallelPixels>;
-  using LineBufferDataBundleOut =
-      fpga_tools::DataBundle<PixelTypeOut, kParallelPixels>;
+  using LineBufferDataBundleIn = std::array<PixelTypeIn, kParallelPixels>;
+  using LineBufferDataBundleOut = std::array<PixelTypeOut, kParallelPixels>;
 
   // public members
   [[intel::fpga_register]]  // NO-FORMAT: Attribute
@@ -44,8 +43,7 @@ class LineBuffer2d {
  private:
   // types used internally
   using PixelWithSignals = PixelWithSignals_<PixelTypeIn>;
-  using BundledPixels =
-      fpga_tools::DataBundle<PixelWithSignals, kParallelPixels>;
+  using BundledPixels = std::array<PixelWithSignals, kParallelPixels>;
   constexpr static short kRowWriteInit = (short)(0 - kStencilSize);
   constexpr static short kColWriteInit = (short)(0 - kStencilSize);
 
@@ -80,7 +78,7 @@ class LineBuffer2d {
   constexpr static short kPreBufferSize = kParallelPixels + kBufferOffset;
 
   [[intel::fpga_register]]  // NO-FORMAT: Attribute
-  fpga_tools::DataBundle<PixelWithSignals, kPreBufferSize>
+  fpga_tools::ShiftReg<PixelWithSignals, kPreBufferSize>
       pre_buffer;
 
   // separate the loop bound calculation so loop iterations are easier to
@@ -171,8 +169,10 @@ class LineBuffer2d {
     // grab the first `kParallelPixels` samples to push into the stencil
     [[intel::fpga_register]]  // NO-FORMAT: Attribute
     BundledPixels input_val;
-    input_val.template ShiftMultiVals<kParallelPixels, kPreBufferSize>(
-        pre_buffer);
+#pragma unroll
+    for (int i = 0; i < kParallelPixels; i++) {
+      input_val[i] = pre_buffer[i];
+    }
 
     [[intel::fpga_register]]  // NO-FORMAT: Attribute
     BundledPixels pixel_column[kStencilSize];

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/convolution2d/include/data_bundle.hpp

@@ -6,7 +6,8 @@
 #ifndef __SHIFT_REG_HPP__
 #define __SHIFT_REG_HPP__
 
-#include "data_bundle.hpp"
+#include <array>
+
 #include "unrolled_loop.hpp"
 
 namespace fpga_tools {
@@ -57,7 +58,7 @@ class ShiftReg {
   }
 
   template <size_t kShiftAmt>
-  void ShiftMultiVals(DataBundle<T, kShiftAmt> in) {
+  void ShiftMultiVals(std::array<T, kShiftAmt> in) {
     fpga_tools::UnrolledLoop<0, (kRegDepth - kShiftAmt)>(
         [&](int i) { registers[i] = registers[i + kShiftAmt]; });
 
@@ -129,7 +130,7 @@ class ShiftReg2d {
   }
 
   template <size_t kShiftAmt>
-  void ShiftCols(DataBundle<T, kShiftAmt> in[kRegRows]) {
+  void ShiftCols(std::array<T, kShiftAmt> in[kRegRows]) {
     fpga_tools::UnrolledLoop<0, kRegRows>(
         [&](int i) { registers[i].template ShiftMultiVals<kShiftAmt>(in[i]); });
   }