11841839434

Committed 14 Nov 2024 04:54PM UTC coverage: 95.154% (-0.02%) from 95.176%

Build # 11841839434

Build Type

push

github

Committed by

fknorr

Commit Message

Update benchmark results for epoch refactoring

Run Details

3018 of 3414 branches covered (88.4%)

Branch coverage included in aggregate %.

6682 of 6780 relevant lines covered (98.55%)

1298952.13 hits per line

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

79.41

/include/task.h

#pragma once

#include <memory>
#include <unordered_map>
#include <unordered_set>
#include <utility>
#include <vector>

#include "graph.h"
#include "grid.h"
#include "hint.h"
#include "intrusive_graph.h"
#include "launcher.h"
#include "range_mapper.h"
#include "reduction.h"
#include "types.h"


namespace celerity {

class handler;

namespace detail {

        class buffer_access_map {
          public:
                void add_access(buffer_id bid, std::unique_ptr<range_mapper_base>&& rm) { m_accesses.emplace_back(bid, std::move(rm)); }

                std::unordered_set<buffer_id> get_accessed_buffers() const;
                std::unordered_set<sycl::access::mode> get_access_modes(buffer_id bid) const;
                size_t get_num_accesses() const { return m_accesses.size(); }
                std::pair<buffer_id, access_mode> get_nth_access(const size_t n) const {
                        const auto& [bid, rm] = m_accesses[n];
                        return {bid, rm->get_access_mode()};
                }

