oneAPI Accelerator

hls4ml/backends/__init__.py

@@ -1,6 +1,7 @@
 from hls4ml.backends.backend import Backend, get_available_backends, get_backend, register_backend  # noqa: F401
 from hls4ml.backends.fpga.fpga_backend import FPGABackend  # noqa: F401
 from hls4ml.backends.oneapi.oneapi_backend import OneAPIBackend
+from hls4ml.backends.oneapi_accelerator.oneapi_accelerator_backend import OneAPIAcceleratorBackend
 from hls4ml.backends.quartus.quartus_backend import QuartusBackend
 from hls4ml.backends.symbolic.symbolic_backend import SymbolicExpressionBackend
 from hls4ml.backends.vivado.vivado_backend import VivadoBackend
@@ -18,3 +19,4 @@
 register_backend('Catapult', CatapultBackend)
 register_backend('SymbolicExpression', SymbolicExpressionBackend)
 register_backend('oneAPI', OneAPIBackend)
+register_backend('oneAPIAccelerator', OneAPIAcceleratorBackend)  # Can only be registered after oneAPI

hls4ml/backends/oneapi/oneapi_backend.py

@@ -17,8 +17,8 @@
 
 
 class OneAPIBackend(FPGABackend):
-    def __init__(self):
-        super().__init__('oneAPI')
+    def __init__(self, name='oneAPI'):  # the default name should be used in most cases
+        super().__init__(name)
         self._register_layer_attributes()
         self._register_flows()
 

hls4ml/backends/oneapi/oneapi_template.py

@@ -52,8 +52,8 @@ def _default_function_params(self, layer):
         params = self._default_params(layer)
         params['name'] = layer.name
         params['config'] = f'config{layer.index}'
-        params['input_pipe'] = layer.get_input_variable().pipe_name
-        params['output_pipe'] = layer.get_output_variable().pipe_name
+        params['input_pipe'] = layer.get_input_variable(layer.inputs[0]).pipe_name
+        params['output_pipe'] = layer.get_output_variable(layer.outputs[0]).pipe_name
 
         return params
 

hls4ml/backends/oneapi_accelerator/oneapi_accelerator_backend.py

@@ -0,0 +1,67 @@
+from hls4ml.backends import OneAPIBackend
+from hls4ml.model.flow import register_flow
+
+
+class OneAPIAcceleratorBackend(OneAPIBackend):
+    """
+    This is the backend to run oneAPI code on an accelerator using the oneAPI framework.
+    """
+
+    def __init__(self):
+        super().__init__(name='oneAPIAccelerator')
+
+    def _register_flows(self):
+        writer_passes = ['make_stamp', 'oneapiaccelerator:write_hls']
+        self._writer_flow = register_flow('write', writer_passes, requires=['oneapi:ip'], backend=self.name)
+
+        oneapi_types = [
+            'oneapiaccelerator:transform_types',
+            'oneapi:register_bram_weights',
+            'oneapi:apply_resource_strategy',
+            'oneapi:apply_winograd_kernel_transformation',
+        ]
+        oneapi_types_flow = register_flow('specific_types', oneapi_types, requires=['oneapi:init_layers'], backend=self.name)
+
+        streaming_passes = [
+            'oneapi:clone_output',
+            'oneapiaccelerator:extract_sideband',
+            'oneapiaccelerator:merge_sideband',
+        ]
+        streaming_flow = register_flow('streaming', streaming_passes, requires=['oneapi:init_layers'], backend=self.name)
+
+        template_flow = register_flow(
+            'apply_templates', self._get_layer_templates, requires=['oneapi:init_layers'], backend=self.name
+        )
+
+        accel_flow_requirements = [
+            'optimize',
+            'oneapi:init_layers',
+            streaming_flow,
+            'oneapi:quantization',
+            'oneapi:optimize',
+            oneapi_types_flow,
+            template_flow,
+        ]
+
+        accel_flow_requirements = list(filter(None, accel_flow_requirements))
+        self._default_flow = register_flow('accel', None, requires=accel_flow_requirements, backend=self.name)
+
+    def create_initial_config(
+        self, part, clock_period=5, hyperopt_handshake=False, io_type='io_parallel', write_tar=False, **_
+    ):
+        """Create initial configuration of the oneAPI backend.
+
+        Args:
+            part (str): The path to the board support package to be used. Can add :<board-variant>
+            clock_period (int, optional): The clock period in ns. Defaults to 5.
+            hyperopt_handshake (bool, optional): Should hyper-optimized handshaking be used? Defaults to False
+            io_type (str, optional): Type of implementation used. One of
+                'io_parallel' or 'io_stream'. Defaults to 'io_parallel'.
+            write_tar (bool, optional): If True, compresses the output directory into a .tar.gz file. Defaults to False.
+
+        Returns:
+            dict: initial configuration.
+        """
+        config = super().create_initial_config(part, clock_period, hyperopt_handshake, io_type, write_tar, **_)
+        config['UseOneAPIBSP'] = True
+        return config

