11631827009

Committed 01 Nov 2024 03:26PM UTC coverage: 95.265% (+0.01%) from 95.251%

Build # 11631827009

Build Type

Pull #292

github

Committed by

fknorr

Commit Message

Shortcut-evaluate region_intersection on box-shaped regions

Pull Request Pull Request #292: Generalize range mappers to regions, add `neighborhood(along_axes)` RM

Run Details

3016 of 3402 branches covered (88.65%)

Branch coverage included in aggregate %.

86 of 87 new or added lines in 8 files covered. (98.85%)

13 existing lines in 4 files now uncovered.

6682 of 6778 relevant lines covered (98.58%)

1304731.31 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

94.17

/include/range_mapper.h

#pragma once

#include <stdexcept>
#include <type_traits>

#include <fmt/format.h>
#include <sycl/sycl.hpp>

#include "grid.h"
#include "ranges.h"
#include "utils.h"


namespace celerity {
namespace detail {

        template <typename Functor, int BufferDims, int KernelDims>
        constexpr bool is_range_mapper_invocable_for_chunk_only = std::is_invocable_r_v<region<BufferDims>, const Functor&, const celerity::chunk<KernelDims>&>;

        template <typename Functor, int BufferDims, int KernelDims>
        constexpr bool is_range_mapper_invocable_for_chunk_and_global_size =
            std::is_invocable_r_v<region<BufferDims>, const Functor&, const celerity::chunk<KernelDims>&, const range<BufferDims>&>;

        template <typename Functor, int BufferDims, int KernelDims>
        constexpr bool is_range_mapper_invocable_for_kernel = is_range_mapper_invocable_for_chunk_only<Functor, BufferDims, KernelDims> //
                                                              || is_range_mapper_invocable_for_chunk_and_global_size<Functor, BufferDims, KernelDims>;

        template <typename Functor, int BufferDims>
        constexpr bool is_range_mapper_invocable = is_range_mapper_invocable_for_kernel<Functor, BufferDims, 0>    //
                                                   || is_range_mapper_invocable_for_kernel<Functor, BufferDims, 1> //
                                                   || is_range_mapper_invocable_for_kernel<Functor, BufferDims, 2> //
                                                   || is_range_mapper_invocable_for_kernel<Functor, BufferDims, 3>;

        [[noreturn]] inline void throw_invalid_range_mapper_args(int expect_kernel_dims, int expect_buffer_dims) {
                throw std::runtime_error(fmt::format("Invalid range mapper dimensionality: {0}-dimensional kernel submitted with a requirement whose range mapper "
                                                     "is neither invocable for chunk<{0}> nor (chunk<{0}>, range<{1}>) to produce subrange<{1}>",
                    expect_kernel_dims, expect_buffer_dims));
        }

        [[noreturn]] inline void throw_invalid_range_mapper_result(int expect_sr_dims, int actual_sr_dims, int kernel_dims) {
                throw std::runtime_error(fmt::format("Range mapper produces subrange of wrong dimensionality: Expecting subrange<{}>, got subrange<{}> for chunk<{}>",
                    expect_sr_dims, actual_sr_dims, kernel_dims));
        }

        template <typename Functor, int KernelDims, int BufferDims>
        region<BufferDims> invoke_range_mapper(Functor&& fn, const celerity::chunk<KernelDims>& chunk, const range<BufferDims>& buffer_size) {
                static_assert(KernelDims >= 0 && KernelDims <= 3 && BufferDims >= 0 && BufferDims <= 3);
                if constexpr(is_range_mapper_invocable_for_chunk_and_global_size<Functor, BufferDims, KernelDims>) {
                        return std::forward<Functor>(fn)(chunk, buffer_size);
                } else if constexpr(is_range_mapper_invocable_for_chunk_only<Functor, BufferDims, KernelDims>) {
                        return std::forward<Functor>(fn)(chunk);
                } else {
                        throw_invalid_range_mapper_args(KernelDims, BufferDims);
                }
        }

        template <int BufferDims>
        region<BufferDims> clamp_region_to_buffer_size(const region<BufferDims>& r, const range<BufferDims>& buffer_size) {
                return region_intersection(r, box<BufferDims>::full_range(buffer_size));
        }