                /**
                 * @brief Computes the combined access-region for a given buffer, mode and subrange.
                 *
                 * @param bid
                 * @param mode
                 * @param sr The subrange to be passed to the range mappers (extended to a chunk using the global size of the task)
                 *
                 * @returns The region obtained by merging the results of all range-mappers for this buffer and mode
                 */
                region<3> get_mode_requirements(
                    const buffer_id bid, const access_mode mode, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;

                region<3> get_requirements_for_nth_access(const size_t n, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;

                std::vector<const range_mapper_base*> get_range_mappers(const buffer_id bid) const {
                        std::vector<const range_mapper_base*> rms;
                        for(const auto& [a_bid, a_rm] : m_accesses) {
                                if(a_bid == bid) { rms.push_back(a_rm.get()); }
                        }
                        return rms;
                }

                box_vector<3> get_required_contiguous_boxes(const buffer_id bid, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;

          private:
                std::vector<std::pair<buffer_id, std::unique_ptr<range_mapper_base>>> m_accesses;
        };

        using reduction_set = std::vector<reduction_info>;

        class side_effect_map : private std::unordered_map<host_object_id, experimental::side_effect_order> {
          private:
                using map_base = std::unordered_map<host_object_id, experimental::side_effect_order>;

          public:
                using typename map_base::const_iterator, map_base::value_type, map_base::key_type, map_base::mapped_type, map_base::const_reference,
                    map_base::const_pointer;
                using iterator = const_iterator;
                using reference = const_reference;
                using pointer = const_pointer;

                using map_base::size, map_base::count, map_base::empty, map_base::cbegin, map_base::cend, map_base::at;

                iterator begin() const { return cbegin(); }
                iterator end() const { return cend(); }
                iterator find(host_object_id key) const { return map_base::find(key); }

                void add_side_effect(host_object_id hoid, experimental::side_effect_order order);
        };

        class task_promise {
          public:
                task_promise() = default;
                task_promise(const task_promise&) = delete;
                task_promise(task_promise&&) = delete;
                task_promise& operator=(const task_promise&) = delete;
                task_promise& operator=(task_promise&&) = delete;
                virtual ~task_promise() = default;

                virtual void fulfill() = 0;
                virtual allocation_id get_user_allocation_id() = 0; // TODO move to struct task instead
        };

        struct task_geometry {
                int dimensions = 0;
                range<3> global_size{1, 1, 1};
                id<3> global_offset{};
                range<3> granularity{1, 1, 1};
        };

        class task : public intrusive_graph_node<task> {
          public:
                task_type get_type() const { return m_type; }

                task_id get_id() const { return m_tid; }

                collective_group_id get_collective_group_id() const { return m_cgid; }

                const buffer_access_map& get_buffer_access_map() const { return m_access_map; }

                const side_effect_map& get_side_effect_map() const { return m_side_effects; }

                const task_geometry& get_geometry() const { return m_geometry; }

                int get_dimensions() const { return m_geometry.dimensions; }

                range<3> get_global_size() const { return m_geometry.global_size; }

                id<3> get_global_offset() const { return m_geometry.global_offset; }

                range<3> get_granularity() const { return m_geometry.granularity; }

                void set_debug_name(const std::string& debug_name) { m_debug_name = debug_name; }
                const std::string& get_debug_name() const { return m_debug_name; }

                bool has_variable_split() const { return m_type == task_type::host_compute || m_type == task_type::device_compute; }

                execution_target get_execution_target() const {
                        switch(m_type) {
                        case task_type::epoch: return execution_target::none;
                        case task_type::device_compute: return execution_target::device;
                        case task_type::host_compute:
                        case task_type::collective:
                        case task_type::master_node: return execution_target::host;
                        case task_type::horizon:
                        case task_type::fence: return execution_target::none;
                        default: utils::unreachable(); // LCOV_EXCL_LINE
                        }
                }

                const reduction_set& get_reductions() const { return m_reductions; }

                epoch_action get_epoch_action() const { return m_epoch_action; }

                task_promise* get_task_promise() const { return m_promise.get(); }

                template <typename Launcher>
                Launcher get_launcher() const {
                        return std::get<Launcher>(m_launcher);
                }

                void add_hint(std::unique_ptr<hint_base>&& h) { m_hints.emplace_back(std::move(h)); }

                template <typename Hint>
                const Hint* get_hint() const {
                        static_assert(std::is_base_of_v<hint_base, Hint>, "Hint must extend hint_base");
                        for(auto& h : m_hints) {
                                if(auto* ptr = dynamic_cast<Hint*>(h.get()); ptr != nullptr) { return ptr; }
                        }
                        return nullptr;
                }

                static std::unique_ptr<task> make_epoch(task_id tid, detail::epoch_action action, std::unique_ptr<task_promise> promise) {
                        return std::unique_ptr<task>(new task(tid, task_type::epoch, non_collective_group_id, task_geometry{}, {}, {}, {}, {}, action, std::move(promise)));
                }

                static std::unique_ptr<task> make_host_compute(task_id tid, task_geometry geometry, host_task_launcher launcher, buffer_access_map access_map,
                    side_effect_map side_effect_map, reduction_set reductions) {
                        return std::unique_ptr<task>(new task(tid, task_type::host_compute, non_collective_group_id, geometry, std::move(launcher), std::move(access_map),
                            std::move(side_effect_map), std::move(reductions), {}, nullptr));
                }

                static std::unique_ptr<task> make_device_compute(
                    task_id tid, task_geometry geometry, device_kernel_launcher launcher, buffer_access_map access_map, reduction_set reductions) {
                        return std::unique_ptr<task>(new task(tid, task_type::device_compute, non_collective_group_id, geometry, std::move(launcher), std::move(access_map),
                            {}, std::move(reductions), {}, nullptr));
                }

                static std::unique_ptr<task> make_collective(task_id tid, collective_group_id cgid, size_t num_collective_nodes, host_task_launcher launcher,
                    buffer_access_map access_map, side_effect_map side_effect_map) {
                        const task_geometry geometry{1, detail::range_cast<3>(range(num_collective_nodes)), {}, {1, 1, 1}};
                        return std::unique_ptr<task>(
                            new task(tid, task_type::collective, cgid, geometry, std::move(launcher), std::move(access_map), std::move(side_effect_map), {}, {}, nullptr));
                }

                static std::unique_ptr<task> make_master_node(task_id tid, host_task_launcher launcher, buffer_access_map access_map, side_effect_map side_effect_map) {
                        return std::unique_ptr<task>(new task(tid, task_type::master_node, non_collective_group_id, task_geometry{}, std::move(launcher),
                            std::move(access_map), std::move(side_effect_map), {}, {}, nullptr));
                }

                static std::unique_ptr<task> make_horizon(task_id tid) {
                        return std::unique_ptr<task>(new task(tid, task_type::horizon, non_collective_group_id, task_geometry{}, {}, {}, {}, {}, {}, nullptr));
                }

                static std::unique_ptr<task> make_fence(
                    task_id tid, buffer_access_map access_map, side_effect_map side_effect_map, std::unique_ptr<task_promise> promise) {
                        return std::unique_ptr<task>(new task(tid, task_type::fence, non_collective_group_id, task_geometry{}, {}, std::move(access_map),
                            std::move(side_effect_map), {}, {}, std::move(promise)));
                }

          private:
                task_id m_tid;
                task_type m_type;
                collective_group_id m_cgid;
                task_geometry m_geometry;
                command_group_launcher m_launcher;
                buffer_access_map m_access_map;
                detail::side_effect_map m_side_effects;
                reduction_set m_reductions;
                std::string m_debug_name;
                detail::epoch_action m_epoch_action;
                std::unique_ptr<task_promise> m_promise; // TODO keep user_allocation_id in struct task instead of inside task_promise
                std::vector<std::unique_ptr<hint_base>> m_hints;

                task(task_id tid, task_type type, collective_group_id cgid, task_geometry geometry, command_group_launcher launcher, buffer_access_map access_map,
                    detail::side_effect_map side_effects, reduction_set reductions, detail::epoch_action epoch_action, std::unique_ptr<task_promise> promise)
                    : m_tid(tid), m_type(type), m_cgid(cgid), m_geometry(geometry), m_launcher(std::move(launcher)), m_access_map(std::move(access_map)),
                      m_side_effects(std::move(side_effects)), m_reductions(std::move(reductions)), m_epoch_action(epoch_action), m_promise(std::move(promise)) {
                        assert(type == task_type::host_compute || type == task_type::device_compute || get_granularity().size() == 1);
                        // Only host tasks can have side effects
                        assert(this->m_side_effects.empty() || type == task_type::host_compute || type == task_type::collective || type == task_type::master_node
                               || type == task_type::fence);
                }
        };

        [[nodiscard]] std::string print_task_debug_label(const task& tsk, bool title_case = false);

        /// Determines which overlapping regions appear between write accesses when the iteration space of `tsk` is split into `chunks`.
        std::unordered_map<buffer_id, region<3>> detect_overlapping_writes(const task& tsk, const box_vector<3>& chunks);

        /// The task graph (TDAG) represents all cluster-wide operations, such as command group submissions and fences, and their interdependencies.
        class task_graph : public graph<task> {}; // inheritance instead of type alias so we can forward declare task_graph

} // namespace detail
} // namespace celerity

1	#pragma once
2
3	#include <memory>
4	#include <unordered_map>
5	#include <unordered_set>
6	#include <utility>
7	#include <vector>
8
9	#include "graph.h"
10	#include "grid.h"
11	#include "hint.h"
12	#include "intrusive_graph.h"
13	#include "launcher.h"
14	#include "range_mapper.h"
15	#include "reduction.h"
16	#include "types.h"
17
18
19	namespace celerity {
20
21	class handler;
22
23	namespace detail {
24
25	class buffer_access_map {
26	public:
27	void add_access(buffer_id bid, std::unique_ptr<range_mapper_base>&& rm) { m_accesses.emplace_back(bid, std::move(rm)); }	2,689✔
28
29	std::unordered_set<buffer_id> get_accessed_buffers() const;
30	std::unordered_set<sycl::access::mode> get_access_modes(buffer_id bid) const;
31	size_t get_num_accesses() const { return m_accesses.size(); }	13,826✔
32	std::pair<buffer_id, access_mode> get_nth_access(const size_t n) const {	1,952✔
33	const auto& [bid, rm] = m_accesses[n];	1,952✔
34	return {bid, rm->get_access_mode()};	1,952✔
35	}
36
37	/**
38	* @brief Computes the combined access-region for a given buffer, mode and subrange.
39	*
40	* @param bid
41	* @param mode
42	* @param sr The subrange to be passed to the range mappers (extended to a chunk using the global size of the task)
43	*
44	* @returns The region obtained by merging the results of all range-mappers for this buffer and mode
45	*/
46	region<3> get_mode_requirements(
47	const buffer_id bid, const access_mode mode, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;
48
49	region<3> get_requirements_for_nth_access(const size_t n, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;
50
51	std::vector<const range_mapper_base*> get_range_mappers(const buffer_id bid) const {
52	std::vector<const range_mapper_base*> rms;
53	for(const auto& [a_bid, a_rm] : m_accesses) {
54	if(a_bid == bid) { rms.push_back(a_rm.get()); }
55	}
56	return rms;
57	}
58
59	box_vector<3> get_required_contiguous_boxes(const buffer_id bid, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const;
60
61	private:
62	std::vector<std::pair<buffer_id, std::unique_ptr<range_mapper_base>>> m_accesses;
63	};
64
65	using reduction_set = std::vector<reduction_info>;
66
67	class side_effect_map : private std::unordered_map<host_object_id, experimental::side_effect_order> {
68	private:
69	using map_base = std::unordered_map<host_object_id, experimental::side_effect_order>;
70
71	public:
72	using typename map_base::const_iterator, map_base::value_type, map_base::key_type, map_base::mapped_type, map_base::const_reference,
73	map_base::const_pointer;
74	using iterator = const_iterator;
75	using reference = const_reference;
76	using pointer = const_pointer;
77
78	using map_base::size, map_base::count, map_base::empty, map_base::cbegin, map_base::cend, map_base::at;
79
80	iterator begin() const { return cbegin(); }	7,744✔
81	iterator end() const { return cend(); }	7,725✔
82	iterator find(host_object_id key) const { return map_base::find(key); }
83
84	void add_side_effect(host_object_id hoid, experimental::side_effect_order order);
85	};
86
87	class task_promise {
88	public:
89	task_promise() = default;	589✔
90	task_promise(const task_promise&) = delete;
91	task_promise(task_promise&&) = delete;
92	task_promise& operator=(const task_promise&) = delete;
93	task_promise& operator=(task_promise&&) = delete;
94	virtual ~task_promise() = default;	589✔
95
96	virtual void fulfill() = 0;
97	virtual allocation_id get_user_allocation_id() = 0; // TODO move to struct task instead
98	};
99
100	struct task_geometry {
101	int dimensions = 0;
102	range<3> global_size{1, 1, 1};
103	id<3> global_offset{};
104	range<3> granularity{1, 1, 1};
105	};
106
107	class task : public intrusive_graph_node<task> {
108	public:
109	task_type get_type() const { return m_type; }	35,960✔
110
111	task_id get_id() const { return m_tid; }	37,827✔
112
113	collective_group_id get_collective_group_id() const { return m_cgid; }	14,953✔
114
115	const buffer_access_map& get_buffer_access_map() const { return m_access_map; }	39,072✔
116
117	const side_effect_map& get_side_effect_map() const { return m_side_effects; }	13,562✔
118
119	const task_geometry& get_geometry() const { return m_geometry; }	1,931✔
120
121	int get_dimensions() const { return m_geometry.dimensions; }	67,446✔
122
123	range<3> get_global_size() const { return m_geometry.global_size; }	95,284✔
124
125	id<3> get_global_offset() const { return m_geometry.global_offset; }	19,368✔
126
127	range<3> get_granularity() const { return m_geometry.granularity; }	6,903✔
128
129	void set_debug_name(const std::string& debug_name) { m_debug_name = debug_name; }	2,707✔
130	const std::string& get_debug_name() const { return m_debug_name; }	7,198✔
131
132	bool has_variable_split() const { return m_type == task_type::host_compute \|\| m_type == task_type::device_compute; }	4,963✔
133
134	execution_target get_execution_target() const {	9,411✔
135	switch(m_type) {	9,411!
136	case task_type::epoch: return execution_target::none;	×
137	case task_type::device_compute: return execution_target::device;	4,118✔
138	case task_type::host_compute:	5,293✔
139	case task_type::collective:
140	case task_type::master_node: return execution_target::host;	5,293✔
141	case task_type::horizon:	×
142	case task_type::fence: return execution_target::none;	×
143	default: utils::unreachable(); // LCOV_EXCL_LINE
144	}
145	}
146
147	const reduction_set& get_reductions() const { return m_reductions; }	68,308✔
148
149	epoch_action get_epoch_action() const { return m_epoch_action; }	3,960✔
150
151	task_promise* get_task_promise() const { return m_promise.get(); }	1,260✔
152
153	template <typename Launcher>
154	Launcher get_launcher() const {	2,606✔
155	return std::get<Launcher>(m_launcher);	2,606✔
156	}
157
158	void add_hint(std::unique_ptr<hint_base>&& h) { m_hints.emplace_back(std::move(h)); }	82✔
159
160	template <typename Hint>
161	const Hint* get_hint() const {	7,183✔
162	static_assert(std::is_base_of_v<hint_base, Hint>, "Hint must extend hint_base");
163	for(auto& h : m_hints) {	7,386✔
164	if(auto* ptr = dynamic_cast<Hint*>(h.get()); ptr != nullptr) { return ptr; }	319!
165	}
166	return nullptr;	7,067✔
167	}
168
169	static std::unique_ptr<task> make_epoch(task_id tid, detail::epoch_action action, std::unique_ptr<task_promise> promise) {	1,326✔
170	return std::unique_ptr<task>(new task(tid, task_type::epoch, non_collective_group_id, task_geometry{}, {}, {}, {}, {}, action, std::move(promise)));	3,978!
171	}
172
173	static std::unique_ptr<task> make_host_compute(task_id tid, task_geometry geometry, host_task_launcher launcher, buffer_access_map access_map,	336✔
174	side_effect_map side_effect_map, reduction_set reductions) {
175	return std::unique_ptr<task>(new task(tid, task_type::host_compute, non_collective_group_id, geometry, std::move(launcher), std::move(access_map),	672✔
176	std::move(side_effect_map), std::move(reductions), {}, nullptr));	1,008!
177	}
178
179	static std::unique_ptr<task> make_device_compute(	1,047✔
180	task_id tid, task_geometry geometry, device_kernel_launcher launcher, buffer_access_map access_map, reduction_set reductions) {
181	return std::unique_ptr<task>(new task(tid, task_type::device_compute, non_collective_group_id, geometry, std::move(launcher), std::move(access_map),	2,094✔
182	{}, std::move(reductions), {}, nullptr));	3,141!
183	}
184
185	static std::unique_ptr<task> make_collective(task_id tid, collective_group_id cgid, size_t num_collective_nodes, host_task_launcher launcher,	61✔
186	buffer_access_map access_map, side_effect_map side_effect_map) {
187	const task_geometry geometry{1, detail::range_cast<3>(range(num_collective_nodes)), {}, {1, 1, 1}};	61✔
188	return std::unique_ptr<task>(
189	new task(tid, task_type::collective, cgid, geometry, std::move(launcher), std::move(access_map), std::move(side_effect_map), {}, {}, nullptr));	122!
190	}
191
192	static std::unique_ptr<task> make_master_node(task_id tid, host_task_launcher launcher, buffer_access_map access_map, side_effect_map side_effect_map) {	1,264✔
193	return std::unique_ptr<task>(new task(tid, task_type::master_node, non_collective_group_id, task_geometry{}, std::move(launcher),	1,264✔
194	std::move(access_map), std::move(side_effect_map), {}, {}, nullptr));	5,056!
195	}
196
197	static std::unique_ptr<task> make_horizon(task_id tid) {	814✔
198	return std::unique_ptr<task>(new task(tid, task_type::horizon, non_collective_group_id, task_geometry{}, {}, {}, {}, {}, {}, nullptr));	2,442!
199	}
200
201	static std::unique_ptr<task> make_fence(	85✔
202	task_id tid, buffer_access_map access_map, side_effect_map side_effect_map, std::unique_ptr<task_promise> promise) {
203	return std::unique_ptr<task>(new task(tid, task_type::fence, non_collective_group_id, task_geometry{}, {}, std::move(access_map),	85✔
204	std::move(side_effect_map), {}, {}, std::move(promise)));	340!
205	}
206
207	private:
208	task_id m_tid;
209	task_type m_type;
210	collective_group_id m_cgid;
211	task_geometry m_geometry;
212	command_group_launcher m_launcher;
213	buffer_access_map m_access_map;
214	detail::side_effect_map m_side_effects;
215	reduction_set m_reductions;
216	std::string m_debug_name;
217	detail::epoch_action m_epoch_action;
218	std::unique_ptr<task_promise> m_promise; // TODO keep user_allocation_id in struct task instead of inside task_promise
219	std::vector<std::unique_ptr<hint_base>> m_hints;
220
221	task(task_id tid, task_type type, collective_group_id cgid, task_geometry geometry, command_group_launcher launcher, buffer_access_map access_map,	4,695✔
222	detail::side_effect_map side_effects, reduction_set reductions, detail::epoch_action epoch_action, std::unique_ptr<task_promise> promise)
223	: m_tid(tid), m_type(type), m_cgid(cgid), m_geometry(geometry), m_launcher(std::move(launcher)), m_access_map(std::move(access_map)),	4,695✔
224	m_side_effects(std::move(side_effects)), m_reductions(std::move(reductions)), m_epoch_action(epoch_action), m_promise(std::move(promise)) {	9,390✔
225	assert(type == task_type::host_compute \|\| type == task_type::device_compute \|\| get_granularity().size() == 1);	4,695✔
226	// Only host tasks can have side effects
227	assert(this->m_side_effects.empty() \|\| type == task_type::host_compute \|\| type == task_type::collective \|\| type == task_type::master_node	4,695✔
228	\|\| type == task_type::fence);
229	}	4,695✔
230	};
231
232	[[nodiscard]] std::string print_task_debug_label(const task& tsk, bool title_case = false);
233
234	/// Determines which overlapping regions appear between write accesses when the iteration space of `tsk` is split into `chunks`.
235	std::unordered_map<buffer_id, region<3>> detect_overlapping_writes(const task& tsk, const box_vector<3>& chunks);
236
237	/// The task graph (TDAG) represents all cluster-wide operations, such as command group submissions and fences, and their interdependencies.
238	class task_graph : public graph<task> {}; // inheritance instead of type alias so we can forward declare task_graph
239
240	} // namespace detail
241	} // namespace celerity

celerity / celerity-runtime / 11841839434

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