hls4ml/backends/oneapi_accelerator/oneapi_accelerator_layers.py

@@ -0,0 +1,45 @@
+from hls4ml.model.attributes import Attribute, TypeAttribute
+from hls4ml.model.layers import Layer, register_layer
+from hls4ml.model.types import IntegerPrecisionType
+
+SIDEBAND_SHAPE = 2
+
+
+class SidebandExtraction(Layer):
+    """This layer extract the sideband and sends it to a different strem"""
+
+    _expected_attributes = [Attribute('n_in'), TypeAttribute('sideband_t', description='The type of the sidbands')]
+
+    def initialize(self):
+        inp = self.get_input_variable()
+        self.set_attr('n_in', inp.size())
+
+        # I think the order of these must be as stated because they each set the result_t type.
+        # We want the second one to be the actual result_t.
+        self.add_output_variable(
+            SIDEBAND_SHAPE,
+            out_name='sideband',
+            var_name='sideband_out',
+            type_name='sideband_t',
+            precision=IntegerPrecisionType(1, False),
+        )
+        self.set_attr('sideband_t', self.get_attr('sideband').type)  # need to manually set this, unlike result_t
+        self.add_output_variable(inp.shape, precision=inp.type.precision)
+
+
+class SidebandMerging(Layer):
+    """This layer gets the sideband from a different input and merges it"""
+
+    _expected_attributes = [
+        Attribute('n_in'),
+    ]
+
+    def initialize(self):
+        inp = self.get_input_variable()
+        self.set_attr('n_in', inp.size())
+        self.add_output_variable(inp.shape, precision=inp.type.precision)
+
+
+# register the layers
+register_layer('SidebandExtraction', SidebandExtraction)
+register_layer('SidebandMerging', SidebandMerging)

hls4ml/backends/oneapi_accelerator/oneapi_accelerator_types.py

@@ -0,0 +1,35 @@
+from hls4ml.backends.fpga.fpga_types import VariableDefinition
+from hls4ml.backends.oneapi.oneapi_types import AggregratedArrayVariableConverter
+
+
+# region InterfaceMemberVariable
+class OneAPIAcceleratorInterfaceVariableDefinition(VariableDefinition):
+    def definition_cpp(self, name_suffix='', as_reference=False):
+        if self.pragma and not isinstance(self.pragma, tuple):
+            return f'[[{self.pragma}]] {self.type.name} {self.name}{name_suffix}'
+        else:
+            return f'{self.type.name} {self.name}{name_suffix}'
+
+    # 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
+
+
+class OneAPIAcceleratorInterfaceVariableConverter(AggregratedArrayVariableConverter):
+    def __init__(self, type_converter):
+        super().__init__(
+            type_converter=type_converter,
+            prefix='OneAPIAccelerator',
+            definition_cls=OneAPIAcceleratorInterfaceVariableDefinition,
+        )

hls4ml/backends/oneapi_accelerator/passes/sideband_templates.py

