stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius > Class Template Reference

A grid class for the tiling architecture. More...

#include <Grid.hpp>

Classes
class	GridAccessor
	An accessor for the monotile grid. More...

Public Member Functions
	Grid (uindex_t grid_width, uindex_t grid_height)
	Create a new, uninitialized grid with the given dimensions.
	Grid (sycl::range< 2 > range)
	Create a new, uninitialized grid with the given dimensions.
	Grid (sycl::buffer< Cell, 2 > input_buffer)
	Create a new grid with the same size and contents as the given SYCL buffer.
	Grid (Grid const &other_grid)
	Create a new reference to the given grid.
void	copy_from_buffer (sycl::buffer< Cell, 2 > input_buffer)
	Copy the contents of the SYCL buffer into the grid.
void	copy_to_buffer (sycl::buffer< Cell, 2 > output_buffer)
	Copy the contents of the grid into the SYCL buffer.
Grid	make_similar () const
	Create an new, uninitialized grid with the same size as the current one.
uindex_t	get_grid_width () const
	Return the width, or number of columns, of the grid.
uindex_t	get_grid_height () const
	Return the height, or number of rows, of the grid.
GenericID< uindex_t >	get_tile_range () const
	Return the range of (central) tiles of the grid.
template<typename in_pipe >
sycl::event	submit_read (sycl::queue &queue, uindex_t tile_c, uindex_t tile_r, Cell halo_value)
	Submit a kernel that sends a tile of the grid into a pipe.
template<typename out_pipe >
sycl::event	submit_write (sycl::queue queue, uindex_t tile_c, uindex_t tile_r)
	Submit a kernel that receives cells from the pipe and writes them to the grid.

Static Public Attributes
static constexpr uindex_t	dimensions = 2
	The number of dimensions of the grid.

Detailed Description

template<typename Cell, uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>
class stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >

A grid class for the tiling architecture.

This grid, which fullfils the Grid concept, contains a two-dimensional buffer of cells to be used together with the StencilUpdate class.

Since this class template requires multiple template arguments that must match the used StencilUpdate instance, it is advised to use the StencilUpdate::GridImpl shorthand instead of re-applying the same arguments to the grid class template.

The contents of the grid can be accessed by a host-side program using the GridAccessor class template. For example, one can write the contents of a grid object as follows:

Grid::GridAccessor<sycl::access::mode::read_write> accessor(grid);
for (uindex_t c = 0; c < grid.get_grid_width(); c++) {
    for (uindex_t r = 0; r < grid.get_grid_height(); r++) {
        accessor[c][r] = foo(c, r);
    }
}

Alternatively, one may write their data into a SYCL buffer and copy it into the grid using the method copy_from_buffer. The method copy_to_buffer does the reverse: It writes the contents of the grid into a SYCL buffer.

On the device side, the data can be read or written with the help of the method templates submit_read and submit_write. Those take a SYCL pipe as a template argument and enqueue kernels that read/write the contents of the grid to/from the pipes.

Template Parameters

Cell	The cell type to store.
tile_width	The width of a grid tile. This has to match the tile width of the used StencilUpdate class instance.
tile_height	The height of a grid tile. This has to match the tile height of the used StencilUpdate class instance.
halo_radius	The halo radius required for input tiles. This has to be the number of PEs in a StencilUpdate times the stencil radius of the implemented transition function.

Constructor & Destructor Documentation

Grid() [1/4]

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::Grid ( uindex_t  grid_width, uindex_t  grid_height  )

inline

Create a new, uninitialized grid with the given dimensions.

Parameters

grid_width	The width, or number of columns, of the new grid.
grid_height	The height, or number of rows, of the new grid.

Grid() [2/4]

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::Grid ( sycl::range< 2 >  range )

inline

Create a new, uninitialized grid with the given dimensions.

Parameters

range

The range of the new grid. The first index will be the width and the second index will be the height of the grid.

Grid() [3/4]

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::Grid ( sycl::buffer< Cell, 2 >  input_buffer )

inline

Create a new grid with the same size and contents as the given SYCL buffer.

The contents of the buffer will be copied to the grid by the host. The SYCL buffer can later be used elsewhere.

Parameters

input_buffer

The buffer with the contents of the new grid.

Grid() [4/4]

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::Grid ( Grid< Cell, tile_width, tile_height, halo_radius > const &  other_grid )

inline

Create a new reference to the given grid.

The newly created grid object will point to the same underlying data as the referenced grid. Changes made via the newly created grid object will also be visible to the old grid object, and vice-versa.

