Grid.hpp

Go to the documentation of this file.

/*
 * Copyright © 2020-2024 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.
 */
#pragma once
#include "../AccessorSubscript.hpp"
#include "../Concepts.hpp"
#include "../GenericID.hpp"
#include <memory>
#include <numeric>
#include <vector>
 
namespace stencil {
namespace tiling {
 
template <typename Cell, uindex_t tile_width = 1024, uindex_t tile_height = 1024,
          uindex_t halo_radius = 1>
class Grid {
    static_assert(2 * halo_radius < tile_height && 2 * halo_radius < tile_width);
 
  public:
    static constexpr uindex_t dimensions = 2;
 
    Grid(uindex_t grid_width, uindex_t grid_height)
        : grid_buffer(sycl::range<2>(grid_width, grid_height)) {}
 
    Grid(sycl::range<2> range) : grid_buffer(range) {}
 
    Grid(sycl::buffer<Cell, 2> input_buffer) : grid_buffer(input_buffer.get_range()) {
        copy_from_buffer(input_buffer);
    }
 
    Grid(Grid const &other_grid) : grid_buffer(other_grid.grid_buffer) {}
 
    template <sycl::access::mode access_mode> class GridAccessor {
      public:
        static constexpr uindex_t dimensions = Grid::dimensions;
 
        GridAccessor(Grid &grid) : accessor(grid.grid_buffer) {}
 
        using BaseSubscript = AccessorSubscript<Cell, GridAccessor, access_mode>;
 
        BaseSubscript operator[](uindex_t i) { return BaseSubscript(*this, i); }
 
        Cell const &operator[](sycl::id<2> id)
            requires(access_mode == sycl::access::mode::read)
        {
            return accessor[id];
        }
 
        Cell &operator[](sycl::id<2> id)
            requires(access_mode != sycl::access::mode::read)
        {
            return accessor[id];
        }
 
      private:
        sycl::host_accessor<Cell, 2, access_mode> accessor;
    };
 
    void copy_from_buffer(sycl::buffer<Cell, 2> input_buffer) {
        if (input_buffer.get_range() != grid_buffer.get_range()) {
            throw std::out_of_range("The target buffer has not the same size as the grid");
        }
 
        sycl::host_accessor grid_ac{grid_buffer, sycl::write_only};
        sycl::host_accessor input_ac{input_buffer, sycl::read_only};
        for (uindex_t c = 0; c < get_grid_width(); c++) {
            for (uindex_t r = 0; r < get_grid_height(); r++) {
                grid_ac[c][r] = input_ac[c][r];
            }
        }
    }
 
    void copy_to_buffer(sycl::buffer<Cell, 2> output_buffer) {
        if (output_buffer.get_range() != grid_buffer.get_range()) {
            throw std::out_of_range("The target buffer has not the same size as the grid");
        }
 
        sycl::host_accessor grid_ac{grid_buffer, sycl::read_only};
        sycl::host_accessor output_ac{output_buffer, sycl::write_only};
        for (uindex_t c = 0; c < get_grid_width(); c++) {
            for (uindex_t r = 0; r < get_grid_height(); r++) {
                output_ac[c][r] = grid_ac[c][r];
            }
        }
    }
 
    Grid make_similar() const { return Grid(grid_buffer.get_range()); }
 
    uindex_t get_grid_width() const { return grid_buffer.get_range()[0]; }
 
    uindex_t get_grid_height() const { return grid_buffer.get_range()[1]; }
 
    GenericID<uindex_t> get_tile_range() const {
        return GenericID<uindex_t>(std::ceil(float(get_grid_width()) / float(tile_width)),
                                   std::ceil(float(get_grid_height()) / float(tile_height)));
    }
 
    template <typename in_pipe>
    sycl::event submit_read(sycl::queue &queue, uindex_t tile_c, uindex_t tile_r, Cell halo_value) {
        if (tile_c >= get_tile_range().c || tile_r >= get_tile_range().r) {
            throw std::out_of_range("Tile index out of range!");
        }
 
        constexpr uindex_t column_bits = 1 + std::bit_width(tile_width + halo_radius);
        constexpr uindex_t row_bits = 1 + std::bit_width(tile_height + halo_radius);
        using index_c_t = ac_int<column_bits, true>;
        using index_r_t = ac_int<row_bits, true>;
 
        return queue.submit([&](sycl::handler &cgh) {
            sycl::accessor grid_ac{grid_buffer, cgh, sycl::read_only};
            index_t grid_width = this->get_grid_width();
            index_t grid_height = this->get_grid_height();
 
            cgh.single_task([=]() {
                index_t c_offset = tile_c * tile_width;
                index_c_t start_tile_c = -halo_radius;
                index_c_t end_tile_c =
                    index_c_t(std::min(grid_width - tile_c * tile_width, tile_width)) + halo_radius;
 
                index_t r_offset = tile_r * tile_height;
                index_r_t start_tile_r = -halo_radius;
                index_r_t end_tile_r =
                    index_r_t(std::min(grid_height - tile_r * tile_height, tile_height)) +
                    halo_radius;
 
                [[intel::loop_coalesce(2)]] for (index_c_t tile_c = start_tile_c;
                                                 tile_c < end_tile_c; tile_c++) {
                    for (index_r_t tile_r = start_tile_r; tile_r < end_tile_r; tile_r++) {
                        index_t c = c_offset + tile_c.to_long();
                        index_t r = r_offset + tile_r.to_long();
 
                        Cell value;
                        if (c >= 0 && r >= 0 && c < grid_width && r < grid_height) {
                            value = grid_ac[c][r];
                        } else {
                            value = halo_value;
                        }
                        in_pipe::write(value);
                    }
                }
            });
        });
    }
 
    template <typename out_pipe>
    sycl::event submit_write(sycl::queue queue, uindex_t tile_c, uindex_t tile_r) {
        if (tile_c >= get_tile_range().c || tile_r >= get_tile_range().r) {
            throw std::out_of_range("Tile index out of range!");
        }
 
        constexpr uindex_t column_bits = std::bit_width(tile_width);
        constexpr uindex_t row_bits = std::bit_width(tile_height);
        using uindex_c_t = ac_int<column_bits, false>;
        using uindex_r_t = ac_int<row_bits, false>;
 
        return queue.submit([&](sycl::handler &cgh) {
            sycl::accessor grid_ac{grid_buffer, cgh, sycl::read_write};
            uindex_t grid_width = this->get_grid_width();
            uindex_t grid_height = this->get_grid_height();
 
            cgh.single_task([=]() {
                uindex_t c_offset = tile_c * tile_width;
                uindex_c_t end_tile_c =
                    uindex_c_t(std::min(grid_width - tile_c * tile_width, tile_width));
 
                uindex_t r_offset = tile_r * tile_height;
                uindex_r_t end_tile_r =
                    uindex_r_t(std::min(grid_height - tile_r * tile_height, tile_height));
 
                [[intel::loop_coalesce(2)]] for (uindex_c_t tile_c = 0; tile_c < end_tile_c;
                                                 tile_c++) {
                    for (uindex_r_t tile_r = 0; tile_r < end_tile_r; tile_r++) {
                        grid_ac[c_offset + tile_c.to_long()][r_offset + tile_r.to_long()] =
                            out_pipe::read();
                    }
                }
            });
        });
    }
 
  private:
    sycl::buffer<Cell, 2> grid_buffer;
};
 
} // namespace tiling
} // namespace stencil