Skip to content

Commit

Permalink
fixup! Remove multi-pass mechanism (invoke CGFs only once)
Browse files Browse the repository at this point in the history
  • Loading branch information
@psalz
psalz committed Apr 17, 2023
1 parent 85ed4f2 commit 56b845c
Show file tree
Hide file tree
Showing 7 changed files with 60 additions and 131 deletions.
100 changes: 24 additions & 76 deletions include/accessor.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -155,7 +155,7 @@ class buffer_allocation_window {
friend class accessor;
};

#define CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR [[deprecated("Creating accessor from const buffer is deprecated, capture buffer by reference instead")]]
#define CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR [[deprecated("Creating accessor from const buffer is deprecated, capture buffer by reference instead")]]

/**
* Celerity wrapper around SYCL accessors.
Expand Down Expand Up @@ -186,14 +186,6 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
const property::no_init& /* no_init */)
: accessor(ctor_internal_tag(), buff, cgh, rmfn) {}

template <typename Functor, access_mode TagMode, access_mode TagModeNoInit>
accessor(buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn, const detail::access_tag<TagMode, TagModeNoInit, target::device> /* tag */,
const property_list& /* prop_list */) {
static_assert(detail::constexpr_false<Functor>,
"Currently it is not accepted to pass a property list to an accessor constructor. Please use the property celerity::no_init "
"as a last argument in the constructor");
}

template <int D = Dims, std::enable_if_t<D == 0, int> = 0>
accessor(buffer<DataT, Dims>& buff, handler& cgh) : accessor(buff, cgh, access::all()) {}

Expand All @@ -210,50 +202,32 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
: accessor(buff, cgh, access::all(), tag, prop_list) {}

template <typename Functor>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn)
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn)
: accessor(ctor_internal_tag(), buff, cgh, rmfn) {}

template <typename Functor, access_mode TagModeNoInit>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn, const detail::access_tag<Mode, TagModeNoInit, target::device> /* tag */)
: accessor(ctor_internal_tag(), buff, cgh, rmfn) {}

template <typename Functor, access_mode TagMode>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
const detail::access_tag<TagMode, Mode, target::device> /* tag */, const property::no_init& /* no_init */)
: accessor(ctor_internal_tag(), buff, cgh, rmfn) {}

template <typename Functor, access_mode TagMode, access_mode TagModeNoInit>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
const detail::access_tag<TagMode, TagModeNoInit, target::device> /* tag */, const property_list& /* prop_list */) {
static_assert(detail::constexpr_false<Functor>,
"Currently it is not accepted to pass a property list to an accessor constructor. Please use the property celerity::no_init "
"as a last argument in the constructor");
}

template <int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh) : accessor(buff, cgh, access::all()) {}

template <access_mode TagModeNoInit, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<Mode, TagModeNoInit, target::device> tag)
: accessor(buff, cgh, access::all(), tag) {}

template <access_mode TagMode, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<TagMode, Mode, target::device> tag, const property::no_init& no_init)
: accessor(buff, cgh, access::all(), tag, no_init) {}

template <access_mode TagMode, access_mode TagModeNoInit, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<TagMode, TagModeNoInit, target::device> tag, const property_list& prop_list)
: accessor(buff, cgh, access::all(), tag, prop_list) {}

// explicitly defaulted because we define operator=(value_type) for Dims == 0
accessor(accessor&&) noexcept = default;
accessor& operator=(accessor&&) noexcept = default;

#if !defined(__SYCL_DEVICE_ONLY__) && !defined(SYCL_DEVICE_ONLY)
#if !defined(__SYCL_DEVICE_ONLY__)
accessor(const accessor& other) { copy_and_hydrate(other); }