        class range_mapper_base {
          public:
                explicit range_mapper_base(sycl::access::mode am) : m_access_mode(am) {}

                range_mapper_base(const range_mapper_base& other) = delete;
                range_mapper_base(range_mapper_base&& other) = delete;
                range_mapper_base& operator=(const range_mapper_base& other) = delete;
                range_mapper_base& operator=(range_mapper_base&& other) = delete;

                sycl::access::mode get_access_mode() const { return m_access_mode; }

                virtual int get_buffer_dimensions() const = 0;

                virtual region<1> map_1(const chunk<0>& chnk) const = 0;
                virtual region<1> map_1(const chunk<1>& chnk) const = 0;
                virtual region<1> map_1(const chunk<2>& chnk) const = 0;
                virtual region<1> map_1(const chunk<3>& chnk) const = 0;
                virtual region<2> map_2(const chunk<0>& chnk) const = 0;
                virtual region<2> map_2(const chunk<1>& chnk) const = 0;
                virtual region<2> map_2(const chunk<2>& chnk) const = 0;
                virtual region<2> map_2(const chunk<3>& chnk) const = 0;
                virtual region<3> map_3(const chunk<0>& chnk) const = 0;
                virtual region<3> map_3(const chunk<1>& chnk) const = 0;
                virtual region<3> map_3(const chunk<2>& chnk) const = 0;
                virtual region<3> map_3(const chunk<3>& chnk) const = 0;

                virtual ~range_mapper_base() = default;

          private:
                sycl::access::mode m_access_mode;
        };

        template <int BufferDims, typename Functor>
        class range_mapper final : public range_mapper_base {
          public:
                range_mapper(Functor rmfn, sycl::access::mode am, range<BufferDims> buffer_size) : range_mapper_base(am), m_rmfn(rmfn), m_buffer_size(buffer_size) {}

                int get_buffer_dimensions() const override { return BufferDims; }

                region<1> map_1(const chunk<0>& chnk) const override { return map<1>(chnk); }
                region<1> map_1(const chunk<1>& chnk) const override { return map<1>(chnk); }
                region<1> map_1(const chunk<2>& chnk) const override { return map<1>(chnk); }
                region<1> map_1(const chunk<3>& chnk) const override { return map<1>(chnk); }
                region<2> map_2(const chunk<0>& chnk) const override { return map<2>(chnk); }
                region<2> map_2(const chunk<1>& chnk) const override { return map<2>(chnk); }
                region<2> map_2(const chunk<2>& chnk) const override { return map<2>(chnk); }
                region<2> map_2(const chunk<3>& chnk) const override { return map<2>(chnk); }
                region<3> map_3(const chunk<0>& chnk) const override { return map<3>(chnk); }
                region<3> map_3(const chunk<1>& chnk) const override { return map<3>(chnk); }
                region<3> map_3(const chunk<2>& chnk) const override { return map<3>(chnk); }
                region<3> map_3(const chunk<3>& chnk) const override { return map<3>(chnk); }

          private:
                Functor m_rmfn;
                range<BufferDims> m_buffer_size;

                template <int OtherBufferDims, int KernelDims>
                region<OtherBufferDims> map(const chunk<KernelDims>& chnk) const {
                        if constexpr(OtherBufferDims == BufferDims) {
                                const auto r = invoke_range_mapper(m_rmfn, chnk, m_buffer_size);
                                return clamp_region_to_buffer_size(r, m_buffer_size);
                        } else {
                                throw_invalid_range_mapper_result(OtherBufferDims, BufferDims, KernelDims);
                        }
                }
        };

        struct range_mapper_testspy;

} // namespace detail


// --------------------------- Convenience range mappers ---------------------------

/// Optional parameter to the constructor of `access::neighborhood` to specify in what shape the accessed region should extend from the work item.
enum class neighborhood_shape {
        along_axes,   ///< The neighborhood extends along each axis separately, but not diagonally (in 2D, into a "+" shape).
        bounding_box, ///< The neighborhood extends along in all dimensions simultaneously to produce a single bounding box.
};

namespace access {

        struct one_to_one {
                template <int Dims>
                subrange<Dims> operator()(const chunk<Dims>& chnk) const {
                        return chnk;
                }
        };

        template <int BufferDims>
        struct fixed {
                fixed(const subrange<BufferDims>& sr) : m_sr(sr) {}