Parameters

other_grid

The other grid the new grid should reference.

Member Function Documentation

copy_from_buffer()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

void stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::copy_from_buffer ( sycl::buffer< Cell, 2 >  input_buffer )

inline

Copy the contents of the SYCL buffer into the grid.

The SYCL buffer will be accessed read-only one the host; It may be used elsewhere too. The buffer however has to have the same size as the grid, otherwise a std::range_error is thrown.

Parameters

input_buffer

The buffer to copy the data from.

Exceptions

std::range_error

The size of the buffer does not match the grid.

copy_to_buffer()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

void stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::copy_to_buffer ( sycl::buffer< Cell, 2 >  output_buffer )

inline

Copy the contents of the grid into the SYCL buffer.

The contents of the SYCL buffer will be overwritten on the host. The buffer also has to have the same size as the grid, otherwise a std::range_error is thrown.

Parameters

output_buffer

The buffer to copy the data to.

Exceptions

std::range_error

The size of the buffer does not match the grid.

get_grid_height()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

uindex_t stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::get_grid_height ( ) const

inline

Return the height, or number of rows, of the grid.

get_grid_width()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

uindex_t stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::get_grid_width ( ) const

inline

Return the width, or number of columns, of the grid.

get_tile_range()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

GenericID< uindex_t > stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::get_tile_range ( ) const

inline

Return the range of (central) tiles of the grid.

This is not the range of a single tile nor is the range of the grid. It is the range of valid arguments for submit_read and submit_write. For example, if the grid is 60 by 60 cells in size and a tile is 32 by 32 cells in size, the tile range would be 2 by 2 tiles.

Returns

The range of tiles of the grid.

make_similar()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

Grid stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::make_similar ( ) const

inline

Create an new, uninitialized grid with the same size as the current one.

submit_read()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

template<typename in_pipe >

sycl::event stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::submit_read ( sycl::queue &  queue, uindex_t  tile_c, uindex_t  tile_r, Cell  halo_value  )

inline

Submit a kernel that sends a tile of the grid into a pipe.

The submitted kernel will send the contents of a tile, along with the required halo, into a pipe in column-major order. This means that the last index (which denotes the row) will change the quickest. The method returns the event of the launched kernel immediately.

The tile column and row indices denote the indices of the tile, not the cells. The starting column index will be tile_c * tile_width - halo_radius and the end index will be (tile_c + 1) * tile_width + halo_radius. The start and end indices for the rows are analogous.

This method is explicitly part of the user-facing API: You are allowed and encouraged to use this method to feed custom kernels.

Template Parameters

in_pipe

The pipe the data is sent into.

Parameters

queue	The queue to submit the kernel to.
tile_c	The column index of the tile to read.
tile_r	The row index of the tile to read.
halo_value	The value to present for cells outside of the grid.

Exceptions

std::out_of_range

The grid does not contain the requested tile; Either the column or row index to high.

Returns

The event object of the submitted kernel.

submit_write()

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

template<typename out_pipe >

sycl::event stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::submit_write ( sycl::queue  queue, uindex_t  tile_c, uindex_t  tile_r  )

inline

Submit a kernel that receives cells from the pipe and writes them to the grid.

The kernel expects that one tile worth of cells can be read from the pipe. Also, it expects that the cells are sent in column-major order, meaning that the last index (which denotes the row) will change the quickest. The method returns the event of the launched kernel immediately.

The tile column and row indices denote the indices of the tile, not the cells. The starting column index will be tile_c * tile_width and the end index will be (tile_c + 1) * tile_width. The start and end indices for the rows are analogous.

This method is explicitly part of the user-facing API: You are allowed and encouraged to use this method to feed custom kernels.

Template Parameters

out_pipe

The pipe the data is received from.

Parameters

queue	The queue to submit the kernel to.
tile_c	The column index of the tile to read.
tile_r	The row index of the tile to read.

Exceptions

std::out_of_range

The grid does not contain the requested tile; Either the column or row index to high.

Returns

The event object of the submitted kernel.

Member Data Documentation

dimensions

template<typename Cell , uindex_t tile_width = 1024, uindex_t tile_height = 1024, uindex_t halo_radius = 1>

constexpr uindex_t stencil::tiling::Grid< Cell, tile_width, tile_height, halo_radius >::dimensions = 2

staticconstexpr

The number of dimensions of the grid.

May be changed in the future when other dimensions are supported.

---

The documentation for this class was generated from the following file:

• StencilStream/tiling/Grid.hpp