accessor& operator=(const accessor& other) {
Expand Down Expand Up @@ -341,13 +315,14 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
accessor(const ctor_internal_tag /* tag */, const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn) {
using range_mapper = detail::range_mapper<Dims, std::decay_t<Functor>>; // decay function type to function pointer
const auto hid = detail::add_requirement(cgh, detail::get_buffer_id(buff), std::make_unique<range_mapper>(rmfn, Mode, buff.get_range()));
detail::extend_lifetime(cgh, detail::get_lifetime_extending_state(buff));
detail::extend_lifetime(cgh, std::move(detail::get_lifetime_extending_state(buff)));
m_device_ptr = detail::embed_hydration_id<DataT*>(hid);
}

// Constructor for tests, called through accessor_testspy.
accessor(DataT* ptr, id<Dims> index_offset, range<Dims> buffer_range) : m_device_ptr(ptr), m_index_offset(index_offset), m_buffer_range(buffer_range) {
#if defined(__SYCL_DEVICE_ONLY__) || defined(SYCL_DEVICE_ONLY)
accessor(DataT* const ptr, const id<Dims>& index_offset, const range<Dims>& buffer_range)
: m_device_ptr(ptr), m_index_offset(index_offset), m_buffer_range(buffer_range) {
#if defined(__SYCL_DEVICE_ONLY__)
#if CELERITY_WORKAROUND_HIPSYCL // hipSYCL does not yet implement is_device_copyable_v
static_assert(std::is_trivially_copyable_v<accessor>);
#else
Expand All @@ -357,15 +332,15 @@ class accessor<DataT, Dims, Mode, target::device> : public detail::accessor_base
}

// Constructor for tests, called through accessor_testspy.
accessor(detail::hydration_id hid, id<Dims> index_offset, range<Dims> buffer_range)
accessor(const detail::hydration_id hid, const id<Dims>& index_offset, const range<Dims>& buffer_range)
: accessor(detail::embed_hydration_id<DataT*>(hid), index_offset, buffer_range) {}

void copy_and_hydrate(const accessor& other) {
m_device_ptr = other.m_device_ptr;
m_index_offset = other.m_index_offset;
m_buffer_range = other.m_buffer_range;

#if !defined(__SYCL_DEVICE_ONLY__) && !defined(SYCL_DEVICE_ONLY)
#if !defined(__SYCL_DEVICE_ONLY__)
if(detail::is_embedded_hydration_id(m_device_ptr)) {
if(detail::cgf_diagnostics::is_available()) {
detail::cgf_diagnostics::get_instance().register_accessor(detail::extract_hydration_id(m_device_ptr), target::device);
Expand Down Expand Up @@ -411,14 +386,6 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
const property::no_init& /* no_init */)
: accessor(ctor_internal_tag{}, buff, cgh, rmfn) {}

template <typename Functor, access_mode TagMode, access_mode TagModeNoInit>
accessor(buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn, const detail::access_tag<TagMode, TagModeNoInit, target::host_task> /* tag */,
const property_list& /* prop_list */) {
static_assert(detail::constexpr_false<Functor>,
"Currently it is not accepted to pass a property list to an accessor constructor. Please use the property celerity::no_init "
"as a last argument in the constructor");
}

template <int D = Dims, std::enable_if_t<D == 0, int> = 0>
accessor(buffer<DataT, Dims>& buff, handler& cgh) : accessor(buff, cgh, access::all()) {}

Expand All @@ -435,11 +402,11 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
: accessor(buff, cgh, access::all(), tag, prop_list) {}

template <typename Functor>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn)
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn)
: accessor(ctor_internal_tag{}, buff, cgh, rmfn) {}

template <typename Functor, access_mode TagModeNoInit>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn, const detail::access_tag<Mode, TagModeNoInit, target::host_task> /* tag */)
: accessor(ctor_internal_tag{}, buff, cgh, rmfn) {}

Expand All @@ -448,36 +415,18 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
* but once they do this should be replace for a constructor that takes a prop list as an argument.
*/
template <typename Functor, access_mode TagMode, access_mode M = Mode, typename = std::enable_if_t<detail::access::mode_traits::is_producer(M)>>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
const detail::access_tag<TagMode, Mode, target::host_task> /* tag */, const property::no_init& /* no_init */)
: accessor(ctor_internal_tag{}, buff, cgh, rmfn) {}

