7783633344

Committed 05 Feb 2024 11:17AM UTC coverage: 93.942% (+0.02%) from 93.92%

Build # 7783633344

Build Type

Pull #230

github

Committed by

fknorr

Commit Message

Re-define transfer_id as struct{task_id, buffer_id, reduction_id}

Pull Request Pull Request #230: Re-define transfer_id as struct{task_id, buffer_id, reduction_id}

Run Details

2240 of 2514 branches covered (0.0%)

Branch coverage included in aggregate %.

63 of 63 new or added lines in 11 files covered. (100.0%)

1 existing line in 1 file now uncovered.

5002 of 5195 relevant lines covered (96.28%)

250164.07 hits per line

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

87.23

/src/task.cc

#include "task.h"

#include "access_modes.h"
#include "utils.h"


namespace celerity {
namespace detail {

        std::unordered_set<buffer_id> buffer_access_map::get_accessed_buffers() const {
                std::unordered_set<buffer_id> result;
                for(const auto& [bid, _] : m_accesses) {
                        result.emplace(bid);
                }
                return result;
        }

        std::unordered_set<cl::sycl::access::mode> buffer_access_map::get_access_modes(buffer_id bid) const {
                std::unordered_set<cl::sycl::access::mode> result;
                for(const auto& [b, rm] : m_accesses) {
                        if(b == bid) { result.insert(rm->get_access_mode()); }
                }
                return result;
        }

        template <int KernelDims>
        subrange<3> apply_range_mapper(const range_mapper_base* rm, const chunk<KernelDims>& chnk) {
                switch(rm->get_buffer_dimensions()) {
                case 0: return subrange_cast<3>(subrange<0>());
                case 1: return subrange_cast<3>(rm->map_1(chnk));
                case 2: return subrange_cast<3>(rm->map_2(chnk));
                case 3: return rm->map_3(chnk);
                default: assert(false);
                }
                return subrange<3>{};
        }

        region<3> buffer_access_map::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 {
                box_vector<3> boxes;
                for(size_t i = 0; i < m_accesses.size(); ++i) {
                        if(m_accesses[i].first != bid || m_accesses[i].second->get_access_mode() != mode) continue;
                        boxes.push_back(get_requirements_for_nth_access(i, kernel_dims, sr, global_size));
                }
                return region(std::move(boxes));
        }

        box<3> buffer_access_map::get_requirements_for_nth_access(const size_t n, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
                const auto& [_, rm] = m_accesses[n];

                chunk<3> chnk{sr.offset, sr.range, global_size};
                subrange<3> req;
                switch(kernel_dims) {
                case 0: req = apply_range_mapper<0>(rm.get(), chunk_cast<0>(chnk)); break;
                case 1: req = apply_range_mapper<1>(rm.get(), chunk_cast<1>(chnk)); break;
                case 2: req = apply_range_mapper<2>(rm.get(), chunk_cast<2>(chnk)); break;
                case 3: req = apply_range_mapper<3>(rm.get(), chunk_cast<3>(chnk)); break;
                default: assert(!"Unreachable");
                }
                return req;
        }

        void side_effect_map::add_side_effect(const host_object_id hoid, const experimental::side_effect_order order) {
                // TODO for multiple side effects on the same hoid, find the weakest order satisfying all of them
                emplace(hoid, order);
        }

        std::string print_task_debug_label(const task& tsk, bool title_case) {
                const auto type_string = [&] {
                        switch(tsk.get_type()) {
                        case task_type::epoch: return "epoch";
                        case task_type::host_compute: return "host-compute task";
                        case task_type::device_compute: return "device kernel";
                        case task_type::collective: return "collective host task";
                        case task_type::master_node: return "master-node host task";
                        case task_type::horizon: return "horizon";
                        case task_type::fence: return "fence";
                        default: return "unknown task";
                        }
                }();

                auto label = fmt::format("{} T{}", type_string, tsk.get_id());
                if(title_case) { label[0] = static_cast<char>(std::toupper(label[0])); }
                if(!tsk.get_debug_name().empty()) { fmt::format_to(std::back_inserter(label), " \"{}\"", tsk.get_debug_name()); }
                return label;
        }

