11665018161

Committed 04 Nov 2024 01:17PM UTC coverage: 95.167% (-0.08%) from 95.251%

Build # 11665018161

Build Type

push

github

Committed by

fknorr

Commit Message

Update benchmark results for region_builder

Run Details

3016 of 3412 branches covered (88.39%)

Branch coverage included in aggregate %.

6691 of 6788 relevant lines covered (98.57%)

1312478.22 hits per line

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

93.6

/src/task.cc

#include "task.h"

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

#include <utility>


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<sycl::access::mode> buffer_access_map::get_access_modes(buffer_id bid) const {
                std::unordered_set<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>
        region<3> apply_range_mapper(const range_mapper_base* rm, const chunk<KernelDims>& chnk) {
                switch(rm->get_buffer_dimensions()) {
                case 0: return region_cast<3>(region(box<0>()));
                case 1: return region_cast<3>(rm->map_1(chnk));
                case 2: return region_cast<3>(rm->map_2(chnk));
                case 3: return rm->map_3(chnk);
                default: utils::unreachable(); // LCOV_EXCL_LINE
                }
        }

        region<3> apply_range_mapper(const range_mapper_base* rm, const chunk<3>& chnk, int kernel_dims) {
                switch(kernel_dims) {
                case 0: return apply_range_mapper<0>(rm, chunk_cast<0>(chnk));
                case 1: return apply_range_mapper<1>(rm, chunk_cast<1>(chnk));
                case 2: return apply_range_mapper<2>(rm, chunk_cast<2>(chnk));
                case 3: return apply_range_mapper<3>(rm, chunk_cast<3>(chnk));
                default: utils::unreachable(); // LCOV_EXCL_LINE
                }
        }

        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 {
                region_builder<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.add(get_requirements_for_nth_access(i, kernel_dims, sr, global_size));
                }
                return std::move(boxes).into_region();
        }

        region<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 {
                return apply_range_mapper(m_accesses[n].second.get(), chunk<3>{sr.offset, sr.range, global_size}, kernel_dims);
        }

        box_vector<3> buffer_access_map::get_required_contiguous_boxes(
            const buffer_id bid, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
                box_vector<3> boxes;
                for(const auto& [a_bid, a_rm] : m_accesses) {
                        if(a_bid == bid) {
                                const auto accessed_region = apply_range_mapper(a_rm.get(), chunk<3>{sr.offset, sr.range, global_size}, kernel_dims);
                                if(!accessed_region.empty()) { boxes.push_back(bounding_box(accessed_region)); }
                        }
                }
                return boxes;
        }

        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) {
                return utils::make_task_debug_label(tsk.get_type(), tsk.get_id(), tsk.get_debug_name(), title_case);
        }

        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 "grid.h"