template <typename Functor, access_mode TagMode, access_mode TagModeNoInit>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh, const Functor& rmfn,
const detail::access_tag<TagMode, TagModeNoInit, target::host_task> /* tag */, const property_list& /* prop_list */) {
static_assert(detail::constexpr_false<Functor>,
"Currently it is not accepted to pass a property list to an accessor constructor. Please use the property celerity::no_init "
"as a last argument in the constructor");
}

template <int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(const buffer<DataT, Dims>& buff, handler& cgh) : accessor(buff, cgh, access::all()) {}

template <access_mode TagModeNoInit, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<Mode, TagModeNoInit, target::host_task> tag)
: accessor(buff, cgh, access::all(), tag) {}

template <access_mode TagMode, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<TagMode, Mode, target::host_task> tag, const property::no_init& no_init)
: accessor(buff, cgh, access::all(), tag, no_init) {}

template <access_mode TagMode, access_mode TagModeNoInit, int D = Dims, std::enable_if_t<D == 0, int> = 0>
CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR accessor(
const buffer<DataT, Dims>& buff, handler& cgh, const detail::access_tag<TagMode, TagModeNoInit, target::host_task> tag, const property_list& prop_list)
: accessor(buff, cgh, access::all(), tag, prop_list) {}

// explicitly defaulted because we define operator=(value_type) for Dims == 0
accessor(accessor&&) noexcept = default;
accessor& operator=(accessor&&) noexcept = default;
Expand Down Expand Up @@ -650,22 +599,21 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b

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()) {
static_assert(!std::is_same_v<Functor, range<Dims>>, "The accessor constructor overload for master-access tasks (now called 'host tasks') has "
"been removed with Celerity 0.2.0. Please provide a range mapper instead.");
using range_mapper = detail::range_mapper<Dims, std::decay_t<Functor>>; // decay function type to function pointer
const auto hid = detail::add_requirement(cgh, detail::get_buffer_id(buff), std::make_unique<range_mapper>(rmfn, Mode, buff.get_range()));
detail::extend_lifetime(cgh, detail::get_lifetime_extending_state(buff));
detail::extend_lifetime(cgh, std::move(detail::get_lifetime_extending_state(buff)));
m_host_ptr = detail::embed_hydration_id<DataT*>(hid);
}

// Constructor for tests, called through accessor_testspy.
accessor(subrange<Dims> mapped_subrange, DataT* ptr, id<Dims> backing_buffer_offset, range<Dims> backing_buffer_range, range<Dims> virtual_buffer_range)
accessor(const subrange<Dims> mapped_subrange, DataT* const ptr, const id<Dims>& backing_buffer_offset, const range<Dims>& backing_buffer_range,
const range<Dims>& virtual_buffer_range)
: m_mapped_subrange(mapped_subrange), m_index_offset(backing_buffer_offset), m_buffer_range(backing_buffer_range),
m_virtual_buffer_range(virtual_buffer_range), m_host_ptr(ptr) {}

// Constructor for tests, called through accessor_testspy.
accessor(subrange<Dims> mapped_subrange, detail::hydration_id hid, id<Dims> backing_buffer_offset, range<Dims> backing_buffer_range,
range<Dims> virtual_buffer_range)
accessor(const subrange<Dims>& mapped_subrange, const detail::hydration_id hid, const id<Dims>& backing_buffer_offset,
const range<Dims>& backing_buffer_range, range<Dims> virtual_buffer_range)
: accessor(mapped_subrange, detail::embed_hydration_id<DataT*>(hid), backing_buffer_offset, backing_buffer_range, virtual_buffer_range) {}

void copy_and_hydrate(const accessor& other) {
Expand Down Expand Up @@ -693,7 +641,7 @@ class accessor<DataT, Dims, Mode, target::host_task> : public detail::accessor_b
size_t get_linear_offset(const id<Dims>& index) const { return detail::get_linear_index(m_buffer_range, index - m_index_offset); }
};

