Skip to content

Commit

Permalink
Add out-of-bounds check for host side accessors
Browse files Browse the repository at this point in the history
  • Loading branch information
@GagaLP
GagaLP committed Sep 14, 2023
1 parent f190da3 commit 3c3c479
Show file tree
Hide file tree
Showing 4 changed files with 150 additions and 19 deletions.
36 changes: 36 additions & 0 deletions include/accessor.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -407,6 +407,24 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b

template <access_mode M = Mode>
inline std::enable_if_t<detail::access::mode_traits::is_producer(M), DataT&> operator[](const id<Dims>& index) const {

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
if(m_oob_indices != nullptr) {
const id<Dims> all_true = detail::id_cast<Dims>(id<3>(true, true, true));
const bool is_within_bounds_lo = (index >= m_accessed_virtual_subrange.offset) == all_true;
const bool is_within_bounds_hi = (index < (m_accessed_virtual_subrange.offset + m_accessed_virtual_subrange.range)) == all_true;

if((!is_within_bounds_lo || !is_within_bounds_hi)) {
for(int d = 0; d < Dims; ++d) {
std::lock_guard<std::mutex> guard (m_oob_mutex);
m_oob_indices[0][d] = std::min(m_oob_indices[0][d], index[d]);
m_oob_indices[1][d] = std::max(m_oob_indices[1][d], index[d] + 1);
}
return m_oob_fallback_value;
}
}
#endif

return m_host_ptr[get_linear_offset(index)];
}

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

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
id<3>* m_oob_indices = nullptr;
// This mutex has to be inline static, since accessors are copyable making the mutex otherwise useless.
// It is a workaround until atomic_ref() can be used on m_oob_indices in c++20.
inline static std::mutex m_oob_mutex;

// This value (or a reference to it) is returned for all out-of-bounds accesses.
mutable DataT m_oob_fallback_value = DataT{};
#endif

template <target Target = target::host_task, typename Functor>
accessor(ctor_internal_tag /* tag */, const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn) : m_virtual_buffer_range(buff.get_range()) {
using range_mapper = detail::range_mapper<Dims, std::decay_t<Functor>>; // decay function type to function pointer
Expand All @@ -548,6 +576,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
m_backing_buffer_range = other.m_backing_buffer_range;
m_virtual_buffer_range = other.m_virtual_buffer_range;

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
m_oob_indices = other.m_oob_indices;
#endif

if(detail::is_embedded_hydration_id(m_host_ptr)) {
if(detail::cgf_diagnostics::is_available() && detail::cgf_diagnostics::get_instance().is_checking()) {
detail::cgf_diagnostics::get_instance().register_accessor(detail::extract_hydration_id(m_host_ptr), target::host_task);
Expand All @@ -559,6 +591,10 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
m_backing_buffer_offset = detail::id_cast<Dims>(info.backing_buffer_offset);
m_backing_buffer_range = detail::range_cast<Dims>(info.backing_buffer_range);
m_accessed_virtual_subrange = detail::subrange_cast<Dims>(info.accessed_virtual_subrange);

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
m_oob_indices = info.out_of_bounds_indices;
#endif
}
}
}
Expand Down
4 changes: 4 additions & 0 deletions include/worker_job.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -156,6 +156,10 @@ namespace detail {
std::future<host_queue::execution_info> m_future;
bool m_submitted = false;

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
std::vector<std::vector<id<3>>> m_oob_indices_per_accessor;
#endif

bool execute(const command_pkg& pkg) override;
std::string get_description(const command_pkg& pkg) override;
};
Expand Down
51 changes: 42 additions & 9 deletions src/worker_job.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -160,7 +160,18 @@ namespace detail {
const auto [bid, mode] = access_map.get_nth_access(i);
const auto sr = grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
const auto info = m_buffer_mngr.access_host_buffer(bid, mode, sr);
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
std::vector<id<3>> oob_indices {2};
constexpr size_t size_t_max = std::numeric_limits<size_t>::max();
const auto buffer_dims = m_buffer_mngr.get_buffer_info(bid).dimensions;
oob_indices[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
oob_indices[1] = id<3>{1, 1, 1};
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_indices.data()});
m_oob_indices_per_accessor.push_back(std::move(oob_indices));
#else
access_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
#endif
}

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

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
const auto data = std::get<execution_data>(pkg.data);
auto tsk = m_task_mngr.get_task(data.tid);

