Skip to content

Commit d8c3332

Browse files
GagaLPGagaLP
committed
Add out-of-bounds check for host side accessors
1 parent e7921a7 commit d8c3332

File tree

4 files changed

+162
-19
lines changed

4 files changed

+162
-19
lines changed

‎include/accessor.h

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -406,6 +406,22 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
406406

407407
template <access_mode M = Mode>
408408
inline std::enable_if_t<detail::access::mode_traits::is_producer(M), DataT&> operator[](const id<Dims>& index) const {
409+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
410+
if(m_oob_indices != nullptr) {
411+
const bool is_within_bounds_lo = all_true(index >= m_accessed_virtual_subrange.offset);
412+
const bool is_within_bounds_hi = all_true(index < (m_accessed_virtual_subrange.offset + m_accessed_virtual_subrange.range));
413+
414+
if((!is_within_bounds_lo || !is_within_bounds_hi)) {
415+
std::lock_guard<std::mutex> guard(m_oob_mutex);
416+
for(int d = 0; d < Dims; ++d) {
417+
m_oob_indices[0][d] = std::min(m_oob_indices[0][d], index[d]);
418+
m_oob_indices[1][d] = std::max(m_oob_indices[1][d], index[d] + 1);
419+
}
420+
return m_oob_fallback_value;
421+
}
422+
}
423+
#endif
424+
409425
return m_host_ptr[get_linear_offset(index)];
410426
}
411427

@@ -521,6 +537,16 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
521537
// m_host_ptr must be defined *last* for it to overlap with the sequence of range and id members in the 0-dimensional case
522538
CELERITY_DETAIL_NO_UNIQUE_ADDRESS DataT* m_host_ptr = nullptr;
523539

540+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
541+
id<3>* m_oob_indices = nullptr;
542+
// This mutex has to be inline static, since accessors are copyable making the mutex otherwise useless.
543+
// It is a workaround until atomic_ref() can be used on m_oob_indices in c++20.
544+
inline static std::mutex m_oob_mutex;
545+
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+
524550
template <target Target = target::host_task, typename Functor>
525551
accessor(ctor_internal_tag /* tag */, const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn) : m_virtual_buffer_range(buff.get_range()) {
526552
using range_mapper = detail::range_mapper<Dims, std::decay_t<Functor>>; // decay function type to function pointer
@@ -547,6 +573,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
547573
m_backing_buffer_range = other.m_backing_buffer_range;
548574
m_virtual_buffer_range = other.m_virtual_buffer_range;
549575

576+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
577+
m_oob_indices = other.m_oob_indices;
578+
#endif
579+
550580
if(detail::is_embedded_hydration_id(m_host_ptr)) {
551581
if(detail::cgf_diagnostics::is_available() && detail::cgf_diagnostics::get_instance().is_checking()) {
552582
detail::cgf_diagnostics::get_instance().register_accessor(detail::extract_hydration_id(m_host_ptr), target::host_task);
@@ -558,6 +588,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
558588
m_backing_buffer_offset = detail::id_cast<Dims>(info.backing_buffer_offset);
559589
m_backing_buffer_range = detail::range_cast<Dims>(info.backing_buffer_range);
560590
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
561595
}
562596
}
563597
}

‎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: 47 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -160,7 +160,21 @@ namespace detail {
160160
const auto [bid, mode] = access_map.get_nth_access(i);
161161
const auto sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
162162
const auto info = m_buffer_mngr.access_host_buffer(bid, mode, sr);
163+
164+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
165+
// oob_indices[0] contains the lower bound oob indices
166+
// oob_indices[1] contains the upper bound oob indices
167+
std::vector<id<3>> oob_indices{2};
168+
constexpr size_t size_t_max = std::numeric_limits<size_t>::max();
169+
const auto buffer_dims = m_buffer_mngr.get_buffer_info(bid).dimensions;
170+
oob_indices[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
171+
oob_indices[1] = id<3>{0, 0, 0};
172+
access_infos.push_back(
173+
closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_indices.data()});
174+
m_oob_indices_per_accessor.push_back(std::move(oob_indices));
175+
#else
163176
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
177+
#endif
164178
}
165179

