Skip to content

[SYCL] Implement cl::sycl::buffer::reinterpret #33

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 2 commits into from
Feb 21, 2019
Merged

[SYCL] Implement cl::sycl::buffer::reinterpret #33

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
2374893
[SYCL] Implement cl::sycl::buffer::reinterpret
Fznamznon Feb 15, 2019
45e1548
Merge branch 'sycl' into implement-reinterpret
vladimirlaz Feb 21, 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
96 changes: 47 additions & 49 deletions sycl/include/CL/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -125,11 +125,12 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions == 0) {
SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions > 0) {
dataT *Data;
range<dimensions> Range;
range<dimensions> BufRange;
id<dimensions> Offset;

accessor_impl(dataT *Data, range<dimensions> Range,
id<dimensions> Offset = {})
: Data(Data), Range(Range), Offset(Offset) {}
accessor_impl(dataT * Data, range<dimensions> Range,
range<dimensions> BufRange, id<dimensions> Offset = {})
: Data(Data), Range(Range), BufRange(BufRange), Offset(Offset) {}

// Returns the number of accessed elements.
size_t get_count() const { return Range.size(); }
Expand All @@ -146,10 +147,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<dataT, 1> *m_Buf = nullptr;

detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
#else
char padding[sizeof(detail::buffer_impl<dataT, dimensions> *)];
char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
#endif // __SYCL_DEVICE_ONLY__

dataT *Data;
Expand Down Expand Up @@ -182,22 +182,23 @@ 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<dataT, dimensions> *m_Buf = nullptr;
detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
#else
char padding[sizeof(detail::buffer_impl<dataT, dimensions> *)];
char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
#endif // __SYCL_DEVICE_ONLY__

dataT *Data;
range<dimensions> Range;
range<dimensions> BufRange;
id<dimensions> Offset;

// Device accessors must be associated with a command group handler.
// The handler though can be nullptr at the creation point if the
// accessor is a placeholder accessor.
accessor_impl(dataT *Data, range<dimensions> Range,
handler *Handler = nullptr, id<dimensions> Offset = {})
: Data(Data), Range(Range), Offset(Offset)
{}
accessor_impl(dataT * Data, range<dimensions> Range,
range<dimensions> BufRange, handler *Handler = nullptr,
id<dimensions> Offset = {})
: Data(Data), Range(Range), BufRange(BufRange), Offset(Offset) {}

// Returns the number of accessed elements.
size_t get_count() const { return Range.size(); }
Expand Down Expand Up @@ -633,8 +634,8 @@ 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,
detail::getSyclObjImpl(bufferRef)->Range,
: __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range(),
&commandGroupHandlerRef) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
Expand Down Expand Up @@ -669,8 +670,8 @@ class accessor
AccessTarget == access::target::constant_buffer))) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef)
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
detail::getSyclObjImpl(bufferRef)->Range) {
: __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range()) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
if (BufImpl->OpenCLInterop) {
Expand Down Expand Up @@ -701,17 +702,17 @@ class accessor
access::target AccessTarget = accessTarget,
access::placeholder IsPlaceholder = isPlaceholder>
accessor(typename std::enable_if<
(IsPlaceholder == access::placeholder::false_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef,
(IsPlaceholder == access::placeholder::false_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef,
handler &commandGroupHandlerRef)
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
detail::getSyclObjImpl(bufferRef)->Range,
: __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range(),
&commandGroupHandlerRef) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
Expand Down Expand Up @@ -739,20 +740,19 @@ class accessor
access::target AccessTarget = accessTarget,
access::placeholder IsPlaceholder = isPlaceholder>
accessor(typename std::enable_if<
((IsPlaceholder == access::placeholder::false_t &&
AccessTarget == access::target::host_buffer) ||
(IsPlaceholder == access::placeholder::true_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0)),
buffer<DataT, Dimensions>>::type &bufferRef,
range<Dimensions> Range,
id<Dimensions> Offset = {}
)
((IsPlaceholder == access::placeholder::false_t &&
AccessTarget == access::target::host_buffer) ||
(IsPlaceholder == access::placeholder::true_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0)),
buffer<DataT, Dimensions>>::type &bufferRef,
range<Dimensions> Range, id<Dimensions> Offset = {})
#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, Offset) {
#else // !__SYCL_DEVICE_ONLY__
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
bufferRef.get_range(), Offset) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
if (BufImpl->OpenCLInterop) {
Expand All @@ -769,7 +769,7 @@ class accessor
"interoperability buffer");
}
}
#endif // !__SYCL_DEVICE_ONLY__
#endif // !__SYCL_DEVICE_ONLY__

