Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add out-of-bounds check for host side accessors #211

Merged
merged 1 commit into from
Sep 21, 2023
Merged

Add out-of-bounds check for host side accessors #211

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 34 additions & 0 deletions include/accessor.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -406,6 +406,22 @@ 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 bool is_within_bounds_lo = all_true(index >= m_accessed_virtual_subrange.offset);
const bool is_within_bounds_hi = all_true(index < (m_accessed_virtual_subrange.offset + m_accessed_virtual_subrange.range));

if((!is_within_bounds_lo || !is_within_bounds_hi)) {
std::lock_guard<std::mutex> guard(m_oob_mutex);
for(int d = 0; d < Dims; ++d) {
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 @@ -521,6 +537,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;
GagaLP marked this conversation as resolved.
Show resolved Hide resolved

// 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 @@ -547,6 +573,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 @@ -558,6 +588,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
Open in desktop
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
55 changes: 47 additions & 8 deletions src/worker_job.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -160,7 +160,21 @@ namespace detail {
const auto [bid, mode] = access_map.get_nth_access(i);
const auto sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
const auto info = m_buffer_mngr.access_host_buffer(bid, mode, sr);

#if CELERITY_ACCESSOR_BOUNDARY_CHECK
// oob_indices[0] contains the lower bound oob indices
// oob_indices[1] contains the upper bound oob indices
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>{0, 0, 0};
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 +189,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>{0, 0, 0}) {
const auto& access_map = tsk->get_buffer_access_map();
const auto acc_sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
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 +252,16 @@ 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);
// oob_indices[0] contains the lower bound oob indices
// oob_indices[1] contains the upper bound oob indices
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>{0, 0, 0};
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 @@ -259,13 +296,15 @@ namespace detail {
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}) {
if(oob_max != id<3>{0, 0, 0}) {
const auto& access_map = tsk->get_buffer_access_map();
const auto acc_sr = access_map.get_requirements_for_nth_access(i, tsk->get_dimensions(), data.sr, tsk->get_global_size()).get_subrange();
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);
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
88 changes: 77 additions & 11 deletions test/accessor_tests.cc
View file
Open in desktop
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) {
PeterTh marked this conversation as resolved.
Show resolved Hide resolved
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
#endif
buffer<int, Dims> buff(test_utils::truncate_range<Dims>({10, 20, 30}));
buffer<int, Dims> unnamed_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
buffer<int, Dims> named_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
const auto accessible_sr = test_utils::truncate_subrange<Dims>({{5, 10, 15}, {1, 2, 3}});
const auto oob_idx_lo = test_utils::truncate_id<Dims>({1, 2, 3});
const auto oob_idx_hi = test_utils::truncate_id<Dims>({7, 13, 25});
const auto buffer_name = "oob";

// we need to be careful about the orderign of the construction and destruction
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;
{
Expand All @@ -668,21 +672,83 @@ 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 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))) - range_cast<3>(range<Dims>(zeros))};
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) {
PeterTh marked this conversation as resolved.
Show resolved Hide resolved
#if !CELERITY_ACCESSOR_BOUNDARY_CHECK
SKIP("CELERITY_ACCESSOR_BOUNDARY_CHECK=0");
#endif
buffer<int, Dims> unnamed_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
buffer<int, Dims> named_buff(test_utils::truncate_range<Dims>({10, 20, 30}));
const auto accessible_sr = test_utils::truncate_subrange<Dims>({{5, 10, 15}, {1, 2, 3}});
const auto oob_idx_lo = test_utils::truncate_id<Dims>({1, 2, 3});
const auto oob_idx_hi = test_utils::truncate_id<Dims>({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))) - range_cast<3>(range<Dims>(zeros))};
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