oneAPI backend update: kernel and layer optimizations

docs/api/configuration.rst

@@ -101,6 +101,20 @@ Finally, one then uses the configuration to create an hls model:
         backend='Vitis'
     )
 
+To target an oneAPI Board Support Package (BSP) enabled FPGA for offload acceleration, you can specify the ``part`` argument to be the path to your BSP and the BSP variant. Then, set ``use_oneapi_bsp=True``.
+
+.. code-block:: python
+
+  hls_model = hls4ml.converters.convert_from_keras_model(
+        model,
+        hls_config=config,
+        output_dir="my_project_dir",
+        io_type="io_parallel",
+        backend="oneAPI",
+        part="/path/to/my/bsp:bsp_variant",
+        use_oneapi_bsp=True
+    )
+
 See :py:class:`~hls4ml.converters.convert_from_keras_model` for more information on the various options. Similar functions exist for ONNX and PyTorch.
 
 ----
@@ -132,6 +146,9 @@ It looks like this:
    ClockPeriod: 5
    IOType: io_parallel # options: io_parallel/io_stream
 
+   # oneAPI Offload Acceleration flag.
+   UseOneAPIBSP: True
+
    HLSConfig:
      Model:
        Precision: fixed<16,6>
@@ -168,6 +185,10 @@ For Vivado backend the options are:
   * **PipelineInterval**\ : Optionally override the desired initiation interval of the design. Only valid in combination with "pipeline" style. If unspecified, it is left to the compiler to decide, ideally matching the largest reuse factor of the network.
   * **Precision**\ : this defines the precision of your inputs, outputs, weights and biases. It is denoted by ``fixed<X,Y>``\ , where ``Y`` is the number of bits representing the signed number above the binary point (i.e. the integer part), and ``X`` is the total number of bits. Additionally, integers in the type (\ ``int<N>``\ , where ``N`` is a bit-size from 1 to 1024) can also be used. The format follows ``ap_fixed`` and ``ap_int`` conventions. You have a chance to further configure this more finely with per-layer configuration described below. In the per-layer configuration (but not globally) one can also use ``'auto'`` precision.
 
+For oneaAPI the options are similar but also include:
+
+* **UseOneAPIBSP**\ : path to the oneAPI Board Support Package (and the BSP variant) to enable offload acceleration with an Altera FPGA. This is only needed if you are using oneAPI in the accelerator style.
+
 2.2 Per-Layer Configuration
 ---------------------------
 

hls4ml/backends/oneapi/oneapi_backend.py

@@ -130,11 +130,13 @@ def get_default_flow(self):
     def get_writer_flow(self):
         return self._writer_flow
 
