diff --git a/Doxyfile b/Doxyfile index 9709067..5452467 100644 --- a/Doxyfile +++ b/Doxyfile @@ -829,7 +829,7 @@ WARN_LOGFILE = # spaces. See also FILE_PATTERNS and EXTENSION_MAPPING # Note: If this tag is empty the current directory is searched. -INPUT = StencilStream README.md docs +INPUT = StencilStream StencilStream/monotile StencilStream/tiling README.md docs # This tag can be used to specify the character encoding of the source files # that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses diff --git a/StencilStream/AbstractExecutor.hpp b/StencilStream/AbstractExecutor.hpp index 1ff43cb..d1b5d87 100644 --- a/StencilStream/AbstractExecutor.hpp +++ b/StencilStream/AbstractExecutor.hpp @@ -18,6 +18,7 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once +#include "GenericID.hpp" #include "Index.hpp" #include diff --git a/StencilStream/MonotileExecutor.hpp b/StencilStream/MonotileExecutor.hpp index 3d73d0e..7b5cbbf 100644 --- a/StencilStream/MonotileExecutor.hpp +++ b/StencilStream/MonotileExecutor.hpp @@ -21,8 +21,8 @@ * SOFTWARE. */ #pragma once -#include "MonotileExecutionKernel.hpp" #include "SingleQueueExecutor.hpp" +#include "monotile/ExecutionKernel.hpp" namespace stencil { template ; using out_pipe = cl::sycl::pipe; using ExecutionKernelImpl = - MonotileExecutionKernel; + monotile::ExecutionKernel; cl::sycl::queue &queue = this->get_queue(); diff --git a/StencilStream/SingleQueueExecutor.hpp b/StencilStream/SingleQueueExecutor.hpp index 6109cf8..1322140 100644 --- a/StencilStream/SingleQueueExecutor.hpp +++ b/StencilStream/SingleQueueExecutor.hpp @@ -27,11 +27,12 @@ namespace stencil { /** * \brief An abstract executor with common code for executors that work with a single SYCL queue. * - * This class contains common code for executors that submit an execution kernel to a single queue - * that computes up to a fixed number of generations. This includes queue selection and management - * as well as event-based runtime analysis and an \ref AbstractExecutor.run implementation that - * incorporates the former. Child classes need to implement \ref SingleQueueExecutor.run_pass, - * submitting their execution kernel to the queue. + * This class contains common code for executors that work with execution kernels running on a + * single queue. This includes queue selection and management as well as event-based runtime + * analysis. + * + * User code that may work with any kind of executor should use pointers to the general \ref + * AbstractExecutor. * * \tparam T The cell type. * \tparam stencil_radius The radius of the stencil buffer supplied to the transition function. @@ -40,10 +41,22 @@ namespace stencil { template class SingleQueueExecutor : public AbstractExecutor { public: + /** + * \brief Create a new executor. + * \param halo_value The value of cells that are outside the grid. + * \param trans_func The instance of the transition function that should be used to calculate + * new generations. + */ SingleQueueExecutor(T halo_value, TransFunc trans_func) : AbstractExecutor(halo_value, trans_func), queue(std::nullopt), runtime_sample() {} + /** + * \brief Return the configured queue. + * + * If no queue has been configured yet, this method will configure and return a queue targeting + * the FPGA emulator, without runtime analysis. + */ cl::sycl::queue &get_queue() { if (!this->queue.has_value()) { select_emulator(false); @@ -59,12 +72,17 @@ class SingleQueueExecutor : public AbstractExecutor()) { throw std::runtime_error( @@ -74,7 +92,21 @@ class SingleQueueExecutor : public AbstractExecutorqueue = queue; } + + /** + * \brief Set up a SYCL queue with the FPGA emulator device, without runtime analysis. * * Note that as of OneAPI Version 2021.1.1, device code is usually built either for CPU/GPU, for * the FPGA emulator or for a specific FPGA. Using the wrong queue with the wrong device will @@ -83,7 +115,7 @@ class SingleQueueExecutor : public AbstractExecutorhas_property(); @@ -128,10 +165,9 @@ class SingleQueueExecutor : public AbstractExecutor /** * \brief Copy the current state of the grid to the buffer. * - * The \ref output_buffer has to have the exact range as returned by \ref + * The `output_buffer` has to have the exact range as returned by \ref * StencilExecutor.get_grid_range. * * \param output_buffer Copy the state of the grid to this buffer. @@ -116,8 +116,8 @@ class StencilExecutor : public SingleQueueExecutor using in_pipe = cl::sycl::pipe; using out_pipe = cl::sycl::pipe; using ExecutionKernelImpl = - TilingExecutionKernel; + tiling::ExecutionKernel; cl::sycl::queue &queue = this->get_queue(); @@ -126,7 +126,7 @@ class StencilExecutor : public SingleQueueExecutor uindex_t grid_height = input_grid.get_grid_range().r; while (this->get_i_generation() < target_i_generation) { - Grid output_grid = input_grid.make_output_grid(); + GridImpl output_grid = input_grid.make_output_grid(); std::vector events; events.reserve(input_grid.get_tile_range().c * input_grid.get_tile_range().r); @@ -166,7 +166,7 @@ class StencilExecutor : public SingleQueueExecutor } private: - using GridImpl = Grid; + using GridImpl = tiling::Grid; GridImpl input_grid; }; } // namespace stencil \ No newline at end of file diff --git a/StencilStream/MonotileExecutionKernel.hpp b/StencilStream/monotile/ExecutionKernel.hpp similarity index 96% rename from StencilStream/MonotileExecutionKernel.hpp rename to StencilStream/monotile/ExecutionKernel.hpp index 061324b..9fc95c1 100644 --- a/StencilStream/MonotileExecutionKernel.hpp +++ b/StencilStream/monotile/ExecutionKernel.hpp @@ -18,13 +18,14 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once -#include "GenericID.hpp" -#include "Helpers.hpp" -#include "Index.hpp" -#include "Stencil.hpp" +#include "../GenericID.hpp" +#include "../Helpers.hpp" +#include "../Index.hpp" +#include "../Stencil.hpp" #include namespace stencil { +namespace monotile { /** * \brief A kernel that executes a stencil transition function using the monotile approach. @@ -46,7 +47,7 @@ namespace stencil { */ template -class MonotileExecutionKernel { +class ExecutionKernel { public: static_assert( std::is_invocable_r const &>::value); @@ -90,8 +91,8 @@ class MonotileExecutionKernel { * \param grid_height The number of cell rows in the grid. * \param halo_value The value of cells outside the grid. */ - MonotileExecutionKernel(TransFunc trans_func, uindex_t i_generation, uindex_t n_generations, - uindex_t grid_width, uindex_t grid_height, T halo_value) + ExecutionKernel(TransFunc trans_func, uindex_t i_generation, uindex_t n_generations, + uindex_t grid_width, uindex_t grid_height, T halo_value) : trans_func(trans_func), i_generation(i_generation), n_generations(n_generations), grid_width(grid_width), grid_height(grid_height), halo_value(halo_value) {} @@ -221,4 +222,5 @@ class MonotileExecutionKernel { T halo_value; }; +} // namespace monotile } // namespace stencil \ No newline at end of file diff --git a/StencilStream/TilingExecutionKernel.hpp b/StencilStream/tiling/ExecutionKernel.hpp similarity index 95% rename from StencilStream/TilingExecutionKernel.hpp rename to StencilStream/tiling/ExecutionKernel.hpp index bc92460..e05fb01 100644 --- a/StencilStream/TilingExecutionKernel.hpp +++ b/StencilStream/tiling/ExecutionKernel.hpp @@ -18,12 +18,14 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once -#include "GenericID.hpp" -#include "Index.hpp" -#include "Stencil.hpp" +#include "../GenericID.hpp" +#include "../Helpers.hpp" +#include "../Index.hpp" +#include "../Stencil.hpp" #include namespace stencil { +namespace tiling { /** * \brief A kernel that executes a stencil transition function on a tile. @@ -42,7 +44,7 @@ namespace stencil { template -class TilingExecutionKernel { +class ExecutionKernel { public: static_assert( std::is_invocable_r const &>::value); @@ -86,9 +88,9 @@ class TilingExecutionKernel { * relative to the grid's origin. See `grid_c_offset` for details. \param grid_width The number * of cell columns in the grid. \param grid_height The number of cell rows in the grid. */ - TilingExecutionKernel(TransFunc trans_func, uindex_t i_generation, uindex_t target_i_generation, - uindex_t grid_c_offset, uindex_t grid_r_offset, uindex_t grid_width, - uindex_t grid_height, T halo_value) + ExecutionKernel(TransFunc trans_func, uindex_t i_generation, uindex_t target_i_generation, + uindex_t grid_c_offset, uindex_t grid_r_offset, uindex_t grid_width, + uindex_t grid_height, T halo_value) : trans_func(trans_func), i_generation(i_generation), target_i_generation(target_i_generation), grid_c_offset(grid_c_offset), grid_r_offset(grid_r_offset), grid_width(grid_width), grid_height(grid_height), @@ -200,4 +202,5 @@ class TilingExecutionKernel { T halo_value; }; +} // namespace tiling } // namespace stencil \ No newline at end of file diff --git a/StencilStream/Grid.hpp b/StencilStream/tiling/Grid.hpp similarity index 98% rename from StencilStream/Grid.hpp rename to StencilStream/tiling/Grid.hpp index 71d350f..9ff9724 100644 --- a/StencilStream/Grid.hpp +++ b/StencilStream/tiling/Grid.hpp @@ -18,7 +18,7 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once -#include "GenericID.hpp" +#include "../GenericID.hpp" #include "IOKernel.hpp" #include "Tile.hpp" #include @@ -28,6 +28,7 @@ #include namespace stencil { +namespace tiling { /** * \brief A rectangular container of cells with a dynamic, arbitrary size, used by the \ref @@ -35,7 +36,7 @@ namespace stencil { * * This class is part of the \ref tiling architecture. It logically contains the grid the transition * function is applied to and it partitions the grid into tiles of static size. These are the units - * the \ref TilingExecutionKernel works on. + * the \ref ExecutionKernel works on. * * Apart from providing copy operations to and from monolithic grid buffers, it also handles the * input and output kernel submission for a given tile. @@ -56,7 +57,7 @@ class Grid { /** * \brief Create a grid with undefined contents. * - * This constructor is used to create the output grid of a \ref TilingExecutionKernel + * This constructor is used to create the output grid of a \ref ExecutionKernel * invocation. It's contents do not need to be initialized or copied from another buffer since * it will override cell values from the execution kernel anyway. * @@ -153,7 +154,7 @@ class Grid { } /** - * \brief Submit the input kernels required for one execution of the \ref TilingExecutionKernel. + * \brief Submit the input kernels required for one execution of the \ref ExecutionKernel. * * This will submit five \ref IOKernel invocations in total, which are executed in order. Those * kernels write the contents of a tile and it's halo to the `in_pipe`. @@ -227,7 +228,7 @@ class Grid { /** * \brief Submit the output kernels required for one execution of the \ref - * TilingExecutionKernel. + * ExecutionKernel. * * This will submit three \ref IOKernel invocations in total, which are executed in order. Those * kernels will write cells from the `out_pipe` to one of the tiles. @@ -356,4 +357,5 @@ class Grid { cl::sycl::range<2> grid_range; }; +} // namespace tiling } // namespace stencil \ No newline at end of file diff --git a/StencilStream/IOKernel.hpp b/StencilStream/tiling/IOKernel.hpp similarity index 96% rename from StencilStream/IOKernel.hpp rename to StencilStream/tiling/IOKernel.hpp index 2e0f534..03c5e1e 100644 --- a/StencilStream/IOKernel.hpp +++ b/StencilStream/tiling/IOKernel.hpp @@ -18,21 +18,22 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once -#include "CounterID.hpp" -#include "Index.hpp" +#include "../CounterID.hpp" +#include "../Index.hpp" #include #include #include namespace stencil { +namespace tiling { /** - * \brief Generic Input/Output kernel for use with the \ref TilingExecutionKernel and \ref Grid. + * \brief Generic Input/Output kernel for use with the \ref ExecutionKernel and \ref Grid. * * This kernel provides IO services to the execution kernel by writing the contents of a input tile * with halo to the input pipe and writing the output of the execution kernel to an output tile. The * input and output code only differs by one line. Therefore, both services are provided by the same - * class. Unlike \ref TilingExecutionKernel, this kernel is supposed to be constructed by a lambda + * class. Unlike \ref ExecutionKernel, this kernel is supposed to be constructed by a lambda * expression that then either executes \ref IOKernel.read or \ref IOKernel.write. * * Logically, an instance of the IO kernel receives access to a horizontal slice of the input or @@ -156,4 +157,5 @@ class IOKernel { uindex_t n_columns; }; +} // namespace tiling } // namespace stencil \ No newline at end of file diff --git a/StencilStream/Tile.hpp b/StencilStream/tiling/Tile.hpp similarity index 98% rename from StencilStream/Tile.hpp rename to StencilStream/tiling/Tile.hpp index e3eca1a..041db2b 100644 --- a/StencilStream/Tile.hpp +++ b/StencilStream/tiling/Tile.hpp @@ -18,18 +18,19 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #pragma once -#include "Helpers.hpp" -#include "Index.hpp" +#include "../Helpers.hpp" +#include "../Index.hpp" #include #include #include namespace stencil { +namespace tiling { /** * \brief A rectangular container of cells with a static size * - * StencilStream tiles the grid and the \ref TilingExecutionKernel works with those tiles: it + * StencilStream tiles the grid and the \ref ExecutionKernel works with those tiles: it * receives the content of a tile (together with it's halo) and emits the contents of a tile. A tile * is partitioned in four corner buffers, four border buffers and one core buffer. This is done to * provide easy access to the halo of a tile. Have a look at the \ref Architecture for more details @@ -298,4 +299,6 @@ class Tile { std::optional> part[3][3]; }; + +} // namespace tiling } // namespace stencil \ No newline at end of file diff --git a/examples/fdtd/src/fdtd.cpp b/examples/fdtd/src/fdtd.cpp index a457c54..cd11a02 100644 --- a/examples/fdtd/src/fdtd.cpp +++ b/examples/fdtd/src/fdtd.cpp @@ -147,7 +147,7 @@ int main(int argc, char **argv) { Executor executor(FDTDKernel::halo(), FDTDKernel(parameters)); executor.set_input(grid_buffer); - executor.set_queue(fpga_queue, true); + executor.set_queue(fpga_queue); uindex_t n_timesteps = parameters.n_timesteps(); @@ -167,7 +167,7 @@ int main(int argc, char **argv) { save_frame(grid_buffer, i_generation, CellField::HZ_SUM, parameters); save_frame(grid_buffer, i_generation, CellField::DISTANCE, parameters); - runtime += executor.get_runtime_sample()->get_total_runtime(); + runtime += executor.get_runtime_sample().get_total_runtime(); double progress = 100.0 * (double(i_generation) / double(n_timesteps)); double speed = double(i_generation) / runtime; double time_remaining = double(n_timesteps - i_generation) / speed; diff --git a/examples/hotspot/hotspot.cpp b/examples/hotspot/hotspot.cpp index 4239871..ac90bba 100644 --- a/examples/hotspot/hotspot.cpp +++ b/examples/hotspot/hotspot.cpp @@ -141,8 +141,7 @@ auto exception_handler = [](cl::sycl::exception_list exceptions) { } }; -double run_simulation(cl::sycl::queue working_queue, buffer temp, - uindex_t sim_time) { +double run_simulation(cl::sycl::queue working_queue, buffer temp, uindex_t sim_time) { uindex_t n_columns = temp.get_range()[0]; uindex_t n_rows = temp.get_range()[1]; @@ -197,11 +196,11 @@ double run_simulation(cl::sycl::queue working_queue, buffer temp, }; #ifdef MONOTILE - using Executor = MonotileExecutor; + using Executor = MonotileExecutor; #else - using Executor = StencilExecutor; + using Executor = StencilExecutor; #endif Executor executor(Cell(0.0, 0.0), kernel); @@ -217,7 +216,7 @@ double run_simulation(cl::sycl::queue working_queue, buffer temp, executor.copy_output(temp); - return executor.get_runtime_sample().value().get_total_runtime(); + return executor.get_runtime_sample().get_total_runtime(); } int main(int argc, char **argv) { @@ -251,8 +250,7 @@ int main(int argc, char **argv) { tfile = argv[4]; pfile = argv[5]; ofile = argv[6]; - buffer temp = - read_input(string(tfile), string(pfile), range<2>(n_columns, n_rows)); + buffer temp = read_input(string(tfile), string(pfile), range<2>(n_columns, n_rows)); printf("Start computing the transient temperature\n"); diff --git a/tests/Makefile b/tests/Makefile index 785e062..6bf3e61 100644 --- a/tests/Makefile +++ b/tests/Makefile @@ -8,7 +8,7 @@ CC = dpcpp STENCIL_PATH = ../ -ARGS = -std=c++17 -g -I$(STENCIL_PATH) -fintelfpga -DSTENCIL_INDEX_WIDTH=32 +ARGS = -std=c++17 -g -I$(STENCIL_PATH) -Isrc/ -fintelfpga -DSTENCIL_INDEX_WIDTH=32 ifdef OPTIMIZE ARGS += -O3 @@ -19,7 +19,8 @@ ifdef EBROOTGCC endif UNIT_ARGS = $(ARGS) -UNIT_OBJECTS = $(patsubst %.cpp,%.o,$(wildcard src/units/*.cpp)) +UNIT_SOURCES = $(wildcard src/units/*.cpp src/units/*/*.cpp) +UNIT_OBJECTS = $(patsubst %.cpp,%.o,$(UNIT_SOURCES)) RESOURCES = src/res/*.hpp ../StencilStream/* Makefile SYNTH_ARGS = $(ARGS) -Xsv diff --git a/tests/src/synthesis/main.cpp b/tests/src/synthesis/main.cpp index 52a042d..9b9d3a6 100644 --- a/tests/src/synthesis/main.cpp +++ b/tests/src/synthesis/main.cpp @@ -72,7 +72,7 @@ int main() { Executor executor(Cell::halo(), TransFunc()); executor.set_input(in_buffer); - executor.set_queue(working_queue, false); + executor.set_queue(working_queue); executor.run(2 * pipeline_length); diff --git a/tests/src/units/CounterID.cpp b/tests/src/units/CounterID.cpp index edf35af..9994189 100644 --- a/tests/src/units/CounterID.cpp +++ b/tests/src/units/CounterID.cpp @@ -17,9 +17,9 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/catch.hpp" -#include "../res/constants.hpp" #include +#include +#include using namespace std; using namespace stencil; diff --git a/tests/src/units/HostPipe.cpp b/tests/src/units/HostPipe.cpp index 359f58c..4127dee 100644 --- a/tests/src/units/HostPipe.cpp +++ b/tests/src/units/HostPipe.cpp @@ -17,8 +17,8 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/HostPipe.hpp" -#include "../res/catch.hpp" +#include +#include using namespace stencil; diff --git a/tests/src/units/MonotileExecutionKernel.cpp b/tests/src/units/MonotileExecutionKernel.cpp deleted file mode 100644 index e1081f4..0000000 --- a/tests/src/units/MonotileExecutionKernel.cpp +++ /dev/null @@ -1,112 +0,0 @@ -/* - * Copyright © 2020-2021 Jan-Oliver Opdenhövel, Paderborn Center for Parallel Computing, Paderborn University - * - * Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the “Software”), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED “AS IS”, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - */ -#include "../res/TransFuncs.hpp" -#include "../res/HostPipe.hpp" -#include "../res/catch.hpp" -#include "../res/constants.hpp" -#include - -using namespace stencil; -using namespace std; -using namespace cl::sycl; - -void test_monotile_kernel(uindex_t n_generations) -{ - using TransFunc = HostTransFunc; - using in_pipe = HostPipe; - using out_pipe = HostPipe; - using TestExecutionKernel = MonotileExecutionKernel; - - for (uindex_t c = 0; c < tile_width; c++) - { - for (uindex_t r = 0; r < tile_height; r++) - { - in_pipe::write(Cell{index_t(c), index_t(r), 0, CellStatus::Normal}); - } - } - - TestExecutionKernel(TransFunc(), 0, n_generations, tile_width, tile_height, Cell::halo())(); - - buffer output_buffer(range<2>(tile_width, tile_height)); - - { - auto output_buffer_ac = output_buffer.get_access(); - for (uindex_t c = 0; c < tile_width; c++) - { - for (uindex_t r = 0; r < tile_height; r++) - { - output_buffer_ac[c][r] = out_pipe::read(); - } - } - } - - REQUIRE(in_pipe::empty()); - REQUIRE(out_pipe::empty()); - - auto output_buffer_ac = output_buffer.get_access(); - for (uindex_t c = 1; c < tile_width; c++) - { - for (uindex_t r = 1; r < tile_height; r++) - { - 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); - } - } -} - -TEST_CASE("MonotileExecutionKernel", "[MonotileExecutionKernel]") -{ - test_monotile_kernel(pipeline_length); -} - -TEST_CASE("MonotileExecutionKernel (partial pipeline)", "[MonotileExecutionKernel]") -{ - static_assert(pipeline_length != 1); - test_monotile_kernel(pipeline_length - 1); -} - -TEST_CASE("MonotileExecutionKernel (noop)", "[MonotileExecutionKernel]") -{ - test_monotile_kernel(0); -} - -TEST_CASE("MonotileExecutionKernel: Incomplete Pipeline with i_generation != 0", "[MonotileExecutionKernel]") -{ - using Cell = uint8_t; - auto trans_func = [](Stencil const &stencil) { - return stencil[ID(0,0)] + 1; - }; - - using in_pipe = HostPipe; - using out_pipe = HostPipe; - using TestExecutionKernel = MonotileExecutionKernel; - - for (int c = 0; c < 64; c++) { - for (int r = 0; r < 64; r++) { - in_pipe::write(0); - } - } - - TestExecutionKernel kernel(trans_func, 16, 20, 64, 64, 0); - kernel.operator()(); - - REQUIRE(in_pipe::empty()); - - for (int c = 0; c < 64; c++) { - for (int r = 0; r < 64; r++) { - REQUIRE(out_pipe::read() == 4); - } - } - - REQUIRE(out_pipe::empty()); -} \ No newline at end of file diff --git a/tests/src/units/Stencil.cpp b/tests/src/units/Stencil.cpp index 2608373..bed26ad 100644 --- a/tests/src/units/Stencil.cpp +++ b/tests/src/units/Stencil.cpp @@ -17,9 +17,9 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/catch.hpp" -#include "../res/constants.hpp" #include +#include +#include using namespace stencil; diff --git a/tests/src/units/StencilExecutor.cpp b/tests/src/units/StencilExecutor.cpp index 93dfd24..dc3ddf2 100644 --- a/tests/src/units/StencilExecutor.cpp +++ b/tests/src/units/StencilExecutor.cpp @@ -18,11 +18,11 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/TransFuncs.hpp" -#include "../res/catch.hpp" -#include "../res/constants.hpp" #include #include +#include +#include +#include using namespace std; using namespace stencil; diff --git a/tests/src/units/main.cpp b/tests/src/units/main.cpp index 7e9b806..cf04247 100644 --- a/tests/src/units/main.cpp +++ b/tests/src/units/main.cpp @@ -18,4 +18,4 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #define CATCH_CONFIG_MAIN -#include "../res/catch.hpp" \ No newline at end of file +#include \ No newline at end of file diff --git a/tests/src/units/monotile/ExecutionKernel.cpp b/tests/src/units/monotile/ExecutionKernel.cpp new file mode 100644 index 0000000..b6db0c0 --- /dev/null +++ b/tests/src/units/monotile/ExecutionKernel.cpp @@ -0,0 +1,113 @@ +/* + * Copyright © 2020-2021 Jan-Oliver Opdenhövel, Paderborn Center for Parallel Computing, Paderborn + * University + * + * Permission is hereby granted, free of charge, to any person obtaining a copy of this software and + * associated documentation files (the “Software”), to deal in the Software without restriction, + * including without limitation the rights to use, copy, modify, merge, publish, distribute, + * sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all copies or + * substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED “AS IS”, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT + * NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND + * NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, + * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ +#include +#include +#include +#include +#include + +using namespace stencil; +using namespace std; +using namespace cl::sycl; + +void test_monotile_kernel(uindex_t n_generations) { + using TransFunc = HostTransFunc; + using in_pipe = HostPipe; + using out_pipe = HostPipe; + using TestExecutionKernel = + monotile::ExecutionKernel; + + for (uindex_t c = 0; c < tile_width; c++) { + for (uindex_t r = 0; r < tile_height; r++) { + in_pipe::write(Cell{index_t(c), index_t(r), 0, CellStatus::Normal}); + } + } + + TestExecutionKernel(TransFunc(), 0, n_generations, tile_width, tile_height, Cell::halo())(); + + buffer output_buffer(range<2>(tile_width, tile_height)); + + { + auto output_buffer_ac = output_buffer.get_access(); + for (uindex_t c = 0; c < tile_width; c++) { + for (uindex_t r = 0; r < tile_height; r++) { + output_buffer_ac[c][r] = out_pipe::read(); + } + } + } + + REQUIRE(in_pipe::empty()); + REQUIRE(out_pipe::empty()); + + auto output_buffer_ac = output_buffer.get_access(); + for (uindex_t c = 1; c < tile_width; c++) { + for (uindex_t r = 1; r < tile_height; r++) { + 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); + } + } +} + +TEST_CASE("monotile::ExecutionKernel", "[monotile::ExecutionKernel]") { + test_monotile_kernel(pipeline_length); +} + +TEST_CASE("monotile::ExecutionKernel (partial pipeline)", "[monotile::ExecutionKernel]") { + static_assert(pipeline_length != 1); + test_monotile_kernel(pipeline_length - 1); +} + +TEST_CASE("monotile::ExecutionKernel (noop)", "[monotile::ExecutionKernel]") { + test_monotile_kernel(0); +} + +TEST_CASE("monotile::ExecutionKernel: Incomplete Pipeline with i_generation != 0", + "[monotile::ExecutionKernel]") { + using Cell = uint8_t; + auto trans_func = [](Stencil const &stencil) { return stencil[ID(0, 0)] + 1; }; + + using in_pipe = HostPipe; + using out_pipe = HostPipe; + using TestExecutionKernel = + monotile::ExecutionKernel; + + for (int c = 0; c < 64; c++) { + for (int r = 0; r < 64; r++) { + in_pipe::write(0); + } + } + + TestExecutionKernel kernel(trans_func, 16, 20, 64, 64, 0); + kernel.operator()(); + + REQUIRE(in_pipe::empty()); + + for (int c = 0; c < 64; c++) { + for (int r = 0; r < 64; r++) { + REQUIRE(out_pipe::read() == 4); + } + } + + REQUIRE(out_pipe::empty()); +} \ No newline at end of file diff --git a/tests/src/units/TilingExecutionKernel.cpp b/tests/src/units/tiling/ExecutionKernel.cpp similarity index 83% rename from tests/src/units/TilingExecutionKernel.cpp rename to tests/src/units/tiling/ExecutionKernel.cpp index 295fc78..3dbd960 100644 --- a/tests/src/units/TilingExecutionKernel.cpp +++ b/tests/src/units/tiling/ExecutionKernel.cpp @@ -17,23 +17,23 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/HostPipe.hpp" -#include "../res/TransFuncs.hpp" -#include "../res/catch.hpp" -#include "../res/constants.hpp" -#include +#include +#include +#include +#include +#include using namespace stencil; using namespace std; +using namespace stencil::tiling; using namespace cl::sycl; void test_tiling_kernel(uindex_t n_generations) { using TransFunc = FPGATransFunc; using in_pipe = HostPipe; using out_pipe = HostPipe; - using TestExecutionKernel = - TilingExecutionKernel; + using TestExecutionKernel = ExecutionKernel; 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++) { @@ -75,18 +75,18 @@ void test_tiling_kernel(uindex_t n_generations) { } } -TEST_CASE("TilingExecutionKernel", "[TilingExecutionKernel]") { +TEST_CASE("tiling::ExecutionKernel", "[tiling::ExecutionKernel]") { test_tiling_kernel(pipeline_length); } -TEST_CASE("TilingExecutionKernel (partial pipeline)", "[TilingExecutionKernel]") { +TEST_CASE("tiling::ExecutionKernel (partial pipeline)", "[tiling::ExecutionKernel]") { static_assert(pipeline_length != 1); test_tiling_kernel(pipeline_length - 1); } -TEST_CASE("TilingExecutionKernel (noop)", "[TilingExecutionKernel]") { test_tiling_kernel(0); } +TEST_CASE("tiling::ExecutionKernel (noop)", "[tiling::ExecutionKernel]") { test_tiling_kernel(0); } -TEST_CASE("Halo values inside the pipeline are handled correctly", "[TilingExecutionKernel]") { +TEST_CASE("Halo values inside the pipeline are handled correctly", "[tiling::ExecutionKernel]") { auto my_kernel = [=](Stencil const &stencil) { ID idx = stencil.id; bool is_valid = true; @@ -108,8 +108,8 @@ TEST_CASE("Halo values inside the pipeline are handled correctly", "[TilingExecu using in_pipe = HostPipe; using out_pipe = HostPipe; using TestExecutionKernel = - TilingExecutionKernel; + ExecutionKernel; 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++) { @@ -129,14 +129,14 @@ TEST_CASE("Halo values inside the pipeline are handled correctly", "[TilingExecu REQUIRE(out_pipe::empty()); } -TEST_CASE("Incomplete Pipeline with i_generation != 0", "[TilingExecutionKernel]") { +TEST_CASE("Incomplete Pipeline with i_generation != 0", "[tiling::ExecutionKernel]") { using Cell = uint8_t; auto trans_func = [](Stencil const &stencil) { return stencil[ID(0, 0)] + 1; }; using in_pipe = HostPipe; using out_pipe = HostPipe; using TestExecutionKernel = - TilingExecutionKernel; + ExecutionKernel; for (int c = -16; c < 16 + 64; c++) { for (int r = -16; r < 16 + 64; r++) { diff --git a/tests/src/units/Grid.cpp b/tests/src/units/tiling/Grid.cpp similarity index 97% rename from tests/src/units/Grid.cpp rename to tests/src/units/tiling/Grid.cpp index 52941c4..6517522 100644 --- a/tests/src/units/Grid.cpp +++ b/tests/src/units/tiling/Grid.cpp @@ -17,14 +17,15 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/catch.hpp" -#include "../res/constants.hpp" #include #include -#include +#include +#include +#include #include using namespace stencil; +using namespace stencil::tiling; using namespace cl::sycl; using namespace std; diff --git a/tests/src/units/IOKernel.cpp b/tests/src/units/tiling/IOKernel.cpp similarity index 98% rename from tests/src/units/IOKernel.cpp rename to tests/src/units/tiling/IOKernel.cpp index 58e7cd4..b2a8971 100644 --- a/tests/src/units/IOKernel.cpp +++ b/tests/src/units/tiling/IOKernel.cpp @@ -17,15 +17,16 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/HostPipe.hpp" -#include "../res/catch.hpp" -#include "../res/constants.hpp" +#include +#include +#include #include #include #include -#include +#include using namespace cl::sycl; +using namespace stencil::tiling; using namespace stencil; using namespace std; diff --git a/tests/src/units/Tile.cpp b/tests/src/units/tiling/Tile.cpp similarity index 97% rename from tests/src/units/Tile.cpp rename to tests/src/units/tiling/Tile.cpp index a04986b..cb693eb 100644 --- a/tests/src/units/Tile.cpp +++ b/tests/src/units/tiling/Tile.cpp @@ -17,14 +17,15 @@ * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "../res/catch.hpp" -#include "../res/constants.hpp" #include #include -#include +#include +#include +#include using namespace stencil; using namespace cl::sycl; +using namespace stencil::tiling; using namespace std; using TileImpl = Tile;