                template <int KernelDims>
                subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */) const {
                        return m_sr;
                }

          private:
                subrange<BufferDims> m_sr;
        };

        template <int Dims>
        struct slice {
                explicit slice(const size_t dim_idx) : m_dim_idx(dim_idx) { assert(dim_idx < Dims && "Invalid slice dimension index (starts at 0)"); }

                subrange<Dims> operator()(const chunk<Dims>& chnk, const range<Dims>& buffer_size) const {
                        subrange<Dims> result = chnk;
                        result.offset[m_dim_idx] = 0;
                        result.range[m_dim_idx] = buffer_size[m_dim_idx];
                        return result;
                }

          private:
                size_t m_dim_idx;
        };

        struct all {
                template <int KernelDims, int BufferDims>
                subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */, const range<BufferDims>& buffer_size) const {
                        return {{}, buffer_size};
                }
        };

        /// Declares a buffer access that extends from the current work item by a symmetric boundary offset, either in all directions in the shape of a bounding
        /// box (default), or along each axis separately (without "diagonal" boundary elements). Buffer and kernel dimensions must both match the `Dims` parameter.
        ///
        /// This is typically used in stencil applications. For bounding-box neighborhoods,
        /// - `neighborhood({1})` declares the read for a 1D 3-point stencil,
        /// - `neighborhood({1, 1})` for a 2D 9-point stencil, and
        /// - `neighborhood({1, 1, 1})` for a 3D 27-point stencil.
        ///
        /// For neighborhoods defined along axes only,
        /// - `neighborhood({1}, neighborhood_shape::along_axes)` declares the read for a 1D 3-point stencil,
        /// - `neighborhood({1, 1}, neighborhood_shape::along_axes)` for a 2D 5-point stencil, and
        /// - `neighborhood({1, 1, 1}, neighborhood_shape::along_axes)` for a 3D 7-point stencil.
        ///
        /// For reads, `neighborhood_shape::bounding_box` is functionally correct whenever `neighborhood_shape::along_axes` is, but will lead to unnecessary copies
        /// and transfers between diagonal neighbors in 2D-split work assignments when the application does not actually read from those buffer elements.
        template <int Dims>
        struct neighborhood {
                explicit neighborhood(const range<Dims>& extent, const neighborhood_shape shape = neighborhood_shape::bounding_box)
                    : m_extent(extent), m_shape(shape) {}

                [[deprecated("Use the neighborhood({a, b} [, shape]) instead of neighborhood(a, b)")]]
                explicit neighborhood(const size_t dim0, const size_t dim1)
                    requires(Dims == 2)
                    : neighborhood({dim0, dim1}) {}

                [[deprecated("Use the neighborhood({a, b, c} [, shape]) instead of neighborhood(a, b, c)")]]
                explicit neighborhood(const size_t dim0, const size_t dim1, const size_t dim2)
                    requires(Dims == 3)
                    : neighborhood({dim0, dim1, dim2}) {}

                detail::region<Dims> operator()(const chunk<Dims>& chnk) const {
                        const detail::box interior(subrange(chnk.offset, chnk.range));
                        detail::box_vector<Dims> boxes;
                        if(m_shape == neighborhood_shape::along_axes) {
                                boxes.push_back(interior);
                                for(int d = 0; d < Dims; ++d) {
                                        boxes.push_back(extend_axis(interior, d));
                                }
                        } else {
                                auto& bounding_box = boxes.emplace_back(interior);
                                for(int d = 0; d < Dims; ++d) {
                                        bounding_box = extend_axis(bounding_box, d);
                                }
                        }
                        return detail::region(std::move(boxes));
                }

          private:
                friend struct celerity::detail::range_mapper_testspy;

                range<Dims> m_extent;
                neighborhood_shape m_shape;

                inline detail::box<Dims> extend_axis(const detail::box<Dims>& box, const int d) const {
                        auto min = box.get_min();
                        auto max = box.get_max();
                        min[d] -= std::min(m_extent[d], min[d]);
                        max[d] += std::min(m_extent[d], std::numeric_limits<size_t>::max() - max[d]);
                        return detail::box(min, max);
                }
        };

        neighborhood(size_t) -> neighborhood<1>;
        neighborhood(size_t, size_t) -> neighborhood<2>;
        neighborhood(size_t, size_t, size_t) -> neighborhood<3>;

        // Explicit CTAD guides allow deducing Dims from `neighborhood{{1, 1}}`.
        neighborhood(range<1>) -> neighborhood<1>;
        neighborhood(range<2>) -> neighborhood<2>;
        neighborhood(range<3>) -> neighborhood<3>;
        neighborhood(range<1>, neighborhood_shape) -> neighborhood<1>;
        neighborhood(range<2>, neighborhood_shape) -> neighborhood<2>;
        neighborhood(range<3>, neighborhood_shape) -> neighborhood<3>;

} // namespace access