        std::unordered_map<buffer_id, region<3>> detect_overlapping_writes(const task& tsk, const box_vector<3>& chunks) {
                const box<3> scalar_reduction_box({0, 0, 0}, {1, 1, 1});

                auto& bam = tsk.get_buffer_access_map();

                // track the union of writes we have checked so far in order to detect an overlap between that union and the next write
                std::unordered_map<buffer_id, region<3>> buffer_write_accumulators;
                // collect overlapping writes in order to report all of them before throwing
                std::unordered_map<buffer_id, region<3>> overlapping_writes;

                for(const auto bid : bam.get_accessed_buffers()) {
                        for(const auto& ck : chunks) {
                                region<3> writes;
                                for(const auto mode : bam.get_access_modes(bid)) {
                                        if(access::mode_traits::is_producer(mode)) {
                                                const auto req = bam.get_mode_requirements(bid, mode, tsk.get_dimensions(), ck.get_subrange(), tsk.get_global_size());
                                                writes = region_union(writes, req);
                                        }
                                }
                                if(!writes.empty()) {
                                        auto& write_accumulator = buffer_write_accumulators[bid]; // allow default-insert
                                        if(const auto overlap = region_intersection(write_accumulator, writes); !overlap.empty()) {
                                                auto& full_overlap = overlapping_writes[bid]; // allow default-insert
                                                full_overlap = region_union(full_overlap, overlap);
                                        }
                                        write_accumulator = region_union(write_accumulator, writes);
                                }
                        }
                }

                // we already check for accessor-reduction overlaps on task generation, but we still repeat the sanity-check here
                for(const auto& rinfo : tsk.get_reductions()) {
                        auto& write_accumulator = buffer_write_accumulators[rinfo.bid]; // allow default-insert
                        if(const auto overlap = region_intersection(write_accumulator, scalar_reduction_box); !overlap.empty()) {
                                auto& full_overlap = overlapping_writes[rinfo.bid]; // allow default-insert
                                full_overlap = region_union(full_overlap, overlap);
                        }
                        write_accumulator = region_union(write_accumulator, scalar_reduction_box);
                }

                return overlapping_writes;
        }

} // namespace detail
} // namespace celerity

