Skip to content

Commit

Permalink
Removing the restriction for cell sizes to divide the burst size
Browse files Browse the repository at this point in the history
  • Loading branch information
@JOOpdenhoevel
JOOpdenhoevel committed Aug 13, 2021
1 parent 370b794 commit 0e2bfa5
Show file tree
Hide file tree
Showing 5 changed files with 47 additions and 39 deletions.
3 changes: 1 addition & 2 deletions StencilStream/StencilExecutor.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -61,8 +61,7 @@ template <typename T, uindex_t stencil_radius, typename TransFunc, uindex_t pipe
class StencilExecutor
{
public:
static_assert(burst_size % sizeof(T) == 0);
static constexpr uindex_t burst_length = burst_size / sizeof(T);
static constexpr uindex_t burst_length = std::min<uindex_t>(1, burst_size / sizeof(T));
static constexpr uindex_t halo_radius = stencil_radius * pipeline_length;

/**
Expand Down
31 changes: 21 additions & 10 deletions tests/src/res/FPGATransFunc.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,17 +13,28 @@
#include <StencilStream/Index.hpp>
#include <StencilStream/Stencil.hpp>

enum class CellStatus
{
Normal,
Invalid,
Halo,
};

struct Cell
{
stencil::index_t c;
stencil::index_t r;
stencil::index_t i_generation;
CellStatus status;
};

template <stencil::uindex_t radius>
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<stencil::uindex_t, 4>;
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<Cell, radius> const &stencil) const
{
Expand All @@ -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;
Expand Down
11 changes: 5 additions & 6 deletions tests/src/synthesis/main.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,6 @@ const uindex_t grid_height = 2 * tile_height;
const uindex_t burst_size = 1024;

using TransFunc = FPGATransFunc<stencil_radius>;
using Cell = TransFunc::Cell;
using Executor = StencilExecutor<Cell, stencil_radius, TransFunc, pipeline_length, tile_width, tile_height, burst_size>;

void exception_handler(cl::sycl::exception_list exceptions)
Expand Down Expand Up @@ -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};
}
}
}
Expand All @@ -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);
}
}
}
Expand Down
20 changes: 10 additions & 10 deletions tests/src/units/ExecutionKernel.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,17 +20,17 @@ using namespace cl::sycl;
void test_kernel(uindex_t n_generations)
{
using TransFunc = FPGATransFunc<stencil_radius>;
using in_pipe = HostPipe<class ExecutionKernelInPipeID, TransFunc::Cell>;
using out_pipe = HostPipe<class ExecutionKernelOutPipeID, TransFunc::Cell>;
using TestExecutionKernel = ExecutionKernel<TransFunc, TransFunc::Cell, stencil_radius, pipeline_length, tile_width, tile_height, in_pipe, out_pipe>;
using in_pipe = HostPipe<class ExecutionKernelInPipeID, Cell>;
using out_pipe = HostPipe<class ExecutionKernelOutPipeID, Cell>;
using TestExecutionKernel = ExecutionKernel<TransFunc, Cell, stencil_radius, pipeline_length, tile_width, tile_height, in_pipe, out_pipe>;

for (index_t c = -halo_radius; c < index_t(halo_radius + tile_width); c++)
{
for (index_t r = -halo_radius; r < index_t(halo_radius + tile_height); r++)
{
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
{
Expand All @@ -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<TransFunc::Cell, 2> output_buffer(range<2>(tile_width, tile_height));
buffer<Cell, 2> output_buffer(range<2>(tile_width, tile_height));

{
auto output_buffer_ac = output_buffer.get_access<access::mode::discard_write>();
Expand All @@ -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);
}
}
}
Expand Down
21 changes: 10 additions & 11 deletions tests/src/units/StencilExecutor.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,6 @@ using namespace stencil;
using namespace cl::sycl;

using TransFunc = FPGATransFunc<stencil_radius>;
using Cell = typename TransFunc::Cell;

TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer<T, 2>)", "[StencilExecutor]")
{
Expand All @@ -29,7 +28,7 @@ TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer<T, 2>)", "[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};
}
}
}
Expand All @@ -46,10 +45,10 @@ TEST_CASE("StencilExecutor::copy_output(cl::sycl::buffer<T, 2>)", "[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);
}
}
}
Expand All @@ -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};
}
}
}
Expand All @@ -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);
}
}
}
Expand Down

0 comments on commit 0e2bfa5

Please sign in to comment.