namespace experimental::access {

        /**
         * For a 1D kernel, splits an nD-buffer evenly along its slowest dimension.
         *
         * This range mapper is unique in the sense that the chunk parameter (i.e. the iteration space) is unrelated to the buffer indices it maps to.
         * It is designed to distribute a buffer in contiguous portions between nodes for collective host tasks, allowing each node to output its portion in
         * I/O operations. See `accessor::get_allocation_window` on how to access the resulting host memory.
         */
        template <int BufferDims>
        class even_split {
                static_assert(BufferDims > 0);

          public:
                even_split() = default;
                explicit even_split(const range<BufferDims>& granularity) : m_granularity(granularity) {}

                subrange<BufferDims> operator()(const chunk<1>& chunk, const range<BufferDims>& buffer_size) const {
                        if(chunk.global_size[0] == 0) { return {}; }

                        // Equal splitting has edge cases when buffer_size is not a multiple of global_size * granularity. Splitting is performed in a manner that
                        // distributes the remainder as equally as possible while adhering to granularity. In case buffer_size is not even a multiple of granularity,
                        // only last chunk should be oddly sized so that only one node needs to deal with misaligned buffers.

                        // 1. Each slice has at least buffer_size / global_size items, rounded down to the nearest multiple of the granularity.
                        // 2. The first chunks in the range receive one additional granularity-sized block each to distribute most of the remainder
                        // 3. The last chunk additionally receives the not-granularity-sized part of the remainder, if any.

                        auto dim0_step = buffer_size[0] / (chunk.global_size[0] * m_granularity[0]) * m_granularity[0];
                        auto dim0_remainder = buffer_size[0] - chunk.global_size[0] * dim0_step;
                        auto dim0_range_in_this_chunk = chunk.range[0] * dim0_step;
                        auto sum_dim0_remainder_in_prev_chunks = std::min(dim0_remainder / m_granularity[0] * m_granularity[0], chunk.offset[0] * m_granularity[0]);
                        if(dim0_remainder > sum_dim0_remainder_in_prev_chunks) {
                                dim0_range_in_this_chunk +=
                                    std::min(chunk.range[0], (dim0_remainder - sum_dim0_remainder_in_prev_chunks) / m_granularity[0]) * m_granularity[0];
                                if(chunk.offset[0] + chunk.range[0] == chunk.global_size[0]) { dim0_range_in_this_chunk += dim0_remainder % m_granularity[0]; }
                        }
                        auto dim0_offset_in_this_chunk = chunk.offset[0] * dim0_step + sum_dim0_remainder_in_prev_chunks;

                        subrange<BufferDims> sr;
                        sr.offset[0] = dim0_offset_in_this_chunk;
                        sr.range = buffer_size;
                        sr.range[0] = dim0_range_in_this_chunk;

                        return sr;
                }

          private:
                range<BufferDims> m_granularity = detail::range_cast<BufferDims>(range<3>(1, 1, 1));
        };

} // namespace experimental::access

} // namespace celerity