for(size_t i = 0; i < m_oob_indices_per_accessor.size(); ++i) {
const id<3>& oob_min = m_oob_indices_per_accessor[i][0];
const id<3>& oob_max = m_oob_indices_per_accessor[i][1];

if(oob_max != id<3>{1, 1, 1}) {
const auto& access_map = tsk->get_buffer_access_map();
const auto acc_sr =
grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
const auto buffer_id = access_map.get_nth_access(i).first;
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
CELERITY_ERROR("Out-of-bounds access in host task detected: Accessor {} for buffer {} attempted to access indices between {} which are "
"outside of mapped subrange {}", i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
}
}
#endif

auto info = m_future.get();
CELERITY_TRACE("Delta time submit -> start: {}us, start -> end: {}us",
std::chrono::duration_cast<std::chrono::microseconds>(info.start_time - info.submit_time).count(),
Expand Down Expand Up @@ -217,14 +249,14 @@ namespace detail {
try {
const auto info = m_buffer_mngr.access_device_buffer(bid, mode, sr);
#if CELERITY_ACCESSOR_BOUNDARY_CHECK
auto* const oob_idx = sycl::malloc_host<id<3>>(2, m_queue.get_sycl_queue());
assert(oob_idx != nullptr);
auto* const oob_indices = sycl::malloc_host<id<3>>(2, m_queue.get_sycl_queue());
assert(oob_indices != nullptr);
constexpr size_t size_t_max = std::numeric_limits<size_t>::max();
const auto buffer_dims = m_buffer_mngr.get_buffer_info(bid).dimensions;
oob_idx[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
oob_idx[1] = id<3>{1, 1, 1};
m_oob_indices_per_accessor.push_back(oob_idx);
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_idx});
oob_indices[0] = id<3>{size_t_max, buffer_dims > 1 ? size_t_max : 0, buffer_dims == 3 ? size_t_max : 0};
oob_indices[1] = id<3>{1, 1, 1};
m_oob_indices_per_accessor.push_back(oob_indices);
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr, oob_indices});
#else
accessor_infos.push_back(closure_hydrator::accessor_info{info.ptr, info.backing_buffer_range, info.backing_buffer_offset, sr});
#endif
Expand Down Expand Up @@ -264,9 +296,10 @@ namespace detail {
const auto acc_sr =
grid_box_to_subrange(access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()));
const auto oob_sr = subrange<3>(oob_min, range_cast<3>(oob_max - oob_min));
const auto buffer_id = access_map.get_nth_access(i).first;
const auto buffer_name = m_buffer_mngr.get_debug_name(buffer_id);
CELERITY_ERROR("Out-of-bounds access in kernel '{}' detected: Accessor {} for buffer {} attempted to access indices between {} which are "
"outside of mapped subrange {}",
tsk->get_debug_name(), i, access_map.get_nth_access(i).first, oob_sr, acc_sr);
"outside of mapped subrange {}", tsk->get_debug_name(), i, (buffer_name.empty() ? fmt::format("{}", buffer_id) : buffer_name), oob_sr, acc_sr);
}
sycl::free(m_oob_indices_per_accessor[i], m_queue.get_sycl_queue());
}
Expand Down
78 changes: 68 additions & 10 deletions test/accessor_tests.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -650,16 +650,20 @@ namespace detail {
template <int>
class acc_out_of_bounds_kernel {};

TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "device accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
#endif
buffer<int, Dims> buff(range_cast<Dims>(range<3>{10, 20, 30}));
buffer<int, Dims> unnamed_buff(range_cast<Dims>(range<3>{10, 20, 30}));
buffer<int, Dims> named_buff(range_cast<Dims>(range<3>{10, 20, 30}));
const auto accessible_sr = subrange_cast<Dims>(subrange<3>{{5, 10, 15}, {1, 2, 3}});
const auto oob_idx_lo = id_cast<Dims>(id<3>{1, 2, 3});
const auto oob_idx_hi = id_cast<Dims>(id<3>{7, 13, 25});
const auto buffer_name = "oob";

celerity::debug::set_buffer_name(named_buff, buffer_name);

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

q.submit([&](handler& cgh) {
accessor acc(buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
accessor unnamed_acc(unnamed_buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);
accessor named_acc(named_buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only, celerity::no_init);

cgh.parallel_for<acc_out_of_bounds_kernel<Dims>>(range<Dims>(ones), [=](item<Dims>) {
acc[oob_idx_lo] = 0;
acc[oob_idx_hi] = 0;
unnamed_acc[oob_idx_lo] = 0;
unnamed_acc[oob_idx_hi] = 0;

named_acc[oob_idx_lo] = 0;
named_acc[oob_idx_hi] = 0;
});
});
q.slow_full_sync();
}

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)))};
const auto error_message = fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for "
"buffer 0 attempted to access indices between {} which are outside of mapped subrange {}",
Dims, attempted_sr, subrange_cast<3>(accessible_sr));
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(error_message));
const auto unnamed_error_message = fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 0 for buffer 0 attempted to "
"access indices between {} which are outside of mapped subrange {}", Dims, attempted_sr, subrange_cast<3>(accessible_sr));
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(unnamed_error_message));