#undef CELERITY_INTERNAL_ACCESSOR_DEPRECATED_CTOR
#undef CELERITY_DETAIL_ACCESSOR_DEPRECATED_CTOR

// TODO: Make buffer non-const once corresponding (deprecated!) constructor overloads are removed
template <typename T, int D, typename Functor, access_mode Mode, access_mode ModeNoInit, target Target>
Expand Down Expand Up @@ -746,7 +694,7 @@ class local_accessor {
template <int D = Dims, typename = std::enable_if_t<D == 0>>
local_accessor(handler& cgh) : local_accessor(range<0>(), cgh) {}

#if !defined(__SYCL_DEVICE_ONLY__) && !defined(SYCL_DEVICE_ONLY)
#if !defined(__SYCL_DEVICE_ONLY__)
local_accessor(const range<Dims>& allocation_size, handler& cgh) : m_sycl_acc{}, m_allocation_size(allocation_size) {}

local_accessor(const local_accessor& other)
Expand Down Expand Up @@ -826,7 +774,7 @@ class local_accessor {
sycl::range<sycl_dims> sycl_allocation_size() const { return sycl::range<sycl_dims>(detail::range_cast<sycl_dims>(m_allocation_size)); }

// Constructor for tests, called through accessor_testspy.
local_accessor(const range<Dims>& allocation_size) : m_sycl_acc{}, m_allocation_size(allocation_size) {}
explicit local_accessor(const range<Dims>& allocation_size) : m_sycl_acc{}, m_allocation_size(allocation_size) {}
};

} // namespace celerity
14 changes: 2 additions & 12 deletions include/buffer.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -85,28 +85,18 @@ class buffer : public detail::lifetime_extending_state_wrapper {
return get_access<Mode, target::device, Functor>(cgh, rmfn);
}

template <access_mode Mode, typename Functor, int D = Dims, std::enable_if_t<D == 0, int> = 0>
[[deprecated("Calling get_access on a const buffer is deprecated")]] accessor<DataT, Dims, Mode, target::device> get_access(handler& cgh) const {
return get_access<Mode, target::device, Functor>(cgh);
}

template <access_mode Mode, target Target, typename Functor, int D = Dims, std::enable_if_t<(D > 0), int> = 0>
[[deprecated("Calling get_access on a const buffer is deprecated")]] accessor<DataT, Dims, Mode, Target> get_access(handler& cgh, Functor rmfn) const {
return accessor<DataT, Dims, Mode, Target>(*this, cgh, rmfn);
}

template <access_mode Mode, target Target, typename Functor, int D = Dims, std::enable_if_t<D == 0, int> = 0>
[[deprecated("Calling get_access on a const buffer is deprecated")]] accessor<DataT, Dims, Mode, Target> get_access(handler& cgh) const {
return accessor<DataT, Dims, Mode, Target>(*this, cgh);
}

const range<Dims>& get_range() const { return m_impl->range; }

protected:
std::shared_ptr<detail::lifetime_extending_state> get_lifetime_extending_state() const override { return m_impl; }

private:
struct impl : public detail::lifetime_extending_state {
struct impl final : public detail::lifetime_extending_state {
impl(range<Dims> rng, const DataT* host_init_ptr) : range(rng) {
if(!detail::runtime::is_initialized()) { detail::runtime::init(nullptr, nullptr); }
id = detail::runtime::get_instance().get_buffer_manager().register_buffer<DataT, Dims>(detail::range_cast<3>(range), host_init_ptr);
Expand All @@ -115,7 +105,7 @@ class buffer : public detail::lifetime_extending_state_wrapper {
impl(impl&&) = delete;
impl& operator=(const impl&) = delete;
impl& operator=(impl&&) = delete;
~impl() noexcept override { detail::runtime::get_instance().get_buffer_manager().unregister_buffer(id); }
~impl() override { detail::runtime::get_instance().get_buffer_manager().unregister_buffer(id); }
detail::buffer_id id;
celerity::range<Dims> range;
std::string debug_name;
Expand Down
Loading

0 comments on commit 56b845c

Please sign in to comment.