1	#pragma once
2
3	#include <stdexcept>
4	#include <type_traits>
5
6	#include <fmt/format.h>
7	#include <sycl/sycl.hpp>
8
9	#include "grid.h"
10	#include "ranges.h"
11	#include "utils.h"
12
13
14	namespace celerity {
15	namespace detail {
16
17	template <typename Functor, int BufferDims, int KernelDims>
18	constexpr bool is_range_mapper_invocable_for_chunk_only = std::is_invocable_r_v<region<BufferDims>, const Functor&, const celerity::chunk<KernelDims>&>;
19
20	template <typename Functor, int BufferDims, int KernelDims>
21	constexpr bool is_range_mapper_invocable_for_chunk_and_global_size =
22	std::is_invocable_r_v<region<BufferDims>, const Functor&, const celerity::chunk<KernelDims>&, const range<BufferDims>&>;
23
24	template <typename Functor, int BufferDims, int KernelDims>
25	constexpr bool is_range_mapper_invocable_for_kernel = is_range_mapper_invocable_for_chunk_only<Functor, BufferDims, KernelDims> //
26	\|\| is_range_mapper_invocable_for_chunk_and_global_size<Functor, BufferDims, KernelDims>;
27
28	template <typename Functor, int BufferDims>
29	constexpr bool is_range_mapper_invocable = is_range_mapper_invocable_for_kernel<Functor, BufferDims, 0> //
30	\|\| is_range_mapper_invocable_for_kernel<Functor, BufferDims, 1> //
31	\|\| is_range_mapper_invocable_for_kernel<Functor, BufferDims, 2> //
32	\|\| is_range_mapper_invocable_for_kernel<Functor, BufferDims, 3>;
33
34	[[noreturn]] inline void throw_invalid_range_mapper_args(int expect_kernel_dims, int expect_buffer_dims) {	2✔
35	throw std::runtime_error(fmt::format("Invalid range mapper dimensionality: {0}-dimensional kernel submitted with a requirement whose range mapper "	4✔
36	"is neither invocable for chunk<{0}> nor (chunk<{0}>, range<{1}>) to produce subrange<{1}>",
37	expect_kernel_dims, expect_buffer_dims));	6✔
38	}
39
40	[[noreturn]] inline void throw_invalid_range_mapper_result(int expect_sr_dims, int actual_sr_dims, int kernel_dims) {	×
41	throw std::runtime_error(fmt::format("Range mapper produces subrange of wrong dimensionality: Expecting subrange<{}>, got subrange<{}> for chunk<{}>",	×
42	expect_sr_dims, actual_sr_dims, kernel_dims));	×
43	}
44
45	template <typename Functor, int KernelDims, int BufferDims>
46	region<BufferDims> invoke_range_mapper(Functor&& fn, const celerity::chunk<KernelDims>& chunk, const range<BufferDims>& buffer_size) {	14,804✔
47	static_assert(KernelDims >= 0 && KernelDims <= 3 && BufferDims >= 0 && BufferDims <= 3);
48	if constexpr(is_range_mapper_invocable_for_chunk_and_global_size<Functor, BufferDims, KernelDims>) {
49	return std::forward<Functor>(fn)(chunk, buffer_size);	6,482✔
50	} else if constexpr(is_range_mapper_invocable_for_chunk_only<Functor, BufferDims, KernelDims>) {
51	return std::forward<Functor>(fn)(chunk);	22,754✔
52	} else {
53	throw_invalid_range_mapper_args(KernelDims, BufferDims);	2✔
54	}
55	}
56
57	template <int BufferDims>
58	region<BufferDims> clamp_region_to_buffer_size(const region<BufferDims>& r, const range<BufferDims>& buffer_size) {	14,802✔
59	return region_intersection(r, box<BufferDims>::full_range(buffer_size));	29,603✔
60	}
61
62	class range_mapper_base {
63	public:
64	explicit range_mapper_base(sycl::access::mode am) : m_access_mode(am) {}	2,711✔
65
66	range_mapper_base(const range_mapper_base& other) = delete;
67	range_mapper_base(range_mapper_base&& other) = delete;
68	range_mapper_base& operator=(const range_mapper_base& other) = delete;
69	range_mapper_base& operator=(range_mapper_base&& other) = delete;
70
71	sycl::access::mode get_access_mode() const { return m_access_mode; }	72,333✔
72
73	virtual int get_buffer_dimensions() const = 0;
74
75	virtual region<1> map_1(const chunk<0>& chnk) const = 0;
76	virtual region<1> map_1(const chunk<1>& chnk) const = 0;
77	virtual region<1> map_1(const chunk<2>& chnk) const = 0;
78	virtual region<1> map_1(const chunk<3>& chnk) const = 0;
79	virtual region<2> map_2(const chunk<0>& chnk) const = 0;
80	virtual region<2> map_2(const chunk<1>& chnk) const = 0;
81	virtual region<2> map_2(const chunk<2>& chnk) const = 0;
82	virtual region<2> map_2(const chunk<3>& chnk) const = 0;
83	virtual region<3> map_3(const chunk<0>& chnk) const = 0;
84	virtual region<3> map_3(const chunk<1>& chnk) const = 0;
85	virtual region<3> map_3(const chunk<2>& chnk) const = 0;
86	virtual region<3> map_3(const chunk<3>& chnk) const = 0;
87
88	virtual ~range_mapper_base() = default;	2,711✔
89
90	private:
91	sycl::access::mode m_access_mode;
92	};
93
94	template <int BufferDims, typename Functor>
95	class range_mapper final : public range_mapper_base {
96	public:
97	range_mapper(Functor rmfn, sycl::access::mode am, range<BufferDims> buffer_size) : range_mapper_base(am), m_rmfn(rmfn), m_buffer_size(buffer_size) {}	2,711✔
98
99	int get_buffer_dimensions() const override { return BufferDims; }	27,990✔
100
101	region<1> map_1(const chunk<0>& chnk) const override { return map<1>(chnk); }	1,330✔
102	region<1> map_1(const chunk<1>& chnk) const override { return map<1>(chnk); }	5,655✔
103	region<1> map_1(const chunk<2>& chnk) const override { return map<1>(chnk); }	228✔
NEW 104	region<1> map_1(const chunk<3>& chnk) const override { return map<1>(chnk); }	×
105	region<2> map_2(const chunk<0>& chnk) const override { return map<2>(chnk); }	455✔
106	region<2> map_2(const chunk<1>& chnk) const override { return map<2>(chnk); }	3,623✔
107	region<2> map_2(const chunk<2>& chnk) const override { return map<2>(chnk); }	2,360✔
108	region<2> map_2(const chunk<3>& chnk) const override { return map<2>(chnk); }	50✔
109	region<3> map_3(const chunk<0>& chnk) const override { return map<3>(chnk); }	278✔
110	region<3> map_3(const chunk<1>& chnk) const override { return map<3>(chnk); }	137✔
111	region<3> map_3(const chunk<2>& chnk) const override { return map<3>(chnk); }	3✔
112	region<3> map_3(const chunk<3>& chnk) const override { return map<3>(chnk); }	685✔
113
114	private:
115	Functor m_rmfn;
116	range<BufferDims> m_buffer_size;
117
118	template <int OtherBufferDims, int KernelDims>
119	region<OtherBufferDims> map(const chunk<KernelDims>& chnk) const {	14,804✔
120	if constexpr(OtherBufferDims == BufferDims) {
121	const auto r = invoke_range_mapper(m_rmfn, chnk, m_buffer_size);	14,804✔
122	return clamp_region_to_buffer_size(r, m_buffer_size);	29,603✔
123	} else {	14,801✔
124	throw_invalid_range_mapper_result(OtherBufferDims, BufferDims, KernelDims);	×
125	}
126	}
127	};
128
129	struct range_mapper_testspy;
130
131	} // namespace detail
132
133
134	// --------------------------- Convenience range mappers ---------------------------
135
136	/// Optional parameter to the constructor of `access::neighborhood` to specify in what shape the accessed region should extend from the work item.
137	enum class neighborhood_shape {
138	along_axes, ///< The neighborhood extends along each axis separately, but not diagonally (in 2D, into a "+" shape).
139	bounding_box, ///< The neighborhood extends along in all dimensions simultaneously to produce a single bounding box.
140	};
141
142	namespace access {
143
144	struct one_to_one {
145	template <int Dims>
146	subrange<Dims> operator()(const chunk<Dims>& chnk) const {	5,039✔
147	return chnk;	5,039✔
148	}
149	};
150
151	template <int BufferDims>
152	struct fixed {
153	fixed(const subrange<BufferDims>& sr) : m_sr(sr) {}	248✔
154
155	template <int KernelDims>
156	subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */) const {	1,848✔
157	return m_sr;	1,848✔
158	}
159
160	private:
161	subrange<BufferDims> m_sr;
162	};
163
164	template <int Dims>
165	struct slice {
166	explicit slice(const size_t dim_idx) : m_dim_idx(dim_idx) { assert(dim_idx < Dims && "Invalid slice dimension index (starts at 0)"); }	3!
167
168	subrange<Dims> operator()(const chunk<Dims>& chnk, const range<Dims>& buffer_size) const {	3✔
169	subrange<Dims> result = chnk;	3✔
170	result.offset[m_dim_idx] = 0;	3✔
171	result.range[m_dim_idx] = buffer_size[m_dim_idx];	3✔
172	return result;	3✔
173	}
174
175	private:
176	size_t m_dim_idx;
177	};
178
179	struct all {
180	template <int KernelDims, int BufferDims>
181	subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */, const range<BufferDims>& buffer_size) const {	3,231✔
182	return {{}, buffer_size};	3,970✔
183	}
184	};
185
186	/// Declares a buffer access that extends from the current work item by a symmetric boundary offset, either in all directions in the shape of a bounding
187	/// box (default), or along each axis separately (without "diagonal" boundary elements). Buffer and kernel dimensions must both match the `Dims` parameter.
188	///
189	/// This is typically used in stencil applications. For bounding-box neighborhoods,
190	/// - `neighborhood({1})` declares the read for a 1D 3-point stencil,
191	/// - `neighborhood({1, 1})` for a 2D 9-point stencil, and
192	/// - `neighborhood({1, 1, 1})` for a 3D 27-point stencil.
193	///
194	/// For neighborhoods defined along axes only,
195	/// - `neighborhood({1}, neighborhood_shape::along_axes)` declares the read for a 1D 3-point stencil,
196	/// - `neighborhood({1, 1}, neighborhood_shape::along_axes)` for a 2D 5-point stencil, and
197	/// - `neighborhood({1, 1, 1}, neighborhood_shape::along_axes)` for a 3D 7-point stencil.
198	///
199	/// For reads, `neighborhood_shape::bounding_box` is functionally correct whenever `neighborhood_shape::along_axes` is, but will lead to unnecessary copies
200	/// and transfers between diagonal neighbors in 2D-split work assignments when the application does not actually read from those buffer elements.
201	template <int Dims>
202	struct neighborhood {
203	explicit neighborhood(const range<Dims>& extent, const neighborhood_shape shape = neighborhood_shape::bounding_box)	28✔
204	: m_extent(extent), m_shape(shape) {}	28✔
205
206	[[deprecated("Use the neighborhood({a, b} [, shape]) instead of neighborhood(a, b)")]]
207	explicit neighborhood(const size_t dim0, const size_t dim1)	1✔
208	requires(Dims == 2)
209	: neighborhood({dim0, dim1}) {}	1✔
210
211	[[deprecated("Use the neighborhood({a, b, c} [, shape]) instead of neighborhood(a, b, c)")]]
212	explicit neighborhood(const size_t dim0, const size_t dim1, const size_t dim2)	1✔
213	requires(Dims == 3)
214	: neighborhood({dim0, dim1, dim2}) {}	1✔
215
216	detail::region<Dims> operator()(const chunk<Dims>& chnk) const {	335✔
217	const detail::box interior(subrange(chnk.offset, chnk.range));	335✔
218	detail::box_vector<Dims> boxes;	335✔
219	if(m_shape == neighborhood_shape::along_axes) {	335✔
220	boxes.push_back(interior);	97✔
221	for(int d = 0; d < Dims; ++d) {	291✔
222	boxes.push_back(extend_axis(interior, d));	194✔
223	}
224	} else {
225	auto& bounding_box = boxes.emplace_back(interior);	238✔
226	for(int d = 0; d < Dims; ++d) {	714✔
227	bounding_box = extend_axis(bounding_box, d);	476✔
228	}
229	}
230	return detail::region(std::move(boxes));	670✔
231	}	335✔
232
233	private:
234	friend struct celerity::detail::range_mapper_testspy;
235
236	range<Dims> m_extent;
237	neighborhood_shape m_shape;
238
239	inline detail::box<Dims> extend_axis(const detail::box<Dims>& box, const int d) const {	670✔
240	auto min = box.get_min();	670✔
241	auto max = box.get_max();	670✔
242	min[d] -= std::min(m_extent[d], min[d]);	670✔
243	max[d] += std::min(m_extent[d], std::numeric_limits<size_t>::max() - max[d]);	670✔
244	return detail::box(min, max);	1,340✔
245	}
246	};
247
248	neighborhood(size_t) -> neighborhood<1>;
249	neighborhood(size_t, size_t) -> neighborhood<2>;
250	neighborhood(size_t, size_t, size_t) -> neighborhood<3>;
251
252	// Explicit CTAD guides allow deducing Dims from `neighborhood{{1, 1}}`.
253	neighborhood(range<1>) -> neighborhood<1>;
254	neighborhood(range<2>) -> neighborhood<2>;
255	neighborhood(range<3>) -> neighborhood<3>;
256	neighborhood(range<1>, neighborhood_shape) -> neighborhood<1>;
257	neighborhood(range<2>, neighborhood_shape) -> neighborhood<2>;
258	neighborhood(range<3>, neighborhood_shape) -> neighborhood<3>;
259
260	} // namespace access
261
262	namespace experimental::access {
263
264	/**
265	* For a 1D kernel, splits an nD-buffer evenly along its slowest dimension.
266	*
267	* This range mapper is unique in the sense that the chunk parameter (i.e. the iteration space) is unrelated to the buffer indices it maps to.
268	* It is designed to distribute a buffer in contiguous portions between nodes for collective host tasks, allowing each node to output its portion in
269	* I/O operations. See `accessor::get_allocation_window` on how to access the resulting host memory.
270	*/
271	template <int BufferDims>
272	class even_split {
273	static_assert(BufferDims > 0);
274
275	public:
276	even_split() = default;	3✔
277	explicit even_split(const range<BufferDims>& granularity) : m_granularity(granularity) {}	4✔
278
279	subrange<BufferDims> operator()(const chunk<1>& chunk, const range<BufferDims>& buffer_size) const {	7✔
280	if(chunk.global_size[0] == 0) { return {}; }	7✔
281
282	// Equal splitting has edge cases when buffer_size is not a multiple of global_size * granularity. Splitting is performed in a manner that
283	// distributes the remainder as equally as possible while adhering to granularity. In case buffer_size is not even a multiple of granularity,
284	// only last chunk should be oddly sized so that only one node needs to deal with misaligned buffers.
285
286	// 1. Each slice has at least buffer_size / global_size items, rounded down to the nearest multiple of the granularity.
287	// 2. The first chunks in the range receive one additional granularity-sized block each to distribute most of the remainder
288	// 3. The last chunk additionally receives the not-granularity-sized part of the remainder, if any.
289
290	auto dim0_step = buffer_size[0] / (chunk.global_size[0] * m_granularity[0]) * m_granularity[0];	7✔
291	auto dim0_remainder = buffer_size[0] - chunk.global_size[0] * dim0_step;	7✔
292	auto dim0_range_in_this_chunk = chunk.range[0] * dim0_step;	7✔
293	auto sum_dim0_remainder_in_prev_chunks = std::min(dim0_remainder / m_granularity[0] * m_granularity[0], chunk.offset[0] * m_granularity[0]);	7✔
294	if(dim0_remainder > sum_dim0_remainder_in_prev_chunks) {	7✔
295	dim0_range_in_this_chunk +=	4✔
296	std::min(chunk.range[0], (dim0_remainder - sum_dim0_remainder_in_prev_chunks) / m_granularity[0]) * m_granularity[0];	4✔
297	if(chunk.offset[0] + chunk.range[0] == chunk.global_size[0]) { dim0_range_in_this_chunk += dim0_remainder % m_granularity[0]; }	4✔
298	}
299	auto dim0_offset_in_this_chunk = chunk.offset[0] * dim0_step + sum_dim0_remainder_in_prev_chunks;	7✔
300
301	subrange<BufferDims> sr;	7✔
302	sr.offset[0] = dim0_offset_in_this_chunk;	7✔
303	sr.range = buffer_size;	7✔
304	sr.range[0] = dim0_range_in_this_chunk;	7✔
305
306	return sr;	7✔
307	}
308
309	private:
310	range<BufferDims> m_granularity = detail::range_cast<BufferDims>(range<3>(1, 1, 1));
311	};
312
313	} // namespace experimental::access
314
315	} // namespace celerity

celerity / celerity-runtime / 11631827009

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous