Skip to content

Commit d7fa0a6

Browse files
GagaLPGagaLP
committed
Add out-of-bounds check for host side accessors
1 parent 0a282f5 commit d7fa0a6

File tree

4 files changed

+142
-13
lines changed

4 files changed

+142
-13
lines changed

‎include/accessor.h

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -407,6 +407,25 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
407407

408408
template <access_mode M = Mode>
409409
inline std::enable_if_t<detail::access::mode_traits::is_producer(M), DataT&> operator[](const id<Dims>& index) const {
410+
411+
// analogue to the device buffer.
412+
// without sycl::memory ...
413+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
414+
if(m_oob_indices != nullptr) {
415+
const id<Dims> all_true = detail::id_cast<Dims>(id<3>(true, true, true));
416+
const bool is_within_bounds_lo = (index >= m_accessed_virtual_subrange.offset) == all_true;
417+
const bool is_within_bounds_hi = (index < (m_accessed_virtual_subrange.offset + m_accessed_virtual_subrange.range)) == all_true;
418+
419+
if((!is_within_bounds_lo || !is_within_bounds_hi)) {
420+
for(int d = 0; d < Dims; ++d) {
421+
m_oob_indices[0][d] = std::min(m_oob_indices[0][d], index[d]);
422+
m_oob_indices[1][d] = std::max(m_oob_indices[1][d], index[d] + 1);
423+
}
424+
return m_oob_fallback_value;
425+
}
426+
}
427+
#endif
428+
410429
return m_host_ptr[get_linear_offset(index)];
411430
}
412431

@@ -522,6 +541,12 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
522541
// m_host_ptr must be defined *last* for it to overlap with the sequence of range and id members in the 0-dimensional case
523542
CELERITY_DETAIL_NO_UNIQUE_ADDRESS DataT* m_host_ptr = nullptr;
524543

544+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
545+
id<3>* m_oob_indices = nullptr;
546+
// This value (or a reference to it) is returned for all out-of-bounds accesses.
547+
mutable DataT m_oob_fallback_value = DataT{};
548+
#endif
549+
525550
template <target Target = target::host_task, typename Functor>
526551
accessor(ctor_internal_tag /* tag */, const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn) : m_virtual_buffer_range(buff.get_range()) {
527552
using range_mapper = detail::range_mapper<Dims, std::decay_t<Functor>>; // decay function type to function pointer
@@ -548,6 +573,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
548573
m_backing_buffer_range = other.m_backing_buffer_range;
549574
m_virtual_buffer_range = other.m_virtual_buffer_range;
550575

576+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
577+
m_oob_indices = other.m_oob_indices;
578+
#endif
579+
551580
if(detail::is_embedded_hydration_id(m_host_ptr)) {
552581
if(detail::cgf_diagnostics::is_available() && detail::cgf_diagnostics::get_instance().is_checking()) {
553582
detail::cgf_diagnostics::get_instance().register_accessor(detail::extract_hydration_id(m_host_ptr), target::host_task);
@@ -559,6 +588,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
559588
m_backing_buffer_offset = detail::id_cast<Dims>(info.backing_buffer_offset);
560589
m_backing_buffer_range = detail::range_cast<Dims>(info.backing_buffer_range);
561590
m_accessed_virtual_subrange = detail::subrange_cast<Dims>(info.accessed_virtual_subrange);
591+
592+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
593+
m_oob_indices = info.out_of_bounds_indices;
594+
#endif
562595
}
563596
}
564597
}

‎include/worker_job.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -156,6 +156,10 @@ namespace detail {
156156
std::future<host_queue::execution_info> m_future;
157157
bool m_submitted = false;
158158

159+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
160+
std::vector<std::vector<id<3>>> m_oob_indices_per_accessor;
161+
#endif
162+
159163
bool execute(const command_pkg& pkg) override;
160164
std::string get_description(const command_pkg& pkg) override;
161165
};

‎src/worker_job.cc

Lines changed: 37 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -160,7 +160,19 @@ namespace detail {
160160
const auto [bid, mode] = access_map.get_nth_access(i);
161161
const auto sr = grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
162162
const auto info = m_buffer_mngr.access_host_buffer(bid, mode, sr);
163-
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
163+
164+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
165+
// for host tasks this can be a vector (RAII)
166+
std::vector<id<3>> oob_idx {2};
167+
constexpr size_t size_t_max = std::numeric_limits<size_t>::max();
168+
const auto buffer_dims = m_buffer_mngr.get_buffer_info(bid).dimensions;
169+
oob_idx[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
170+
oob_idx[1] = id<3>{1, 1, 1};
171+
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_idx.data()});
172+
m_oob_indices_per_accessor.push_back(std::move(oob_idx));
173+
#else
174+
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
175+
#endif
164176
}
165177

