Skip to content

[SYCL] Implement missed cl::sycl::buffer functionality #56

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

Merged
merged 4 commits into from
Apr 1, 2019
Merged

[SYCL] Implement missed cl::sycl::buffer functionality #56

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
e37640e
[SYCL] Add support for custom allocators in buffer
Fznamznon Mar 29, 2019
78abab4
[SYCL] Implement buffer::get_access with range and offset
Fznamznon Mar 29, 2019
639d51f
[SYCL] Implement buffer::has_property and buffer::get_property
Fznamznon Mar 29, 2019
3034366
[SYCL] Implement buffer::set_write_back
Fznamznon Mar 29, 2019
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
10 changes: 5 additions & 5 deletions sycl/include/CL/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -150,9 +150,9 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
// reinterpret casting while setting kernel arguments in order to get cl_mem
// value from the buffer regardless of the accessor's dimensionality.
#ifndef __SYCL_DEVICE_ONLY__
detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
#else
char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
#endif // __SYCL_DEVICE_ONLY__

dataT *Data;
Expand Down Expand Up @@ -185,9 +185,9 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
// reinterpret casting while setting kernel arguments in order to get cl_mem
// value from the buffer regardless of the accessor's dimensionality.
#ifndef __SYCL_DEVICE_ONLY__
detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
#else
char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
#endif // __SYCL_DEVICE_ONLY__

dataT *Data;
Expand Down Expand Up @@ -763,7 +763,7 @@ class accessor
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else // !__SYCL_DEVICE_ONLY__
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
: __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
bufferRef.get_range(), Offset) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
Expand Down
98 changes: 54 additions & 44 deletions sycl/include/CL/sycl/buffer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ class queue;
template <int dimentions> class range;

template <typename T, int dimensions = 1,
typename AllocatorT = cl::sycl::buffer_allocator<char>>
typename AllocatorT = cl::sycl::buffer_allocator>
class buffer {
public:
using value_type = T;
Expand All @@ -36,11 +36,11 @@ class buffer {
get_count() * sizeof(T), propList);
}

// buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
// const property_list &propList = {}) {
// impl = std::make_shared<detail::buffer_impl>(bufferRange, allocator,
// propList);
// }
buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
const property_list &propList = {}) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
get_count() * sizeof(T), propList, allocator);
}

buffer(T *hostData, const range<dimensions> &bufferRange,
const property_list &propList = {})
Expand All @@ -49,11 +49,11 @@ class buffer {
hostData, get_count() * sizeof(T), propList);
}

// buffer(T *hostData, const range<dimensions> &bufferRange,
// AllocatorT allocator, const property_list &propList = {}) {
// impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
// allocator, propList);
// }
buffer(T *hostData, const range<dimensions> &bufferRange,
AllocatorT allocator, const property_list &propList = {}) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList, allocator);
}

buffer(const T *hostData, const range<dimensions> &bufferRange,
const property_list &propList = {})
Expand All @@ -62,18 +62,18 @@ class buffer {
hostData, get_count() * sizeof(T), propList);
}

// buffer(const T *hostData, const range<dimensions> &bufferRange,
// AllocatorT allocator, const property_list &propList = {}) {
// impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
// allocator, propList);
// }
buffer(const T *hostData, const range<dimensions> &bufferRange,
AllocatorT allocator, const property_list &propList = {}) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList, allocator);
}

// buffer(const shared_ptr_class<T> &hostData,
// const range<dimensions> &bufferRange, AllocatorT allocator,
// const property_list &propList = {}) {
// impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
// allocator, propList);
// }
buffer(const shared_ptr_class<T> &hostData,
const range<dimensions> &bufferRange, AllocatorT allocator,
const property_list &propList = {}) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList, allocator);
}

buffer(const shared_ptr_class<T> &hostData,
const range<dimensions> &bufferRange,
Expand All @@ -83,12 +83,13 @@ class buffer {
hostData, get_count() * sizeof(T), propList);
}

// template <class InputIterator>
// buffer<T, 1>(InputIterator first, InputIterator last, AllocatorT allocator,
// const property_list &propList = {}) {
// impl = std::make_shared<detail::buffer_impl>(first, last, allocator,
// propList);
// }
template <class InputIterator>
buffer(InputIterator first, InputIterator last, AllocatorT allocator,
const property_list &propList = {})
: Range(range<1>(std::distance(first, last))) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
first, last, get_count() * sizeof(T), propList, allocator);
}