-    def create_initial_config(self, part='Arria10', clock_period=5, io_type='io_parallel', write_tar=False, **_):
+    def create_initial_config(
+        self, part='Agilex7', clock_period=5, io_type='io_parallel', write_tar=False, use_oneapi_bsp=False, **_
+    ):
         """Create initial configuration of the oneAPI backend.
 
         Args:
-            part (str, optional): The FPGA part to be used. Defaults to 'Arria10'.
+            part (str, optional): The FPGA part to be used. Defaults to 'Agilex7'.
             clock_period (int, optional): The clock period. Defaults to 5.
             io_type (str, optional): Type of implementation used. One of
                 'io_parallel' or 'io_stream'. Defaults to 'io_parallel'.
@@ -146,15 +148,15 @@ def create_initial_config(self, part='Arria10', clock_period=5, io_type='io_para
 
         config = {}
 
-        config['Part'] = part if part is not None else 'Arria10'
+        config['Part'] = part if part is not None else 'Agilex7'
         config['ClockPeriod'] = clock_period
         config['IOType'] = io_type
         config['HLSConfig'] = {}
         config['WriterConfig'] = {
             # TODO:  add namespace
             'WriteTar': write_tar,
         }
-
+        config['UseOneAPIBSP'] = use_oneapi_bsp
         return config
 
     def compile(self, model):

hls4ml/backends/oneapi/oneapi_types.py

@@ -170,11 +170,18 @@ def definition_cpp(self, name_suffix='', as_reference=False):
         else:
             return f'{self.type.name} {self.name}{name_suffix}'
 
-    def declare_cpp(self, pipe_min_size=0, indent=''):
-        lines = indent + f'class {self.pipe_id};\n'
-        lines += indent + (
-            f'using {self.pipe_name} = sycl::ext::intel::experimental::pipe<{self.pipe_id}, '
-            + f'{self.type.name}, {pipe_min_size}, PipeProps>;\n'
+    # Updated pipe min size to be 32 for simulation.
+    def declare_cpp(self, pipe_min_size=32, indent=''):
+        # Updated to use streaming beat for restartable streaming kernel.
+        # Streaming beat is a wrapper type of the actual type with sideband control signals.
+        # Syntax: using BeatT = sycl::ext::intel::experimental::StreamingBeat<DataT, eop, empty>;
+        streaming_beat_t = f"{self.pipe_name}BeatT"
+        lines = (
+            f"{indent}class {self.pipe_id};\n"
+            f"{indent}using {streaming_beat_t} = "
+            f"sycl::ext::intel::experimental::StreamingBeat<{self.type.name}, true, true>;\n"
+            f"{indent}using {self.pipe_name} = sycl::ext::intel::experimental::pipe<"
+            f"{self.pipe_id}, {streaming_beat_t}, {pipe_min_size}, HostPipePropertiesT>;\n"
         )
         return lines
 
@@ -193,10 +200,13 @@ def definition_cpp(self, name_suffix='', as_reference=True):
         return f'{self.name}{name_suffix}'
 
     def declare_cpp(self, indent=''):
-        lines = indent + f'class {self.pipe_id};\n'
-        lines += indent + (
-            f'using {self.pipe_name} = sycl::ext::intel::experimental::pipe<{self.pipe_id}, '
-            + f'{self.type.name}, {self.pragma[-1]}>;\n'
+        streaming_beat_t = f"{self.pipe_name}BeatT"
+        lines = (
+            f"{indent}class {self.pipe_id};\n"
+            f"{indent}using {streaming_beat_t} = "
+            f"sycl::ext::intel::experimental::StreamingBeat<{self.type.name}, true, true>;\n"
+            f"{indent}using {self.pipe_name} = "
+            f"sycl::ext::intel::experimental::pipe<{self.pipe_id}, {streaming_beat_t}, {self.pragma[-1]}>;\n"
         )
         return lines
 

hls4ml/templates/oneapi/CMakeLists.txt

@@ -38,15 +38,20 @@ set(LIBRARY_NAME myproject-${LIB_STAMP})
 # You can also specify a device family (E.g. "Arria10" or "Stratix10") or a
 # specific part number (E.g. "10AS066N3F40E2SG") to generate a standalone IP.
 if(NOT DEFINED FPGA_DEVICE)
-    set(FPGA_DEVICE "Arria10")
+    set(FPGA_DEVICE "Agilex7")
+    set(BSP_FLAG "")
 endif()
 
+# Set the target to a BSP if we target an actual accelerator board.
+# hls-fpga-machine-learning insert oneapi_bsp_cmake_flag
+
 # Use cmake -DUSER_FPGA_FLAGS=<flags> to set extra flags for FPGA backend
 # compilation.
-set(USER_FPGA_FLAGS -Wno-unused-label ${USER_FPGA_FLAGS})
+# -Xsoptimize=latency Turns off the hyper-optimized handshake
+set(USER_FPGA_FLAGS -Wno-unused-label;${USER_FPGA_FLAGS};-Xsoptimize=latency)
 
 # Use cmake -DUSER_FLAGS=<flags> to set extra flags for general compilation.
-set(USER_FLAGS -Wno-unused-label -fconstexpr-steps=134217728 ${USER_FLAGS})
+set(USER_FLAGS -Wno-unused-label -fconstexpr-steps=134217728 ${USER_FLAGS} ${BSP_FLAG})
 
 # Use cmake -DUSER_INCLUDE_PATHS=<paths> to set extra paths for general
 # compilation.

hls4ml/templates/oneapi/firmware/myproject.h

@@ -5,8 +5,12 @@
 
 // This file defines the interface to the kernel
 
-// currently this is fixed
-using PipeProps = decltype(sycl::ext::oneapi::experimental::properties(sycl::ext::intel::experimental::ready_latency<0>));
+// Pipe properties for host pipes. Host pipes connect to the data source DMA and sink DMA.
+// They are connected to the first and the last layer to stream data into and out from the kernel.
+using HostPipePropertiesT = decltype(sycl::ext::oneapi::experimental::properties(
+    sycl::ext::intel::experimental::ready_latency<0>, sycl::ext::intel::experimental::bits_per_symbol<16>,
+    sycl::ext::intel::experimental::uses_valid<true>, sycl::ext::intel::experimental::first_symbol_in_high_order_bits<true>,
+    sycl::ext::intel::experimental::protocol_avalon_streaming_uses_ready));
 
 // Need to declare the input and output pipes
 
@@ -16,12 +20,15 @@ using PipeProps = decltype(sycl::ext::oneapi::experimental::properties(sycl::ext
 class MyProjectID;
 
 struct MyProject {
-
+#ifndef IS_BSP
     // kernel property method to config invocation interface
     auto get(sycl::ext::oneapi::experimental::properties_tag) {
         return sycl::ext::oneapi::experimental::properties{sycl::ext::intel::experimental::streaming_interface<>,
                                                            sycl::ext::intel::experimental::pipelined<>};
     }
+#else
+    // kernel properties and pipelining is not supported in BSP (accelerator style).
+#endif
 
     SYCL_EXTERNAL void operator()() const;
 };