166178
closure_hydrator::get_instance().arm(target::host_task, std::move(access_infos));
@@ -175,6 +187,27 @@ namespace detail {
175187
if(m_future.wait_for(std::chrono::seconds(0)) == std::future_status::ready) {
176188
m_buffer_mngr.unlock(pkg.cid);
177189

190+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
191+
const auto data = std::get<execution_data>(pkg.data);
192+
auto tsk = m_task_mngr.get_task(data.tid);
193+
194+
for(size_t i = 0; i < m_oob_indices_per_accessor.size(); ++i) {
195+
const id<3>& oob_min = m_oob_indices_per_accessor[i][0];
196+
const id<3>& oob_max = m_oob_indices_per_accessor[i][1];
197+
198+
if(oob_max != id<3>{1, 1, 1}) {
199+
const auto& access_map = tsk->get_buffer_access_map();
200+
const auto acc_sr =
201+
grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
202+
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
203+
const auto buffer_id = access_map.get_nth_access(i).first;
204+
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
205+
CELERITY_ERROR("Out-of-bounds access in host kernel detected: Accessor {} for buffer {} attempted to access indices between {} which are "
206+
"outside of mapped subrange {}", i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
207+
}
208+
}
209+
#endif
210+
178211
auto info = m_future.get();
179212
CELERITY_TRACE("Delta time submit -> start: {}us, start -> end: {}us",
180213
std::chrono::duration_cast<std::chrono::microseconds>(info.start_time - info.submit_time).count(),
@@ -264,9 +297,10 @@ namespace detail {
264297
const auto acc_sr =
265298
grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
266299
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
300+
const auto buffer_id = access_map.get_nth_access(i).first;
301+
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
267302
CELERITY_ERROR("Out-of-bounds access in kernel '{}' detected: Accessor {} for buffer {} attempted to access indices between {} which are "
268-
"outside of mapped subrange {}",
269-
tsk->get_debug_name(), i, access_map.get_nth_access(i).first, oob_sr, acc_sr);
303+
"outside of mapped subrange {}", tsk->get_debug_name(), i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
270304
}
271305
sycl::free(m_oob_indices_per_accessor[i], m_queue.get_sycl_queue());
272306
}

‎test/accessor_tests.cc

Lines changed: 68 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -650,16 +650,20 @@ namespace detail {
650650
template <int>
651651
class acc_out_of_bounds_kernel {};
652652

653-
TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
653+
TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "device accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
654654
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
655655
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
656656
#endif
657-
buffer<int, Dims> buff(range_cast<Dims>(range<3>{10, 20, 30}));
657+
buffer<int, Dims> buff_nl(range_cast<Dims>(range<3>{10, 20, 30}));
658+
buffer<int, Dims> buff_nf(range_cast<Dims>(range<3>{10, 20, 30}));
658659
const auto accessible_sr = subrange_cast<Dims>(subrange<3>{{5, 10, 15}, {1, 2, 3}});
659660
const auto oob_idx_lo = id_cast<Dims>(id<3>{1, 2, 3});
660661
const auto oob_idx_hi = id_cast<Dims>(id<3>{7, 13, 25});
662+
const auto buffer_name = "oob";
663+
664+
celerity::debug::set_buffer_name(buff_nf, buffer_name);
661665

662-
// we need to be careful about the orderign of the construction and destruction
666+
// we need to be careful about the ordering of the construction and destruction
663667
// of the Celerity queue and the log capturing utility here
664668
std::unique_ptr<celerity::test_utils::log_capture> lc;
665669
{
@@ -668,21 +672,75 @@ namespace detail {
668672
lc = std::make_unique<celerity::test_utils::log_capture>();
669673

670674
q.submit([&](handler& cgh) {
671-
accessor acc(buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
675+
accessor acc_nl(buff_nl, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
676+
accessor acc_nf(buff_nf, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
677+
672678
cgh.parallel_for<acc_out_of_bounds_kernel<Dims>>(range<Dims>(unit_range), [=](item<Dims>) {
673-
acc[oob_idx_lo] = 0;
674-
acc[oob_idx_hi] = 0;
679+
acc_nl[oob_idx_lo] = 0;
680+
acc_nl[oob_idx_hi] = 0;
681+
682+
acc_nf[oob_idx_lo] = 0;
683+
acc_nf[oob_idx_hi] = 0;
675684
});
676685
});
677686
q.slow_full_sync();
678687
}
679688

680689
const auto attempted_sr = subrange<3>{id_cast<3>(oob_idx_lo), range_cast<3>(oob_idx_hi - oob_idx_lo + id_cast<Dims>(range<Dims>(unit_range)))};
681-
const auto error_message = fmt::format("Out-of-bounds access in kernel 'acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for buffer 0 attempted to "
682-
"access indices between {} which are outside of mapped subrange {}",
683-
Dims, attempted_sr, subrange_cast<3>(accessible_sr));
684-
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message));
690+
const auto error_message_nl = fmt::format("Out-of-bounds access in kernel 'acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for buffer 0 attempted to "
691+
"access indices between {} which are outside of mapped subrange {}", Dims, attempted_sr, subrange_cast<3>(accessible_sr));
692+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message_nl));
693+
694+
const auto error_message_nf = fmt::format("Out-of-bounds access in kernel 'acc_out_of_bounds_kernel<{}>' detected: Accessor 1 for buffer {} attempted to "
695+
"access indices between {} which are outside of mapped subrange {}", Dims, buffer_name, attempted_sr, subrange_cast<3>(accessible_sr));
696+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message_nf));
685697
}
686698

699+
TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "host accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
700+
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
701+
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
702+
#endif
703+
buffer<int, Dims> buff_nl(range_cast<Dims>(range<3>{10, 20, 30}));
704+
buffer<int, Dims> buff_nf(range_cast<Dims>(range<3>{10, 20, 30}));
705+
const auto accessible_sr = subrange_cast<Dims>(subrange<3>{{5, 10, 15}, {1, 2, 3}});
706+
const auto oob_idx_lo = id_cast<Dims>(id<3>{1, 2, 3});
707+
const auto oob_idx_hi = id_cast<Dims>(id<3>{7, 13, 25});
708+
const auto buffer_name = "oob";
709+
710+
celerity::debug::set_buffer_name(buff_nf, buffer_name);
711+
712+
// we need to be careful about the ordering of the construction and destruction
713+
// of the Celerity queue and the log capturing utility here
714+
std::unique_ptr<celerity::test_utils::log_capture> lc;
715+
{
716+
distr_queue q;
717+
718+
lc = std::make_unique<celerity::test_utils::log_capture>();
719+
720+
q.submit([&](handler& cgh){
721+
accessor acc_nl(buff_nl, cgh, celerity::access::fixed(accessible_sr), celerity::write_only_host_task, celerity::no_init);
722+
accessor acc_nf(buff_nf, cgh, celerity::access::fixed(accessible_sr), celerity::write_only_host_task, celerity::no_init);
723+
724+
cgh.host_task(range<Dims>(unit_range), [=](partition<Dims>) {
725+
acc_nl[oob_idx_lo] = 0;
726+
acc_nl[oob_idx_hi] = 0;
727+
728+
acc_nf[oob_idx_lo] = 0;
729+
acc_nf[oob_idx_hi] = 0;
730+
});
731+
});
732+
733+
q.slow_full_sync();
734+
}
735+
736+
const auto attempted_sr = subrange<3>{id_cast<3>(oob_idx_lo), range_cast<3>(oob_idx_hi - oob_idx_lo + id_cast<Dims>(range<Dims>(unit_range)))};
737+
const auto error_message_nl = fmt::format("Out-of-bounds access in host kernel detected: Accessor 0 for buffer 0 attempted to "
738+
"access indices between {} which are outside of mapped subrange {}", attempted_sr, subrange_cast<3>(accessible_sr));
739+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message_nl));
740+
741+
const auto error_message_nf = fmt::format("Out-of-bounds access in host kernel detected: Accessor 1 for buffer {} attempted to "
742+
"access indices between {} which are outside of mapped subrange {}", buffer_name, attempted_sr, subrange_cast<3>(accessible_sr));
743+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message_nf));
744+
}
687745
} // namespace detail
688746
} // namespace celerity

0 commit comments

Comments
 (0)