Skip to content

Commit 7e028e6

Browse files
jmitrevsjmitrevs
committed
pre-commit fixes
1 parent 454d556 commit 7e028e6

File tree

8 files changed

+59
-78
lines changed

8 files changed

+59
-78
lines changed

‎hls4ml/backends/oneapi/oneapi_backend.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,7 @@ def create_initial_config(self, part='Arria10', clock_period=5, io_type='io_para
153153
# TODO: add namespace
154154
'WriteTar': write_tar,
155155
}
156-
156+
157157
if 'use_bsp' in _:
158158
config['IS_BSP'] = True
159159

‎hls4ml/backends/oneapi/oneapi_types.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -200,7 +200,7 @@ def definition_cpp(self, name_suffix='', as_reference=True):
200200
return f'{self.name}{name_suffix}'
201201

202202
def declare_cpp(self, indent=''):
203-
streaming_beat_t = f"{self.pipe_name}BeatT";
203+
streaming_beat_t = f"{self.pipe_name}BeatT"
204204
lines = (
205205
f"{indent}class {self.pipe_id};\n"
206206
f"{indent}using {streaming_beat_t} = "

‎hls4ml/templates/oneapi/firmware/myproject.h‎

Lines changed: 27 additions & 39 deletions
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,9 @@ using PipeProps = decltype(sycl::ext::oneapi::experimental::properties(sycl::ext
1111
// Pipe properties for host pipes. Host pipes connect to the data source DMA and sink DMA.
1212
// They are connected to the first and the last layer to stream data into and out from the kernel.
1313
using HostPipePropertiesT = decltype(sycl::ext::oneapi::experimental::properties(
14-
sycl::ext::intel::experimental::ready_latency<0>,
15-
sycl::ext::intel::experimental::bits_per_symbol<8>,
16-
sycl::ext::intel::experimental::uses_valid<true>,
17-
sycl::ext::intel::experimental::first_symbol_in_high_order_bits<true>,
18-
sycl::ext::intel::experimental::protocol_avalon_streaming_uses_ready
19-
));
20-
14+
sycl::ext::intel::experimental::ready_latency<0>, sycl::ext::intel::experimental::bits_per_symbol<8>,
15+
sycl::ext::intel::experimental::uses_valid<true>, sycl::ext::intel::experimental::first_symbol_in_high_order_bits<true>,
16+
sycl::ext::intel::experimental::protocol_avalon_streaming_uses_ready));
2117

2218
namespace nnet {
2319

@@ -27,30 +23,27 @@ inline constexpr unsigned kInputBufferLocation = 0;
2723
inline constexpr unsigned kOutputBufferLocation = 1;
2824
#endif
2925

30-
// Implementation of a direct memory access kernel. Move data from source, convert,
26+
// Implementation of a direct memory access kernel. Move data from source, convert,
3127
// and send to the sink. Adaptive to SYCL HLS and hardware acceleration flow.
32-
template <class src_T, class dest_pipe>
33-
struct DMA_convert_data {
28+
template <class src_T, class dest_pipe> struct DMA_convert_data {
3429
#if !defined(IS_BSP)
35-
// When targeting a device family, we instantiate an Avalon Memory Mapped Host for
30+
// When targeting a device family, we instantiate an Avalon Memory Mapped Host for
3631
// data transaction between host and the DMA kernel during emulation and simulation.
37-
sycl::ext::oneapi::experimental::annotated_arg<src_T *,
38-
decltype(sycl::ext::oneapi::experimental::properties{
39-
sycl::ext::intel::experimental::latency<0>,
40-
sycl::ext::intel::experimental::dwidth<16>,
41-
sycl::ext::intel::experimental::buffer_location<kInputBufferLocation>,
42-
sycl::ext::intel::experimental::read_write_mode_read,
43-
sycl::ext::intel::experimental::wait_request_requested})>
32+
sycl::ext::oneapi::experimental::annotated_arg<
33+
src_T *,
34+
decltype(sycl::ext::oneapi::experimental::properties{
35+
sycl::ext::intel::experimental::latency<0>, sycl::ext::intel::experimental::dwidth<16>,
36+
sycl::ext::intel::experimental::buffer_location<kInputBufferLocation>,
37+
sycl::ext::intel::experimental::read_write_mode_read, sycl::ext::intel::experimental::wait_request_requested})>
4438
#else
4539
// When targeting oneAPI BSP, we can use USM pointer to access host memory.
4640
src_T *const
4741
#endif
4842
src;
4943
size_t num_iteration;
5044

51-
[[intel::kernel_args_restrict]]
52-
void operator()() const {
53-
45+
[[intel::kernel_args_restrict]] void operator()() const {
46+
5447
#if defined(IS_BSP)
5548
// Access data using host pointer.
5649
sycl::ext::intel::host_ptr<src_T> src_ptr(src);
@@ -64,8 +57,7 @@ struct DMA_convert_data {
6457
using DstDataType = typename nnet::ExtractDataType<PipeDataType>::value_type;
6558
constexpr auto dstTypeSize = std::tuple_size<DstDataType>{};
6659

67-
[[intel::fpga_register]]
68-
typename nnet::ExtractPipeType<dest_pipe>::value_type packet;
60+
[[intel::fpga_register]] typename nnet::ExtractPipeType<dest_pipe>::value_type packet;
6961

7062
// Keep sending data to the input layer and keep the kernels running.
7163
for (size_t i = 0; i < num_iteration; i++) {
@@ -82,28 +74,25 @@ struct DMA_convert_data {
8274
}
8375
};
8476

85-
// Symmetrical to the DMA_convert_data above, this DMA drains the output pipe and
77+
// Symmetrical to the DMA_convert_data above, this DMA drains the output pipe and
8678
// writes result to memory.
87-
template <class src_pipe, class dst_T>
88-
struct DMA_convert_data_back {
79+
template <class src_pipe, class dst_T> struct DMA_convert_data_back {
8980
#if !defined(IS_BSP)
9081
// Without BSP, instantiate an Avalon Memory Mapped Host to write to host.
91-
sycl::ext::oneapi::experimental::annotated_arg<dst_T *,
92-
decltype(sycl::ext::oneapi::experimental::properties{
93-
sycl::ext::intel::experimental::latency<0>,
94-
sycl::ext::intel::experimental::dwidth<16>,
95-
sycl::ext::intel::experimental::buffer_location<kOutputBufferLocation>,
96-
sycl::ext::intel::experimental::read_write_mode_write,
97-
sycl::ext::intel::experimental::wait_request_requested})>
82+
sycl::ext::oneapi::experimental::annotated_arg<
83+
dst_T *,
84+
decltype(sycl::ext::oneapi::experimental::properties{
85+
sycl::ext::intel::experimental::latency<0>, sycl::ext::intel::experimental::dwidth<16>,
86+
sycl::ext::intel::experimental::buffer_location<kOutputBufferLocation>,
87+
sycl::ext::intel::experimental::read_write_mode_write, sycl::ext::intel::experimental::wait_request_requested})>
9888
#else
9989
// USM pointer, otherwise.
10090
dst_T *const
10191
#endif
10292
dst;
10393
size_t num_iteration;
10494

105-
[[intel::kernel_args_restrict]]
106-
void operator()() const {
95+
[[intel::kernel_args_restrict]] void operator()() const {
10796
#if defined(IS_BSP)
10897
sycl::ext::intel::host_ptr<dst_T> dst_ptr(dst);
10998
#else
@@ -115,9 +104,8 @@ struct DMA_convert_data_back {
115104
using SrcDataType = typename nnet::ExtractDataType<PipeDataType>::value_type;
116105
constexpr auto srcTypeSize = std::tuple_size<SrcDataType>{};
117106

118-
[[intel::fpga_register]]
119-
typename nnet::ExtractPipeType<src_pipe>::value_type packet;
120-
107+
[[intel::fpga_register]] typename nnet::ExtractPipeType<src_pipe>::value_type packet;
108+
121109
// Drain the output pipe and write result to memory.
122110
for (size_t i = 0; i < num_iteration; i++) {
123111
packet = src_pipe::read();
@@ -129,7 +117,7 @@ struct DMA_convert_data_back {
129117
}
130118
};
131119

132-
} // namespace nnet
120+
} // namespace nnet
133121

134122
// Need to declare the input and output pipes
135123

‎hls4ml/templates/oneapi/firmware/nnet_utils/nnet_activation_stream.h‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -29,19 +29,17 @@ template <class data_pipe, class res_pipe, typename CONFIG_T> void linear_stream
2929
// *************************************************
3030
// ReLU Activation
3131
// *************************************************
32-
template <class data_pipe, class res_pipe, typename CONFIG_T>
33-
[[intel::use_stall_enable_clusters]] void relu_stream() {
32+
template <class data_pipe, class res_pipe, typename CONFIG_T> [[intel::use_stall_enable_clusters]] void relu_stream() {
3433
using namespace nnet;
3534
using ResT = typename ExtractDataType<typename ExtractPipeType<res_pipe>::value_type>::value_type;
3635
[[intel::fpga_register]] typename ExtractPipeType<res_pipe>::value_type out_data;
37-
36+
3837
bool keep_going = true;
3938
ReLUActLoop:
40-
[[intel::initiation_interval(1)]]
41-
while(keep_going) {
39+
[[intel::initiation_interval(1)]] while (keep_going) {
4240
for (int i = 0; i < CONFIG_T::n_in / std::tuple_size<ResT>{}; i++) {
4341
[[intel::fpga_register]] auto in_data = data_pipe::read();
44-
ReLUPackLoop:
42+
ReLUPackLoop:
4543
#pragma unroll
4644
for (int j = 0; j < std::tuple_size<ResT>{}; j++) {
4745
if (in_data.data[j] > 0)

‎hls4ml/templates/oneapi/firmware/nnet_utils/nnet_dense_stream.h‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,17 +11,17 @@ namespace nnet {
1111
// Computation is carried out in a while-1 loop as long as there is valid input.
1212
// The loop breaks when the end-of-packet signal is asserted by upstream task.
1313
template <class data_pipe, class res_pipe, typename CONFIG_T>
14-
[[intel::use_stall_enable_clusters]] void dense_resource_stream(const typename CONFIG_T::weight_t weights, const typename CONFIG_T::bias_t biases) {
14+
[[intel::use_stall_enable_clusters]] void dense_resource_stream(const typename CONFIG_T::weight_t weights,
15+
const typename CONFIG_T::bias_t biases) {
1516
using namespace nnet;
1617
using DataT = typename ExtractDataType<typename ExtractPipeType<data_pipe>::value_type>::value_type;
1718
using ResT = typename ExtractDataType<typename ExtractPipeType<res_pipe>::value_type>::value_type;
18-
19+
1920
[[intel::fpga_register]] typename ExtractPipeType<res_pipe>::value_type resbeat;
2021

2122
bool keep_going = true;
2223
bool did_read_input;
23-
[[intel::initiation_interval(1)]]
24-
while (keep_going) {
24+
[[intel::initiation_interval(1)]] while (keep_going) {
2525
did_read_input = false;
2626
[[intel::fpga_register]] auto databeat = data_pipe::read(did_read_input);
2727

‎hls4ml/templates/oneapi/firmware/nnet_utils/nnet_types.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
#include <tuple>
99
#include <utility>
1010

11-
#include <sycl/ext/intel/prototype/pipes_ext.hpp> // Streaming Beat and pipe properties.
11+
#include <sycl/ext/intel/prototype/pipes_ext.hpp> // Streaming Beat and pipe properties.
1212

1313
namespace nnet {
1414

‎hls4ml/templates/oneapi/myproject_test.cpp‎

Lines changed: 14 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -28,23 +28,19 @@ using sycl::ext::intel::experimental::property::usm::buffer_location;
2828
// Functions that reads input and prediction data from files.
2929
// Returns `true` if files are read successfully and not empty.
3030
// Returns `false` otherwise.
31-
bool prepare_data_from_file(
32-
std::string &fin_path,
33-
std::string &fpr_path,
34-
std::vector<std::vector<float>> &inputs,
35-
std::vector<std::vector<float>> &predictions
36-
) {
31+
bool prepare_data_from_file(std::string &fin_path, std::string &fpr_path, std::vector<std::vector<float>> &inputs,
32+
std::vector<std::vector<float>> &predictions) {
3733
// load input data from text file
3834
std::ifstream fin(fin_path.c_str());
3935
// load predictions from text file
4036
std::ifstream fpr(fpr_path.c_str());
41-
37+
4238
std::string iline;
4339
std::string pline;
4440

4541
if (fin.is_open() && fpr.is_open()) {
4642
size_t num_iterations = 0;
47-
43+
4844
// Prepare input data from file. Load predictions from file.
4945
for (; std::getline(fin, iline) && std::getline(fpr, pline); num_iterations++) {
5046
if (num_iterations % CHECKPOINT == 0) {
@@ -106,8 +102,8 @@ int main(int argc, char **argv) {
106102

107103
std::cout << "Running on device: " << device.get_info<sycl::info::device::name>().c_str() << std::endl;
108104

109-
std::string INPUT_FILE = "tb_data/tb_input_features.dat";
110-
std::string PRED_FILE = "tb_data/tb_output_predictions.dat";
105+
std::string INPUT_FILE = "tb_data/tb_input_features.dat";
106+
std::string PRED_FILE = "tb_data/tb_output_predictions.dat";
111107
std::string RESULTS_LOG = "tb_data/results.log";
112108
std::ofstream fout(RESULTS_LOG);
113109

@@ -138,10 +134,12 @@ int main(int argc, char **argv) {
138134
std::cerr << "ERROR: host allocation failed for output\n";
139135
fout.close();
140136
return 1;
141-
}
137+
}
142138
#else
143-
float *vals = sycl::malloc_shared<float>(kInputSz, q, sycl::property_list{buffer_location(nnet::kInputBufferLocation)});
144-
float *outputs = sycl::malloc_shared<float>(kOutputSz, q, sycl::property_list{buffer_location(nnet::kOutputBufferLocation)});
139+
float *vals =
140+
sycl::malloc_shared<float>(kInputSz, q, sycl::property_list{buffer_location(nnet::kInputBufferLocation)});
141+
float *outputs =
142+
sycl::malloc_shared<float>(kOutputSz, q, sycl::property_list{buffer_location(nnet::kOutputBufferLocation)});
145143
#endif
146144

147145
if (file_valid) {
@@ -175,7 +173,7 @@ int main(int argc, char **argv) {
175173
}
176174
} else {
177175
std::cout << "INFO: Unable to open input/predictions file, using default input with " << num_iterations
178-
<< " invocations." << std::endl;
176+
<< " invocations." << std::endl;
179177
q.single_task(MyProject{});
180178
// hls-fpga-machine-learning insert top-level-function
181179
// hls-fpga-machine-learning insert zero
@@ -195,12 +193,10 @@ int main(int argc, char **argv) {
195193
std::cout << "INFO: Saved inference results to file: " << RESULTS_LOG << std::endl;
196194
} catch (sycl::exception const &e) {
197195
// Catches exceptions in the host code.
198-
std::cerr << "Caught a SYCL host exception:\n"
199-
<< e.what() << "\n";
196+
std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";
200197

201198
// Most likely the runtime couldn't find FPGA hardware!
202-
if (e.code().value() == CL_DEVICE_NOT_FOUND)
203-
{
199+
if (e.code().value() == CL_DEVICE_NOT_FOUND) {
204200
std::cerr << "If you are targeting an FPGA, please ensure that your "
205201
"system has a correctly configured FPGA board.\n";
206202
std::cerr << "Run sys_check in the oneAPI root directory to verify.\n";

‎hls4ml/writer/oneapi_writer.py‎

Lines changed: 7 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -137,10 +137,7 @@ def write_project_cpp(self, model):
137137
elif '// hls-fpga-machine-learning read in' in line:
138138
newline = line
139139
if io_type == 'io_parallel':
140-
restartable_kernel_loop = (
141-
f"bool keep_going = true;\n\n"
142-
f"{indent}while (keep_going) {{\n"
143-
)
140+
restartable_kernel_loop = f"bool keep_going = true;\n\n" f"{indent}while (keep_going) {{\n"
144141
newline += indent + restartable_kernel_loop
145142
for inp in model_inputs:
146143
newline += indent * 2 + f'auto {inp.name}_beat = {inp.pipe_name}::read();\n'
@@ -202,7 +199,9 @@ def write_project_cpp(self, model):
202199
newline = indent + newline
203200
for out in model_outputs:
204201
out_beat = f"{out.name}_beat"
205-
newline += indent * 2 + f'typename nnet::ExtractPipeType<{out.pipe_name}>::value_type {out_beat};\n'
202+
newline += (
203+
indent * 2 + f'typename nnet::ExtractPipeType<{out.pipe_name}>::value_type {out_beat};\n'
204+
)
206205
newline += indent * 2 + f'{out_beat}.data = {out.name};\n'
207206
newline += indent * 2 + f'{out.pipe_name}::write({out_beat});\n'
208207
newline += indent * 2 + '// stops the kernel when the last input seen.\n'
@@ -422,7 +421,7 @@ def write_test_bench(self, model):
422421
f'{indent}const size_t kInputLayerSize = {model_inputs[0].size_cpp()};\n'
423422
f'{indent}const size_t kOutLayerSize = {model_outputs[0].size_cpp()};\n'
424423
)
425-
newline += insert_constant_lines;
424+
newline += insert_constant_lines
426425
elif '// hls-fpga-machine-learning insert zero' in line:
427426
newline = line
428427
inp = model_inputs[0]
@@ -445,8 +444,8 @@ def write_test_bench(self, model):
445444
elif '// hls-fpga-machine-learning convert output' in line:
446445
newline = line
447446
out = model_outputs[0]
448-
newline += \
449-
f'{indent}q.single_task(nnet::DMA_convert_data_back<{out.pipe_name}, float>{{outputs, num_iterations}}).wait();\n'
447+
newline += f'{indent}q.single_task(nnet::DMA_convert_data_back<{out.pipe_name}, float>'
448+
newline += '{outputs, num_iterations}).wait();\n'
450449
else:
451450
newline = line
452451

0 commit comments

Comments
 (0)