Altera HLS compatibility

hls4ml/backends/oneapi/oneapi_backend.py

@@ -210,7 +210,11 @@ def build(self, model, build_type='fpga_emu', run=False):
         try:
             subprocess.run('which icpx', shell=True, cwd=builddir, check=True)
         except subprocess.CalledProcessError:
-            raise RuntimeError('Could not find icpx. Please configure oneAPI appropriately')
+            print('Could not find icpx, trying ahls instead.')
+            try:
+                subprocess.run('which ahls', shell=True, cwd=builddir, check=True)
+            except subprocess.CalledProcessError:
+                raise RuntimeError('Could not find icpx or ahls. Please configure oneAPI appropriately')
         subprocess.run('cmake ..', shell=True, cwd=builddir, check=True)
         subprocess.run(f'make {build_type}', shell=True, cwd=builddir, check=True)
 

hls4ml/backends/oneapi/oneapi_types.py

@@ -2,6 +2,8 @@
 This package includes oneAPI-specific customizations to the variable types
 """
 
+import shutil
+
 import numpy as np
 
 from hls4ml.backends.fpga.fpga_types import (
@@ -180,9 +182,10 @@ def definition_cpp(self, name_suffix='', as_reference=False):
             return f'{self.type.name} {self.name}{name_suffix}'
 
     def declare_cpp(self, pipe_min_size=0, indent=''):
+        compiler_name = 'altera' if shutil.which('ahls') else 'intel'
         lines = indent + f'class {self.pipe_id};\n'
         lines += indent + (
-            f'using {self.pipe_name} = sycl::ext::intel::experimental::pipe<{self.pipe_id}, '
+            f'using {self.pipe_name} = sycl::ext::{compiler_name}::experimental::pipe<{self.pipe_id}, '
             + f'{self.type.name}, {pipe_min_size}, PipeProps>;\n'
         )
         return lines
@@ -202,9 +205,10 @@ def definition_cpp(self, name_suffix='', as_reference=True):
         return f'{self.name}{name_suffix}'
 
     def declare_cpp(self, indent=''):
+        compiler_name = 'altera' if shutil.which('ahls') else 'intel'
         lines = indent + f'class {self.pipe_id};\n'
         lines += indent + (
-            f'using {self.pipe_name} = sycl::ext::intel::experimental::pipe<{self.pipe_id}, '
+            f'using {self.pipe_name} = sycl::ext::{compiler_name}::experimental::pipe<{self.pipe_id}, '
             + f'{self.type.name}, {self.pragma[-1]}>;\n'
         )
         return lines

hls4ml/templates/oneapi/CMakeLists.txt

@@ -1,7 +1,12 @@
 # Direct CMake to use icpx rather than the default C++ compiler/linker on Linux
 # and icx-cl on Windows
 if(UNIX)
-    set(CMAKE_CXX_COMPILER icpx)
+    find_program(CXX_COMPILER NAMES icpx ahls)
+    if(NOT CXX_COMPILER)
+        message(FATAL_ERROR "Neither icpx nor ahls was found in PATH.")
+    endif()
+    set(CMAKE_CXX_COMPILER "${CXX_COMPILER}")
+    message(STATUS "Using C++ compiler: ${CXX_COMPILER}")
 else() # Windows
     include (CMakeForceCompiler)
     CMAKE_FORCE_CXX_COMPILER (icx-cl IntelDPCPP)
@@ -12,6 +17,10 @@ cmake_minimum_required (VERSION 3.7.2)
 
 project(myproject CXX)
 
+if(CXX_COMPILER MATCHES "ahls$")
+    add_compile_definitions(AHLS)
+endif()
+
 set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
 set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})

hls4ml/templates/oneapi/firmware/defines.h

@@ -1,9 +1,15 @@
 #ifndef DEFINES_H_
 #define DEFINES_H_
 
+#ifdef AHLS
+#include <sycl/ext/altera/ac_types/ac_fixed.hpp>
+#include <sycl/ext/altera/ac_types/ac_int.hpp>
+#include <sycl/ext/altera/fpga_extensions.hpp>
+#else
 #include <sycl/ext/intel/ac_types/ac_fixed.hpp>
 #include <sycl/ext/intel/ac_types/ac_int.hpp>
 #include <sycl/ext/intel/fpga_extensions.hpp>
+#endif
 #include <sycl/sycl.hpp>
 
 // Include nnet::array - a custom array-like struct, mainly used with io_stream

hls4ml/templates/oneapi/firmware/myproject.cpp

@@ -1,13 +1,21 @@
 #include "myproject.h"
 #include "parameters.h"
+#ifdef AHLS
+#include <sycl/ext/altera/experimental/task_sequence.hpp>
+#else
 #include <sycl/ext/intel/experimental/task_sequence.hpp>
+#endif
 
 // hls-fpga-machine-learning insert weights
 
 // The inter-task pipes need to be declared in the global scope
 // hls-fpga-machine-learning insert inter-task pipes
 
+#ifdef AHLS
+using sycl::ext::altera::experimental::task_sequence;
+#else
 using sycl::ext::intel::experimental::task_sequence;
+#endif
 
 void MyProject::operator()() const {
     // ****************************************

hls4ml/templates/oneapi/firmware/myproject.h

@@ -6,7 +6,14 @@
 // 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>));
+
+using PipeProps = decltype(sycl::ext::oneapi::experimental::properties(
+#ifdef AHLS
+    sycl::ext::altera::experimental::ready_latency<0>
+#else
+    sycl::ext::intel::experimental::ready_latency<0>
+#endif
+    ));
 
 // Need to declare the input and output pipes
 
@@ -19,8 +26,13 @@ struct MyProject {
 
     // 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<>};
+        return sycl::ext::oneapi::experimental::properties{
+#ifdef AHLS
+            sycl::ext::altera::experimental::streaming_interface<>, sycl::ext::altera::experimental::pipelined<>
+#else
+            sycl::ext::intel::experimental::streaming_interface<>, sycl::ext::intel::experimental::pipelined<>
+#endif
+        };
     }
 
     SYCL_EXTERNAL void operator()() const;

hls4ml/templates/oneapi/firmware/nnet_utils/nnet_common.h

@@ -2,9 +2,16 @@
 #define NNET_COMMON_H_
 
 #include "nnet_helpers.h"
+
+#ifdef AHLS
+#include <sycl/ext/altera/ac_types/ac_fixed.hpp>
+#include <sycl/ext/altera/ac_types/ac_fixed_math.hpp>
+#include <sycl/ext/altera/ac_types/ac_int.hpp>
+#else
 #include <sycl/ext/intel/ac_types/ac_fixed.hpp>
 #include <sycl/ext/intel/ac_types/ac_fixed_math.hpp>
 #include <sycl/ext/intel/ac_types/ac_int.hpp>
+#endif
 
 typedef ac_fixed<16, 6> table_default_t;
 

hls4ml/templates/oneapi/myproject_bridge.cpp

@@ -56,7 +56,11 @@ void collect_trace_output(struct trace_data *c_trace_outputs) {
 void myproject_float(
     // hls-fpga-machine-learning insert header #float
 ) {
+#ifdef AHLS
+    auto selector = sycl::ext::altera::fpga_emulator_selector_v;
+#else
     auto selector = sycl::ext::intel::fpga_emulator_selector_v;
+#endif
     static sycl::queue q(selector, fpga_tools::exception_handler, sycl::property::queue::enable_profiling{});
 
     // hls-fpga-machine-learning insert wrapper #float
@@ -67,7 +71,11 @@ void myproject_float(
 void myproject_double(
     // hls-fpga-machine-learning insert header #double
 ) {
+#ifdef AHLS
+    auto selector = sycl::ext::altera::fpga_emulator_selector_v;
+#else
     auto selector = sycl::ext::intel::fpga_emulator_selector_v;
+#endif
     static sycl::queue q(selector, fpga_tools::exception_handler, sycl::property::queue::enable_profiling{});
 
     // hls-fpga-machine-learning insert wrapper #double

hls4ml/templates/oneapi/myproject_test.cpp

@@ -9,9 +9,13 @@
 #include "firmware/myproject.h"
 #include "firmware/parameters.h"
 
+#ifdef AHLS
+#include <sycl/ext/altera/fpga_extensions.hpp>
+#else
 #include <sycl/ext/intel/fpga_extensions.hpp>
+#endif
 
-#if (__INTEL_CLANG_COMPILER < 20250000)
+#if (__INTEL_CLANG_COMPILER < 20250000) && !defined(AHLS)
 #include <sycl/ext/intel/prototype/interfaces.hpp>
 #endif
 
@@ -22,12 +26,22 @@
 
 int main(int argc, char **argv) {
 
+#ifdef AHLS
+#if FPGA_SIMULATOR
+    auto selector = sycl::ext::altera::fpga_simulator_selector_v;
+#elif FPGA_HARDWARE
+    auto selector = sycl::ext::altera::fpga_selector_v;
+#else // #if FPGA_EMULATOR
+    auto selector = sycl::ext::altera::fpga_emulator_selector_v;
+#endif
+#else
 #if FPGA_SIMULATOR
     auto selector = sycl::ext::intel::fpga_simulator_selector_v;
 #elif FPGA_HARDWARE
     auto selector = sycl::ext::intel::fpga_selector_v;
 #else // #if FPGA_EMULATOR
     auto selector = sycl::ext::intel::fpga_emulator_selector_v;
+#endif
 #endif
 
     sycl::queue q(selector, fpga_tools::exception_handler, sycl::property::queue::enable_profiling{});