1	#include "task.h"
2
3	#include "access_modes.h"
4	#include "utils.h"
5
6
7	namespace celerity {
8	namespace detail {
9
10	std::unordered_set<buffer_id> buffer_access_map::get_accessed_buffers() const {	26,992✔
11	std::unordered_set<buffer_id> result;	26,992✔
12	for(const auto& [bid, _] : m_accesses) {	49,505✔
13	result.emplace(bid);	22,528✔
14	}
15	return result;	26,969✔
16	}	×
17
18	std::unordered_set<cl::sycl::access::mode> buffer_access_map::get_access_modes(buffer_id bid) const {	18,594✔
19	std::unordered_set<cl::sycl::access::mode> result;	18,594✔
20	for(const auto& [b, rm] : m_accesses) {	44,408✔
21	if(b == bid) { result.insert(rm->get_access_mode()); }	25,812✔
22	}
23	return result;	18,593✔
24	}	×
25
26	template <int KernelDims>
27	subrange<3> apply_range_mapper(const range_mapper_base* rm, const chunk<KernelDims>& chnk) {	24,831✔
28	switch(rm->get_buffer_dimensions()) {	24,831!
29	case 0: return subrange_cast<3>(subrange<0>());	156✔
30	case 1: return subrange_cast<3>(rm->map_1(chnk));	31,699✔
31	case 2: return subrange_cast<3>(rm->map_2(chnk));	17,320✔
32	case 3: return rm->map_3(chnk);	233✔
UNCOV 33	default: assert(false);	×
34	}
35	return subrange<3>{};
36	}
37
38	region<3> buffer_access_map::get_mode_requirements(	47,136✔
39	const buffer_id bid, const access_mode mode, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
40	box_vector<3> boxes;	47,136✔
41	for(size_t i = 0; i < m_accesses.size(); ++i) {	107,530✔
42	if(m_accesses[i].first != bid \|\| m_accesses[i].second->get_access_mode() != mode) continue;	60,410✔
43	boxes.push_back(get_requirements_for_nth_access(i, kernel_dims, sr, global_size));	21,896✔
44	}
45	return region(std::move(boxes));	94,172✔
46	}	47,098✔
47
48	box<3> buffer_access_map::get_requirements_for_nth_access(const size_t n, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {	24,837✔
49	const auto& [_, rm] = m_accesses[n];	24,837✔
50
51	chunk<3> chnk{sr.offset, sr.range, global_size};	24,834✔
52	subrange<3> req;	24,834✔
53	switch(kernel_dims) {	24,835!
54	case 0: req = apply_range_mapper<0>(rm.get(), chunk_cast<0>(chnk)); break;	13,879✔
55	case 1: req = apply_range_mapper<1>(rm.get(), chunk_cast<1>(chnk)); break;	5,817✔
56	case 2: req = apply_range_mapper<2>(rm.get(), chunk_cast<2>(chnk)); break;	5,005✔
57	case 3: req = apply_range_mapper<3>(rm.get(), chunk_cast<3>(chnk)); break;	134✔
58	default: assert(!"Unreachable");	×
59	}
60	return req;	49,599✔
61	}
62
63	void side_effect_map::add_side_effect(const host_object_id hoid, const experimental::side_effect_order order) {	189✔
64	// TODO for multiple side effects on the same hoid, find the weakest order satisfying all of them
65	emplace(hoid, order);	189✔
66	}	189✔
67
68	std::string print_task_debug_label(const task& tsk, bool title_case) {	21✔
69	const auto type_string = [&] {	21✔
70	switch(tsk.get_type()) {	21!
71	case task_type::epoch: return "epoch";	×
72	case task_type::host_compute: return "host-compute task";	1✔
73	case task_type::device_compute: return "device kernel";	13✔
74	case task_type::collective: return "collective host task";	6✔
75	case task_type::master_node: return "master-node host task";	1✔
76	case task_type::horizon: return "horizon";	×
77	case task_type::fence: return "fence";	×
78	default: return "unknown task";	×
79	}
80	}();	21✔
81
82	auto label = fmt::format("{} T{}", type_string, tsk.get_id());	42✔
83	if(title_case) { label[0] = static_cast<char>(std::toupper(label[0])); }	21✔
84	if(!tsk.get_debug_name().empty()) { fmt::format_to(std::back_inserter(label), " \"{}\"", tsk.get_debug_name()); }	21✔
85	return label;	42✔
86	}	×
87
88	std::unordered_map<buffer_id, region<3>> detect_overlapping_writes(const task& tsk, const box_vector<3>& chunks) {	4,582✔
89	const box<3> scalar_reduction_box({0, 0, 0}, {1, 1, 1});	4,582✔
90
91	auto& bam = tsk.get_buffer_access_map();	4,582✔
92
93	// track the union of writes we have checked so far in order to detect an overlap between that union and the next write
94	std::unordered_map<buffer_id, region<3>> buffer_write_accumulators;	4,582✔
95	// collect overlapping writes in order to report all of them before throwing
96	std::unordered_map<buffer_id, region<3>> overlapping_writes;	4,582✔
97
98	for(const auto bid : bam.get_accessed_buffers()) {	8,256✔
99	for(const auto& ck : chunks) {	8,173✔
100	region<3> writes;	4,499✔
101	for(const auto mode : bam.get_access_modes(bid)) {	9,319✔
102	if(access::mode_traits::is_producer(mode)) {	4,820✔
103	const auto req = bam.get_mode_requirements(bid, mode, tsk.get_dimensions(), ck.get_subrange(), tsk.get_global_size());	3,549✔
104	writes = region_union(writes, req);	3,549✔
105	}	3,549✔
106	}	4,499✔
107	if(!writes.empty()) {	4,499✔
108	auto& write_accumulator = buffer_write_accumulators[bid]; // allow default-insert	3,520✔
109	if(const auto overlap = region_intersection(write_accumulator, writes); !overlap.empty()) {	3,520✔
110	auto& full_overlap = overlapping_writes[bid]; // allow default-insert	18✔
111	full_overlap = region_union(full_overlap, overlap);	18✔
112	}	3,520✔
113	write_accumulator = region_union(write_accumulator, writes);	3,520✔
114	}
115	}	4,499✔
116	}	4,582✔
117
118	// we already check for accessor-reduction overlaps on task generation, but we still repeat the sanity-check here
119	for(const auto& rinfo : tsk.get_reductions()) {	4,684✔
120	auto& write_accumulator = buffer_write_accumulators[rinfo.bid]; // allow default-insert	102✔
121	if(const auto overlap = region_intersection(write_accumulator, scalar_reduction_box); !overlap.empty()) {	102!
122	auto& full_overlap = overlapping_writes[rinfo.bid]; // allow default-insert	×
123	full_overlap = region_union(full_overlap, overlap);	×
124	}	102✔
125	write_accumulator = region_union(write_accumulator, scalar_reduction_box);	102✔
126	}
127
128	return overlapping_writes;	4,582✔
129	}	4,582✔
130
131	} // namespace detail
132	} // namespace celerity

celerity / celerity-runtime / 7783633344

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