diff --git a/StencilStream/StencilExecutor.hpp b/StencilStream/StencilExecutor.hpp index 41b4545..824d662 100644 --- a/StencilStream/StencilExecutor.hpp +++ b/StencilStream/StencilExecutor.hpp @@ -61,8 +61,7 @@ template (1, burst_size / sizeof(T)); static constexpr uindex_t halo_radius = stencil_radius * pipeline_length; /** diff --git a/tests/src/res/FPGATransFunc.hpp b/tests/src/res/FPGATransFunc.hpp index e07d0d1..59806ac 100644 --- a/tests/src/res/FPGATransFunc.hpp +++ b/tests/src/res/FPGATransFunc.hpp @@ -13,17 +13,28 @@ #include #include +enum class CellStatus +{ + Normal, + Invalid, + Halo, +}; + +struct Cell +{ + stencil::index_t c; + stencil::index_t r; + stencil::index_t i_generation; + CellStatus status; +}; + template class FPGATransFunc { public: - // first element: column of the cell - // second element: row of the cell - // third element: generation - // fourth element: status (0: normal, 1: invalid, 2: halo) - using Cell = cl::sycl::vec; + using Cell = Cell; - static Cell halo() { return Cell(0, 0, 0, 2); } + static Cell halo() { return Cell{0, 0, 0, CellStatus::Halo}; } Cell operator()(stencil::Stencil const &stencil) const { @@ -40,17 +51,17 @@ class FPGATransFunc for (stencil::index_t r = -stencil::index_t(radius); r <= stencil::index_t(radius); r++) { Cell old_cell = stencil[stencil::ID(c, r)]; - is_valid &= (old_cell[0] == c + center_column && old_cell[1] == r + center_row && old_cell[2] == stencil.generation) || (old_cell[3] == 2); + is_valid &= (old_cell.c == c + center_column && old_cell.r == r + center_row && old_cell.i_generation == stencil.generation) || (old_cell.status == CellStatus::Halo); } } - if (new_cell[3] == 0) + if (new_cell.status == CellStatus::Normal) { if (!is_valid) { - new_cell[3] = 1; + new_cell.status = CellStatus::Invalid; } - new_cell[2] += 1; + new_cell.i_generation += 1; } return new_cell; diff --git a/tests/src/synthesis/main.cpp b/tests/src/synthesis/main.cpp index 3986ea3..8d76adb 100644 --- a/tests/src/synthesis/main.cpp +++ b/tests/src/synthesis/main.cpp @@ -25,7 +25,6 @@ const uindex_t grid_height = 2 * tile_height; const uindex_t burst_size = 1024; using TransFunc = FPGATransFunc; -using Cell = TransFunc::Cell; using Executor = StencilExecutor; void exception_handler(cl::sycl::exception_list exceptions) @@ -62,7 +61,7 @@ int main() { for (uindex_t r = 0; r < grid_height; r++) { - in_buffer_ac[c][r] = Cell(c, r, 0, 0); + in_buffer_ac[c][r] = Cell{index_t(c), index_t(r), 0, CellStatus::Normal}; } } } @@ -87,10 +86,10 @@ int main() { for (uindex_t r = 0; r < grid_height; r++) { - assert(out_buffer_ac[c][r][0] == c); - assert(out_buffer_ac[c][r][1] == r); - assert(out_buffer_ac[c][r][2] == 2 * pipeline_length); - assert(out_buffer_ac[c][r][3] == 0); + assert(out_buffer_ac[c][r].c == c); + assert(out_buffer_ac[c][r].r == r); + assert(out_buffer_ac[c][r].i_generation == 2 * pipeline_length); + assert(out_buffer_ac[c][r].status == CellStatus::Normal); } } } diff --git a/tests/src/units/ExecutionKernel.cpp b/tests/src/units/ExecutionKernel.cpp index ffcdece..eb83976 100644 --- a/tests/src/units/ExecutionKernel.cpp +++ b/tests/src/units/ExecutionKernel.cpp @@ -20,9 +20,9 @@ using namespace cl::sycl; void test_kernel(uindex_t n_generations) { using TransFunc = FPGATransFunc; - using in_pipe = HostPipe; - using out_pipe = HostPipe; - using TestExecutionKernel = ExecutionKernel; + using in_pipe = HostPipe; + using out_pipe = HostPipe; + using TestExecutionKernel = ExecutionKernel; for (index_t c = -halo_radius; c < index_t(halo_radius + tile_width); c++) { @@ -30,7 +30,7 @@ void test_kernel(uindex_t n_generations) { if (c >= index_t(0) && c < index_t(tile_width) && r >= index_t(0) && r < index_t(tile_height)) { - in_pipe::write(TransFunc::Cell(c, r, 0, 0)); + in_pipe::write(Cell{c, r, 0, CellStatus::Normal}); } else { @@ -41,7 +41,7 @@ void test_kernel(uindex_t n_generations) TestExecutionKernel(TransFunc(), 0, n_generations, 0, 0, tile_width, tile_height, TransFunc::halo())(); - buffer output_buffer(range<2>(tile_width, tile_height)); + buffer output_buffer(range<2>(tile_width, tile_height)); { auto output_buffer_ac = output_buffer.get_access(); @@ -62,11 +62,11 @@ void test_kernel(uindex_t n_generations) { for (uindex_t r = 1; r < tile_height; r++) { - TransFunc::Cell cell = output_buffer_ac[c][r]; - REQUIRE(cell[0] == c); - REQUIRE(cell[1] == r); - REQUIRE(cell[2] == n_generations); - REQUIRE(cell[3] == 0); + Cell cell = output_buffer_ac[c][r]; + REQUIRE(cell.c == c); + REQUIRE(cell.r == r); + REQUIRE(cell.i_generation == n_generations); + REQUIRE(cell.status == CellStatus::Normal); } } } diff --git a/tests/src/units/StencilExecutor.cpp b/tests/src/units/StencilExecutor.cpp index e0faa66..1459e71 100644 --- a/tests/src/units/StencilExecutor.cpp +++ b/tests/src/units/StencilExecutor.cpp @@ -18,7 +18,6 @@ using namespace stencil; using namespace cl::sycl; using TransFunc = FPGATransFunc; -using Cell = typename TransFunc::Cell; TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer)", "[StencilExecutor]") { @@ -29,7 +28,7 @@ TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer)", "[StencilExecu { for (uindex_t r = 0; r < grid_height; r++) { - in_buffer_ac[c][r] = Cell(c, r, 0, 0); + in_buffer_ac[c][r] = Cell{index_t(c), index_t(r), 0, CellStatus::Normal}; } } } @@ -46,10 +45,10 @@ TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer)", "[StencilExecu { for (uindex_t r = 0; r < grid_height; r++) { - REQUIRE(out_buffer_ac[c][r][0] == c); - REQUIRE(out_buffer_ac[c][r][1] == r); - REQUIRE(out_buffer_ac[c][r][2] == 0); - REQUIRE(out_buffer_ac[c][r][3] == 0); + REQUIRE(out_buffer_ac[c][r].c == c); + REQUIRE(out_buffer_ac[c][r].r == r); + REQUIRE(out_buffer_ac[c][r].i_generation == 0); + REQUIRE(out_buffer_ac[c][r].status == CellStatus::Normal); } } } @@ -66,7 +65,7 @@ TEST_CASE("StencilExecutor::run(uindex_t)", "[StencilExecutor]") { for (uindex_t r = 0; r < grid_height; r++) { - in_buffer_ac[c][r] = Cell(c, r, 0, 0); + in_buffer_ac[c][r] = Cell{index_t(c), index_t(r), 0, CellStatus::Normal}; } } } @@ -90,10 +89,10 @@ TEST_CASE("StencilExecutor::run(uindex_t)", "[StencilExecutor]") { for (uindex_t r = 0; r < grid_height; r++) { - REQUIRE(out_buffer_ac[c][r][0] == c); - REQUIRE(out_buffer_ac[c][r][1] == r); - REQUIRE(out_buffer_ac[c][r][2] == n_generations); - REQUIRE(out_buffer_ac[c][r][3] == 0); + REQUIRE(out_buffer_ac[c][r].c == c); + REQUIRE(out_buffer_ac[c][r].r == r); + REQUIRE(out_buffer_ac[c][r].i_generation == n_generations); + REQUIRE(out_buffer_ac[c][r].status == CellStatus::Normal); } } }