template <class InputIterator, int N = dimensions,
typename = std::enable_if<N == 1>>
Expand Down Expand Up @@ -135,7 +136,7 @@ class buffer {

size_t get_size() const { return impl->get_size(); }

// AllocatorT get_allocator() const { return impl->get_allocator(); }
AllocatorT get_allocator() const { return impl->get_allocator(); }

template <access::mode mode,
access::target target = access::target::global_buffer>
Expand All @@ -152,28 +153,29 @@ class buffer {
return impl->template get_access<T, dimensions, mode>(*this);
}

// template <access::mode mode, access::target target =
// access::target::global_buffer> accessor<T, dimensions, mode, target,
// access::placeholder::false_t> get_access( handler &commandGroupHandler,
// range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
// return impl->get_access(commandGroupHandler, accessRange,
// accessOffset);
// }
template <access::mode mode,
access::target target = access::target::global_buffer>
accessor<T, dimensions, mode, target, access::placeholder::false_t>
get_access(handler &commandGroupHandler, range<dimensions> accessRange,
id<dimensions> accessOffset = {}) {
return impl->template get_access<T, dimensions, mode, target>(
*this, commandGroupHandler, accessRange, accessOffset);
}

// template <access::mode mode>
// accessor<T, dimensions, mode, access::target::host_buffer,
// access::placeholder::false_t> get_access( range<dimensions> accessRange,
// id<dimensions> accessOffset = {}) {
// return impl->get_access(accessRange, accessOffset);
// }
template <access::mode mode>
accessor<T, dimensions, mode, access::target::host_buffer,
access::placeholder::false_t>
get_access(range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
return impl->template get_access<T, dimensions, mode>(*this, accessRange,
accessOffset);
}

template <typename Destination = std::nullptr_t>
void set_final_data(Destination finalData = nullptr) {
impl->set_final_data(finalData);
}

// void set_write_back(bool flag = true) { return impl->set_write_back(flag);
// }
void set_write_back(bool flag = true) { return impl->set_write_back(flag); }

// bool is_sub_buffer() const { return impl->is_sub_buffer(); }

Expand All @@ -189,6 +191,14 @@ class buffer {
reinterpretRange);
}

template <typename propertyT> bool has_property() const {
return impl->template has_property<propertyT>();
}

template <typename propertyT> propertyT get_property() const {
return impl->template get_property<propertyT>();
}