@@ -0,0 +1,71 @@
+"""The sideband handling templates are needed for oneAPI accelerator when using io_stream.
+They are not used in io_paralle.
+"""
+
+from hls4ml.backends.oneapi.oneapi_template import StreamFunctionCallTemplate, TaskSequenceTemplate
+from hls4ml.backends.oneapi_accelerator.oneapi_accelerator_layers import SidebandExtraction, SidebandMerging
+from hls4ml.backends.template import FunctionCallTemplate, LayerConfigTemplate
+
+sideband_config_template = """struct config{index} : nnet::sideband_config {{
+    static constexpr unsigned n_in = {n_in};
+}};\n"""
+sideband_stream_function_template = '{name}.async();'
+sideband_extract_task_sequence_template = (
+    'task_sequence<nnet::extract_sideband_stream<{input_pipe}, {output_pipe}, {skip_pipe}, {config}>> {name};'
+)
+sideband_merge_task_sequence_template = (
+    'task_sequence<nnet::merge_sideband_stream<{input_pipe}, {output_pipe}, {skip_pipe}, {config}>> {name};'
+)
+sideband_include_list = ['nnet_utils/nnet_stream_beat.h']
+
+
+class SidebandConfigTemplate(LayerConfigTemplate):
+    def __init__(self):
+        super().__init__((SidebandExtraction, SidebandMerging))
+        self.template = sideband_config_template
+
+    def format(self, node):
+        params = self._default_config_params(node)
+        return self.template.format(**params)
+
+
+class SidebandFunctionTemplate(FunctionCallTemplate):
+    """Only used to add the include list"""
+
+    def __init__(self):
+        super().__init__((SidebandExtraction, SidebandMerging), include_header=sideband_include_list)
+
+    def format(self, node):
+        return ''
+
+
+class SidebandStreamFunctionTemplate(StreamFunctionCallTemplate):
+    def __init__(self):
+        super().__init__((SidebandExtraction, SidebandMerging))
+        self.template = sideband_stream_function_template
+
+    def format(self, node):
+        params = self._default_function_params(node)
+        return self.template.format(**params)
+
+
+class SidebandExtractionTaskSequenceTemplate(TaskSequenceTemplate):
+    def __init__(self):
+        super().__init__(SidebandExtraction)
+        self.template = sideband_extract_task_sequence_template
+
+    def format(self, node):
+        params = self._default_function_params(node)
+        params['skip_pipe'] = node.get_output_variable('sideband').pipe_name
+        return self.template.format(**params)
+
+
+class SidebandMergeTaskSequenceTemplate(TaskSequenceTemplate):
+    def __init__(self):
+        super().__init__(SidebandMerging)
+        self.template = sideband_merge_task_sequence_template
+
+    def format(self, node):
+        params = self._default_function_params(node)
+        params['skip_pipe'] = node.get_input_variable('sideband').pipe_name
+        return self.template.format(**params)

hls4ml/backends/oneapi_accelerator/passes/sidebands.py

@@ -0,0 +1,83 @@
+"""
+This file contains optimizers to add layers to extract and merge the sidebands. This is useful
+for the accelerator flow when using io_stream.
+
+Warning:  current version only works for network with single inputs and outputs.
+
+"""
+
+import warnings
+from collections import OrderedDict
+
+from hls4ml.backends.oneapi_accelerator.oneapi_accelerator_layers import SidebandExtraction, SidebandMerging
+from hls4ml.model.layers import Input
+from hls4ml.model.optimizer import OptimizerPass
+
+
+class ExtractSideband(OptimizerPass):
+    """Add a layer after the input to extract the sideband signals."""
+
+    def match(self, node):
+        if not (isinstance(node, Input) and node.model.config.get_config_value('IOType') == 'io_stream'):
+            return False
+        # now check that not already converted
+        output_nodes = node.get_output_nodes()
+        if len(output_nodes) == 1 and isinstance(output_nodes[0], SidebandExtraction):
+            # already transformed
+            return False
+        return True
+
+    def transform(self, model, node):
+        if len(model.inputs) > 1:
+            warnings.warn('Current sideband extraction scheme only tested on models with one input', stacklevel=1)
+
+        attributes = {'input_shape': node.get_attr('input_shape')}
+        new_node = model.make_node(
+            SidebandExtraction,
+            f'{node.name}_extract_sb',
+            attributes,
+            inputs=[node.outputs[0]],
+            outputs=[f'{node.name}_extract_sb', 'sideband'],
+        )
+        model.insert_node(new_node)
+        return True
+
+
+class MergeSideband(OptimizerPass):
+    """Add a layer after the last layer to merge the sideband signals."""
+
+    def match(self, node):
+        if node.model.config.get_config_value('IOType') == 'io_stream' and not isinstance(node, SidebandMerging):
+            for node_out in node.outputs:
+                if node_out in node.model.outputs:  # if the node output is a model output
+                    return True
+        return False
+
+    def transform(self, model, node):
+        if len(model.outputs) > 1:
+            warnings.warn('Current sideband extraction scheme only tested on models with one output', stacklevel=1)
+
+        attributes = {}
+
+        inputs = [out for out in node.outputs if out in model.outputs]
+
+        if len(inputs) != 1:
+            raise RuntimeError('Unsupported number of outputs found')
+
+        inputs.append('sideband')
+
+        new_name = f'{node.name}_merge_sb'
+        new_node = model.make_node(SidebandMerging, new_name, attributes, inputs=inputs)
+
+        # note that model.insert_node fails here because of the two input nodes, so using a custom version below
+        model.outputs[0] = new_name
+
+        new_graph = OrderedDict()
+        for k, v in model.graph.items():
+            new_graph[k] = v
+            if k == node.name:
+                new_graph[new_node.name] = new_node
+
+        model.graph = new_graph
+
+        return True