// buffer ctor #6:
// accessor(buffer &, handler &, range Range, id Offset);
Expand All @@ -784,20 +784,18 @@ class accessor
access::target AccessTarget = accessTarget,
access::placeholder IsPlaceholder = isPlaceholder>
accessor(typename std::enable_if<
(IsPlaceholder == access::placeholder::false_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef,
handler &commandGroupHandlerRef,
range<Dimensions> Range,
id<Dimensions> Offset = {}
)
(IsPlaceholder == access::placeholder::false_t &&
(AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef,
handler &commandGroupHandlerRef, range<Dimensions> Range,
id<Dimensions> Offset = {})
#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,
&commandGroupHandlerRef, Offset) {
#else // !__SYCL_DEVICE_ONLY__
: __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
bufferRef.get_range(), &commandGroupHandlerRef, Offset) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
throw cl::sycl::runtime_error(
Expand All @@ -807,7 +805,7 @@ class accessor
commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
__impl.m_Buf = BufImpl.get();
}
#endif // !__SYCL_DEVICE_ONLY__
#endif // !__SYCL_DEVICE_ONLY__

// TODO:
// local accessor ctor #1
Expand Down
78 changes: 48 additions & 30 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<T>>
typename AllocatorT = cl::sycl::buffer_allocator<char>>
class buffer {
public:
using value_type = T;
Expand All @@ -30,9 +30,10 @@ class buffer {
using allocator_type = AllocatorT;

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

// buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
Expand All @@ -42,9 +43,10 @@ class buffer {
// }

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

// buffer(T *hostData, const range<dimensions> &bufferRange,
Expand All @@ -54,9 +56,10 @@ class buffer {
// }

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

// buffer(const T *hostData, const range<dimensions> &bufferRange,
Expand All @@ -74,9 +77,10 @@ class buffer {

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

// template <class InputIterator>
Expand All @@ -89,9 +93,10 @@ class buffer {
template <class InputIterator, int N = dimensions,
typename = std::enable_if<N == 1>>
buffer(InputIterator first, InputIterator last,
const property_list &propList = {}) {
impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
first, last, propList);
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);
}

// buffer(buffer<T, dimensions, AllocatorT> b, const id<dimensions>
Expand All @@ -102,7 +107,7 @@ class buffer {
template <int N = dimensions, typename = std::enable_if<N == 1>>
buffer(cl_mem MemObject, const context &SyclContext,
event AvailableEvent = {}) {
impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
MemObject, SyclContext, AvailableEvent);
}

Expand All @@ -124,26 +129,27 @@ class buffer {

/* -- property interface members -- */

range<dimensions> get_range() const { return impl->get_range(); }
range<dimensions> get_range() const { return Range; }

size_t get_count() const { return impl->get_count(); }
size_t get_count() const { return Range.size(); }

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>
accessor<T, dimensions, mode, target, access::placeholder::false_t>
get_access(handler &commandGroupHandler) {
return impl->template get_access<mode, target>(*this, commandGroupHandler);
return impl->template get_access<T, dimensions, mode, target>(
*this, commandGroupHandler);
}

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

// template <access::mode mode, access::target target =
Expand Down Expand Up @@ -171,16 +177,29 @@ class buffer {

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

// template <typename ReinterpretT, int ReinterpretDim>
// buffer<ReinterpretT, ReinterpretDim, AllocatorT>
// reinterpret(range<ReinterpretDim> reinterpretRange) const {
// return impl->reinterpret((reinterpretRange));
// }
template <typename ReinterpretT, int ReinterpretDim>
buffer<ReinterpretT, ReinterpretDim, AllocatorT>
reinterpret(range<ReinterpretDim> reinterpretRange) const {
if (sizeof(ReinterpretT) * reinterpretRange.size() != get_size())
throw cl::sycl::invalid_object_error(
"Total size in bytes represented by the type and range of the "
"reinterpreted SYCL buffer does not equal the total size in bytes "
"represented by the type and range of this SYCL buffer");
return buffer<ReinterpretT, ReinterpretDim, AllocatorT>(impl,
reinterpretRange);
}

private:
shared_ptr_class<detail::buffer_impl<T, dimensions, AllocatorT>> impl;
shared_ptr_class<detail::buffer_impl<AllocatorT>> impl;
template <class Obj>
friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
template <typename A, int dims, typename C> friend class buffer;
range<dimensions> Range;

// Reinterpret contructor
buffer(shared_ptr_class<detail::buffer_impl<AllocatorT>> Impl,
range<dimensions> reinterpretRange)
: impl(Impl), Range(reinterpretRange){};
};
} // namespace sycl
} // namespace cl
Expand All @@ -190,8 +209,7 @@ template <typename T, int dimensions, typename AllocatorT>
struct hash<cl::sycl::buffer<T, dimensions, AllocatorT>> {
size_t
operator()(const cl::sycl::buffer<T, dimensions, AllocatorT> &b) const {
return hash<std::shared_ptr<
cl::sycl::detail::buffer_impl<T, dimensions, AllocatorT>>>()(
return hash<std::shared_ptr<cl::sycl::detail::buffer_impl<AllocatorT>>>()(
cl::sycl::detail::getSyclObjImpl(b));
}
};
Expand Down
Loading