private:
shared_ptr_class<detail::buffer_impl<AllocatorT>> impl;
template <class Obj>
Expand Down
65 changes: 52 additions & 13 deletions sycl/include/CL/sycl/detail/buffer_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,16 +37,18 @@ class handler;
class queue;
template <int dimentions> class id;
template <int dimentions> class range;
template <class T> using buffer_allocator = std::allocator<T>;
using buffer_allocator = std::allocator<char>;
namespace detail {
template <typename AllocatorT> class buffer_impl {
public:
buffer_impl(const size_t sizeInBytes, const property_list &propList)
: buffer_impl((void *)nullptr, sizeInBytes, propList) {}
buffer_impl(const size_t sizeInBytes, const property_list &propList,
AllocatorT allocator = AllocatorT())
: buffer_impl((void *)nullptr, sizeInBytes, propList, allocator) {}

buffer_impl(void *hostData, const size_t sizeInBytes,
const property_list &propList)
: SizeInBytes(sizeInBytes), Props(propList) {
const property_list &propList,
AllocatorT allocator = AllocatorT())
: SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
if (Props.has_property<property::buffer::use_host_ptr>()) {
BufPtr = hostData;
} else {
Expand All @@ -62,8 +64,9 @@ template <typename AllocatorT> class buffer_impl {

// TODO temporary solution for allowing initialisation with const data
buffer_impl(const void *hostData, const size_t sizeInBytes,
const property_list &propList)
: SizeInBytes(sizeInBytes), Props(propList) {
const property_list &propList,
AllocatorT allocator = AllocatorT())
: SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
if (Props.has_property<property::buffer::use_host_ptr>()) {
// TODO make this buffer read only
BufPtr = const_cast<void *>(hostData);
Expand All @@ -79,8 +82,9 @@ template <typename AllocatorT> class buffer_impl {

template <typename T>
buffer_impl(const shared_ptr_class<T> &hostData, const size_t sizeInBytes,
const property_list &propList)
: SizeInBytes(sizeInBytes), Props(propList) {
const property_list &propList,
AllocatorT allocator = AllocatorT())
: SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
if (Props.has_property<property::buffer::use_host_ptr>()) {
BufPtr = hostData.get();
} else {
Expand All @@ -97,8 +101,9 @@ template <typename AllocatorT> class buffer_impl {

template <class InputIterator>
buffer_impl(InputIterator first, InputIterator last, const size_t sizeInBytes,
const property_list &propList)
: SizeInBytes(sizeInBytes), Props(propList) {
const property_list &propList,
AllocatorT allocator = AllocatorT())
: SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
if (Props.has_property<property::buffer::use_host_ptr>()) {
// TODO next line looks unsafe
BufPtr = &*first;
Expand Down Expand Up @@ -140,7 +145,7 @@ template <typename AllocatorT> class buffer_impl {
.copyBack<access::mode::read_write, access::target::host_buffer>(
*this);

if (uploadData != nullptr) {
if (uploadData != nullptr && NeedWriteBack) {
uploadData();
}

Expand Down Expand Up @@ -170,7 +175,7 @@ template <typename AllocatorT> class buffer_impl {
throw cl::sycl::runtime_error(
"set_final_data could not be used with interoperability buffer");
static_assert(!std::is_const<Destination>::value,
"Сan not write in a constant Destination. Destination should "
"Can not write in a constant Destination. Destination should "
"not be const.");
uploadData = [this, final_data]() mutable {
auto *Ptr =
Expand All @@ -182,6 +187,10 @@ template <typename AllocatorT> class buffer_impl {
};
}

void set_write_back(bool flag) { NeedWriteBack = flag; }

AllocatorT get_allocator() const { return MAllocator; }

template <typename T, int dimensions, access::mode mode,
access::target target = access::target::global_buffer>
accessor<T, dimensions, mode, target, access::placeholder::false_t>
Expand All @@ -199,6 +208,34 @@ template <typename AllocatorT> class buffer_impl {
access::placeholder::false_t>(Buffer);
}

template <typename T, int dimensions, access::mode mode,
access::target target = access::target::global_buffer>
accessor<T, dimensions, mode, target, access::placeholder::false_t>
get_access(buffer<T, dimensions, AllocatorT> &Buffer,
handler &commandGroupHandler, range<dimensions> accessRange,
id<dimensions> accessOffset) {
return accessor<T, dimensions, mode, target, access::placeholder::false_t>(
Buffer, commandGroupHandler, accessRange, accessOffset);
}

template <typename T, int dimensions, access::mode mode>
accessor<T, dimensions, mode, access::target::host_buffer,
access::placeholder::false_t>
get_access(buffer<T, dimensions, AllocatorT> &Buffer,
range<dimensions> accessRange, id<dimensions> accessOffset) {
return accessor<T, dimensions, mode, access::target::host_buffer,
access::placeholder::false_t>(Buffer, accessRange,
accessOffset);
}

template <typename propertyT> bool has_property() const {
return Props.has_property<propertyT>();
}

template <typename propertyT> propertyT get_property() const {
return Props.get_property<propertyT>();
}

public:
void moveMemoryTo(QueueImplPtr Queue, std::vector<cl::sycl::event> DepEvents,
EventImplPtr Event);
Expand Down Expand Up @@ -243,8 +280,10 @@ template <typename AllocatorT> class buffer_impl {
// This field must be the first to guarantee that it's safe to use
// reinterpret casting while setting kernel arguments in order to get cl_mem
// value from the buffer regardless of its dimensionality.
AllocatorT MAllocator;
OpenCLMemState OCLState;
bool OpenCLInterop = false;
bool NeedWriteBack = true;
event AvailableEvent;
cl_context OpenCLContext = nullptr;
void *BufPtr = nullptr;
Expand Down
11 changes: 5 additions & 6 deletions sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,8 +61,7 @@ template <typename dataT, int dimensions, access::mode accessMode,
void Node::addAccRequirement(
accessor<dataT, dimensions, accessMode, accessTarget, isPlaceholder> &&Acc,
int argIndex) {
detail::buffer_impl<buffer_allocator<char>> *buf =
Acc.__get_impl()->m_Buf;
detail::buffer_impl<buffer_allocator> *buf = Acc.__get_impl()->m_Buf;
addBufRequirement<accessMode, accessTarget>(*buf);
addInteropArg(nullptr, buf->get_size(), argIndex,
getReqForBuffer(m_Bufs, *buf));
Expand Down Expand Up @@ -134,7 +133,7 @@ void Node::addExplicitMemOp(
auto *DestBase = Dest.__get_impl();
assert(DestBase != nullptr &&
"Accessor should have an initialized accessor_base");
detail::buffer_impl<buffer_allocator<char>> *Buf = DestBase->m_Buf;
detail::buffer_impl<buffer_allocator> *Buf = DestBase->m_Buf;

range<Dimensions> Range = DestBase->AccessRange;
id<Dimensions> Offset = DestBase->Offset;
Expand Down Expand Up @@ -162,10 +161,10 @@ void Node::addExplicitMemOp(
assert(DestBase != nullptr &&
"Accessor should have an initialized accessor_base");

detail::buffer_impl<buffer_allocator<char>> *SrcBuf = SrcBase->m_Buf;
detail::buffer_impl<buffer_allocator> *SrcBuf = SrcBase->m_Buf;
assert(SrcBuf != nullptr &&
"Accessor should have an initialized buffer_impl");
detail::buffer_impl<buffer_allocator<char>> *DestBuf = DestBase->m_Buf;
detail::buffer_impl<buffer_allocator> *DestBuf = DestBase->m_Buf;
assert(DestBuf != nullptr &&
"Accessor should have an initialized buffer_impl");

Expand Down Expand Up @@ -195,7 +194,7 @@ void Scheduler::updateHost(
auto *AccBase = Acc.__get_impl();
assert(AccBase != nullptr &&
"Accessor should have an initialized accessor_base");
detail::buffer_impl<buffer_allocator<char>> *Buf = AccBase->m_Buf;
detail::buffer_impl<buffer_allocator> *Buf = AccBase->m_Buf;

updateHost<mode, tgt>(*Buf, Event);
}
Expand Down
Loading