5	#include "ranges.h"
6	#include "utils.h"
7
8	#include <utility>
9
10
11	namespace celerity {
12	namespace detail {
13
14	std::unordered_set<buffer_id> buffer_access_map::get_accessed_buffers() const {	27,993✔
15	std::unordered_set<buffer_id> result;	27,993✔
16	for(const auto& [bid, _] : m_accesses) {	56,966✔
17	result.emplace(bid);	28,973✔
18	}
19	return result;	27,993✔
20	}	×
21
22	std::unordered_set<sycl::access::mode> buffer_access_map::get_access_modes(buffer_id bid) const {	26,603✔
23	std::unordered_set<sycl::access::mode> result;	26,603✔
24	for(const auto& [b, rm] : m_accesses) {	65,084✔
25	if(b == bid) { result.insert(rm->get_access_mode()); }	38,481✔
26	}
27	return result;	26,603✔
28	}	×
29
30	template <int KernelDims>
31	region<3> apply_range_mapper(const range_mapper_base* rm, const chunk<KernelDims>& chnk) {	35,871✔
32	switch(rm->get_buffer_dimensions()) {	35,871!
33	case 0: return region_cast<3>(region(box<0>()));	22,908✔
34	case 1: return region_cast<3>(rm->map_1(chnk));	7,315✔
35	case 2: return region_cast<3>(rm->map_2(chnk));	16,019✔
36	case 3: return rm->map_3(chnk);	1,083✔
37	default: utils::unreachable(); // LCOV_EXCL_LINE
38	}
39	}
40
41	region<3> apply_range_mapper(const range_mapper_base* rm, const chunk<3>& chnk, int kernel_dims) {	35,871✔
42	switch(kernel_dims) {	35,871!
43	case 0: return apply_range_mapper<0>(rm, chunk_cast<0>(chnk));	28,132✔
44	case 1: return apply_range_mapper<1>(rm, chunk_cast<1>(chnk));	19,049✔
45	case 2: return apply_range_mapper<2>(rm, chunk_cast<2>(chnk));	23,106✔
46	case 3: return apply_range_mapper<3>(rm, chunk_cast<3>(chnk));	1,453✔
47	default: utils::unreachable(); // LCOV_EXCL_LINE
48	}
49	}
50
51	region<3> buffer_access_map::get_mode_requirements(	61,135✔
52	const buffer_id bid, const access_mode mode, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
53	region_builder<3> boxes;	61,135✔
54	for(size_t i = 0; i < m_accesses.size(); ++i) {	144,019✔
55	if(m_accesses[i].first != bid \|\| m_accesses[i].second->get_access_mode() != mode) continue;	82,887✔
56	boxes.add(get_requirements_for_nth_access(i, kernel_dims, sr, global_size));	29,754✔
57	}
58	return std::move(boxes).into_region();	122,255✔
59	}	61,135✔
60
61	region<3> buffer_access_map::get_requirements_for_nth_access(	32,783✔
62	const size_t n, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
63	return apply_range_mapper(m_accesses[n].second.get(), chunk<3>{sr.offset, sr.range, global_size}, kernel_dims);	65,564✔
64	}
65
66	box_vector<3> buffer_access_map::get_required_contiguous_boxes(	3,244✔
67	const buffer_id bid, const int kernel_dims, const subrange<3>& sr, const range<3>& global_size) const {
68	box_vector<3> boxes;	3,244✔
69	for(const auto& [a_bid, a_rm] : m_accesses) {	7,814✔
70	if(a_bid == bid) {	4,570✔
71	const auto accessed_region = apply_range_mapper(a_rm.get(), chunk<3>{sr.offset, sr.range, global_size}, kernel_dims);	3,088✔
72	if(!accessed_region.empty()) { boxes.push_back(bounding_box(accessed_region)); }	3,088✔
73	}	3,088✔
74	}
75	return boxes;	3,244✔
76	}	×
77
78	void side_effect_map::add_side_effect(const host_object_id hoid, const experimental::side_effect_order order) {	228✔
79	// TODO for multiple side effects on the same hoid, find the weakest order satisfying all of them
80	emplace(hoid, order);	228✔
81	}	228✔
82
83	std::string print_task_debug_label(const task& tsk, bool title_case) {	34✔
84	return utils::make_task_debug_label(tsk.get_type(), tsk.get_id(), tsk.get_debug_name(), title_case);	34✔
85	}
86
87	std::unordered_map<buffer_id, region<3>> detect_overlapping_writes(const task& tsk, const box_vector<3>& chunks) {	5,104✔
88	const box<3> scalar_reduction_box({0, 0, 0}, {1, 1, 1});	5,104✔
89
90	auto& bam = tsk.get_buffer_access_map();	5,104✔
91
92	// track the union of writes we have checked so far in order to detect an overlap between that union and the next write
93	std::unordered_map<buffer_id, region<3>> buffer_write_accumulators;	5,104✔
94	// collect overlapping writes in order to report all of them before throwing
95	std::unordered_map<buffer_id, region<3>> overlapping_writes;	5,104✔
96
97	for(const auto bid : bam.get_accessed_buffers()) {	10,497✔
98	for(const auto& ck : chunks) {	12,835✔
99	region<3> writes;	7,442✔
100	for(const auto mode : bam.get_access_modes(bid)) {	15,219✔
101	if(access::mode_traits::is_producer(mode)) {	7,777✔
102	const auto req = bam.get_mode_requirements(bid, mode, tsk.get_dimensions(), ck.get_subrange(), tsk.get_global_size());	4,958✔
103	writes = region_union(writes, req);	4,958✔
104	}	4,958✔
105	}	7,442✔
106	if(!writes.empty()) {	7,442✔
107	auto& write_accumulator = buffer_write_accumulators[bid]; // allow default-insert	4,905✔
108	if(const auto overlap = region_intersection(write_accumulator, writes); !overlap.empty()) {	4,905✔
109	auto& full_overlap = overlapping_writes[bid]; // allow default-insert	23✔
110	full_overlap = region_union(full_overlap, overlap);	23✔
111	}	4,905✔
112	write_accumulator = region_union(write_accumulator, writes);	4,905✔
113	}
114	}	7,442✔
115	}	5,104✔
116
117	// we already check for accessor-reduction overlaps on task generation, but we still repeat the sanity-check here
118	for(const auto& rinfo : tsk.get_reductions()) {	5,382✔
119	auto& write_accumulator = buffer_write_accumulators[rinfo.bid]; // allow default-insert	278✔
120	if(const auto overlap = region_intersection(write_accumulator, scalar_reduction_box); !overlap.empty()) {	278!
121	auto& full_overlap = overlapping_writes[rinfo.bid]; // allow default-insert	×
122	full_overlap = region_union(full_overlap, overlap);	×
123	}	278✔
124	write_accumulator = region_union(write_accumulator, scalar_reduction_box);	278✔
125	}
126
127	return overlapping_writes;	5,104✔
128	}	5,104✔
129
130	} // namespace detail
131	} // namespace celerity

celerity / celerity-runtime / 11665018161

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