const auto named_error_message = fmt::format("Out-of-bounds access in kernel 'celerity::detail::acc_out_of_bounds_kernel<{}>' detected: Accessor 1 for buffer {} attempted to "
"access indices between {} which are outside of mapped subrange {}", Dims, buffer_name, attempted_sr, subrange_cast<3>(accessible_sr));
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(named_error_message));
}

TEMPLATE_TEST_CASE_METHOD_SIG(oob_fixture, "host accessor reports out-of-bounds accesses", "[accessor][oob]", ((int Dims), Dims), 1, 2, 3) {
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
#endif
buffer<int, Dims> unnamed_buff(range_cast<Dims>(range<3>{10, 20, 30}));
buffer<int, Dims> named_buff(range_cast<Dims>(range<3>{10, 20, 30}));
const auto accessible_sr = subrange_cast<Dims>(subrange<3>{{5, 10, 15}, {1, 2, 3}});
const auto oob_idx_lo = id_cast<Dims>(id<3>{1, 2, 3});
const auto oob_idx_hi = id_cast<Dims>(id<3>{7, 13, 25});
const auto buffer_name = "oob";

celerity::debug::set_buffer_name(named_buff, buffer_name);

// we need to be careful about the ordering of the construction and destruction
// of the Celerity queue and the log capturing utility here
std::unique_ptr<celerity::test_utils::log_capture> lc;
{
distr_queue q;

lc = std::make_unique<celerity::test_utils::log_capture>();

q.submit([&](handler& cgh){
accessor unnamed_acc(unnamed_buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only_host_task, celerity::no_init);
accessor nambed_acc(named_buff, cgh, celerity::access::fixed(accessible_sr), celerity::write_only_host_task, celerity::no_init);

cgh.host_task(range<Dims>(ones), [=](partition<Dims>) {
unnamed_acc[oob_idx_lo] = 0;
unnamed_acc[oob_idx_hi] = 0;

nambed_acc[oob_idx_lo] = 0;
nambed_acc[oob_idx_hi] = 0;
});
});

q.slow_full_sync();
}

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)))};
const auto unnamed_error_message = fmt::format("Out-of-bounds access in host task detected: Accessor 0 for buffer 0 attempted to "
"access indices between {} which are outside of mapped subrange {}", attempted_sr, subrange_cast<3>(accessible_sr));
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(unnamed_error_message));

const auto named_error_message = fmt::format("Out-of-bounds access in host task detected: Accessor 1 for buffer {} attempted to "
"access indices between {} which are outside of mapped subrange {}", buffer_name, attempted_sr, subrange_cast<3>(accessible_sr));
CHECK_THAT(lc->get_log(), Catch::Matchers::ContainsSubstring(named_error_message));
}
} // namespace detail
} // namespace celerity

0 comments on commit 3c3c479

Please sign in to comment.