• Home
  • Features
  • Pricing
  • Docs
  • Announcements
  • Sign In

celerity / celerity-runtime / 7783633344

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

Pull #230

github

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

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
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
STATUS · Troubleshooting · Open an Issue · Sales · Support · CAREERS · ENTERPRISE · START FREE · SCHEDULE DEMO
ANNOUNCEMENTS · TWITTER · TOS & SLA · Supported CI Services · What's a CI service? · Automated Testing

© 2026 Coveralls, Inc