166180
closure_hydrator::get_instance().arm(target::host_task, std::move(access_infos));
@@ -175,6 +189,27 @@ namespace detail {
175189
if(m_future.wait_for(std::chrono::seconds(0)) == std::future_status::ready) {
176190
m_buffer_mngr.unlock(pkg.cid);
177191

192+
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
193+
const auto data = std::get<execution_data>(pkg.data);
194+
auto tsk = m_task_mngr.get_task(data.tid);
195+
196+
for(size_t i = 0; i < m_oob_indices_per_accessor.size(); ++i) {
197+
const id<3>& oob_min = m_oob_indices_per_accessor[i][0];
198+
const id<3>& oob_max = m_oob_indices_per_accessor[i][1];
199+
200+
if(oob_max != id<3>{0, 0, 0}) {
201+
const auto& access_map = tsk->get_buffer_access_map();
202+
const auto acc_sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
203+
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
204+
const auto buffer_id = access_map.get_nth_access(i).first;
205+
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
206+
CELERITY_ERROR("Out-of-bounds access in host task detected: Accessor {} for buffer {} attempted to access indices between {} which are "
207+
"outside of mapped subrange {}",
208+
i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
209+
}
210+
}
211+
#endif
212+
178213
auto info = m_future.get();
179214
CELERITY_TRACE("Delta time submit -> start: {}us, start -> end: {}us",
180215
std::chrono::duration_cast<std::chrono::microseconds>(info.start_time - info.submit_time).count(),
@@ -217,14 +252,16 @@ namespace detail {
217252
try {
218253
const auto info = m_buffer_mngr.access_device_buffer(bid, mode, sr);
219254
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
220-
auto* const oob_idx = sycl::malloc_host<id<3>>(2, m_queue.get_sycl_queue());
221-
assert(oob_idx != nullptr);
255+
// oob_indices[0] contains the lower bound oob indices
256+
// oob_indices[1] contains the upper bound oob indices
257+
auto* const oob_indices = sycl::malloc_host<id<3>>(2, m_queue.get_sycl_queue());
258+
assert(oob_indices != nullptr);
222259
constexpr size_t size_t_max = std::numeric_limits<size_t>::max();
223260
const auto buffer_dims = m_buffer_mngr.get_buffer_info(bid).dimensions;
224-
oob_idx[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
225-
oob_idx[1] = id<3>{1, 1, 1};
226-
m_oob_indices_per_accessor.push_back(oob_idx);
227-
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_idx});
261+
oob_indices[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
262+
oob_indices[1] = id<3>{0, 0, 0};
263+
m_oob_indices_per_accessor.push_back(oob_indices);
264+
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_indices});
228265
#else
229266
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
230267
#endif
@@ -259,13 +296,15 @@ namespace detail {
259296
const id<3>& oob_min = m_oob_indices_per_accessor[i][0];
260297
const id<3>& oob_max = m_oob_indices_per_accessor[i][1];
261298

262-
if(oob_max != id<3>{1, 1, 1}) {
299+
if(oob_max != id<3>{0, 0, 0}) {
263300
const auto& access_map = tsk->get_buffer_access_map();
264301
const auto acc_sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
265302
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
303+
const auto buffer_id = access_map.get_nth_access(i).first;
304+
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
266305
CELERITY_ERROR("Out-of-bounds access in kernel '{}' detected: Accessor {} for buffer {} attempted to access indices between {} which are "
267306
"outside of mapped subrange {}",
268-
tsk->get_debug_name(), i, access_map.get_nth_access(i).first, oob_sr, acc_sr);
307+
tsk->get_debug_name(), i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
269308
}
270309
sycl::free(m_oob_indices_per_accessor[i], m_queue.get_sycl_queue());
271310
}

‎test/accessor_tests.cc

Lines changed: 77 additions & 11 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(test_utils::truncate_range<Dims>({10, 20, 30}));
657+
buffer<int, Dims> unnamed_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
658+
buffer<int, Dims> named_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
658659
const auto accessible_sr = test_utils::truncate_subrange<Dims>({{5, 10, 15}, {1, 2, 3}});
659660
const auto oob_idx_lo = test_utils::truncate_id<Dims>({1, 2, 3});
660661
const auto oob_idx_hi = test_utils::truncate_id<Dims>({7, 13, 25});
662+
const auto buffer_name = "oob";
661663

662-
// we need to be careful about the orderign of the construction and destruction
664+
celerity::debug::set_buffer_name(named_buff, buffer_name);
665+
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,83 @@ 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 unnamed_acc(unnamed_buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
676+
accessor named_acc(named_buff, 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>(ones), [=](item<Dims>) {
673-
acc[oob_idx_lo] = 0;
674-
acc[oob_idx_hi] = 0;
679+
unnamed_acc[oob_idx_lo] = 0;
680+
unnamed_acc[oob_idx_hi] = 0;
681+
682+
named_acc[oob_idx_lo] = 0;
683+
named_acc[oob_idx_hi] = 0;
675684
});
676685
});
677686
q.slow_full_sync();
678687
}
679688

680-
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>(ones)))};
681-
const auto error_message = fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for "
682-
"buffer 0 attempted to 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));
689+
const auto attempted_sr =
690+
subrange<3>{id_cast<3>(oob_idx_lo), range_cast<3>(oob_idx_hi - oob_idx_lo + id_cast<Dims>(range<Dims>(ones))) - range_cast<3>(range<Dims>(zeros))};
691+
const auto unnamed_error_message =
692+
fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for buffer 0 attempted to "
693+
"access indices between {} which are outside of mapped subrange {}",
694+
Dims, attempted_sr, subrange_cast<3>(accessible_sr));
695+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(unnamed_error_message));
696+
697+
const auto named_error_message =
698+
fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 1 for buffer {} attempted to "
699+
"access indices between {} which are outside of mapped subrange {}",
700+
Dims, buffer_name, attempted_sr, subrange_cast<3>(accessible_sr));
701+
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(named_error_message));
685702
}
686703

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

0 commit comments

Comments
 (0)