Skip to content

Commit

Permalink
[SYCL] Implement basic sub-buffers support
Browse files Browse the repository at this point in the history
Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
  • Loading branch information
Fznamznon authored and vladimirlaz committed Apr 5, 2019
1 parent 7e5a7aa commit 82fead6
Show file tree
Hide file tree
Showing 3 changed files with 136 additions and 24 deletions.
16 changes: 8 additions & 8 deletions sycl/include/CL/sycl/accessor.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -810,11 +810,11 @@ class accessor
buffer<DataT, Dimensions>>::type &bufferRef)
#ifdef __SYCL_DEVICE_ONLY__
: impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range()) {
bufferRef.MemRange, bufferRef.MemRange) {
#else
: impl(std::make_shared<_ImplT>(
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range())) {
bufferRef.MemRange, bufferRef.MemRange)) {
#endif
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
Expand Down Expand Up @@ -858,12 +858,12 @@ class accessor
// Pass nullptr as a pointer to mem and use buffers from the ctor
// arguments to avoid the need in adding utility functions for
// dummy/default initialization of range fields.
: impl(nullptr, bufferRef.get_range(), bufferRef.get_range(),
: impl(nullptr, bufferRef.MemRange, bufferRef.MemRange,
&commandGroupHandlerRef) {}
#else
: impl(std::make_shared<_ImplT>(
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
bufferRef.get_range(), bufferRef.get_range(),
bufferRef.MemRange, bufferRef.MemRange,
&commandGroupHandlerRef)) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
Expand Down Expand Up @@ -906,11 +906,11 @@ class accessor
// arguments to avoid the need in adding utility functions for
// dummy/default initialization of range<Dimensions> and
// id<Dimension> fields.
: impl(nullptr, Range, bufferRef.get_range(), Offset) {}
: impl(nullptr, Range, bufferRef.MemRange, Offset) {}
#else // !__SYCL_DEVICE_ONLY__
: impl(std::make_shared<_ImplT>(
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
bufferRef.get_range(), Offset)) {
bufferRef.MemRange, Offset)) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
if (BufImpl->OpenCLInterop) {
Expand Down Expand Up @@ -956,12 +956,12 @@ class accessor
// arguments to avoid the need in adding utility functions for
// dummy/default initialization of range<Dimensions> and
// id<Dimension> fields.
: impl(nullptr, Range, bufferRef.get_range(),
: impl(nullptr, Range, bufferRef.MemRange,
&commandGroupHandlerRef, Offset) {}
#else // !__SYCL_DEVICE_ONLY__
: impl(std::make_shared<_ImplT>(
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
bufferRef.get_range(), &commandGroupHandlerRef, Offset)) {
bufferRef.MemRange, &commandGroupHandlerRef, Offset)) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
throw cl::sycl::runtime_error(
Expand Down
55 changes: 39 additions & 16 deletions sycl/include/CL/sycl/buffer.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -33,54 +33,58 @@ class buffer {

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

buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
const property_list &propList = {}) {
const property_list &propList = {})
: Range(bufferRange), MemRange(bufferRange) {
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 = {})
: Range(bufferRange) {
: Range(bufferRange), MemRange(bufferRange) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList);
}

buffer(T *hostData, const range<dimensions> &bufferRange,
AllocatorT allocator, const property_list &propList = {}) {
AllocatorT allocator, const property_list &propList = {})
: Range(bufferRange), MemRange(bufferRange) {
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 = {})
: Range(bufferRange) {
: Range(bufferRange), MemRange(bufferRange) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList);
}

buffer(const T *hostData, const range<dimensions> &bufferRange,
AllocatorT allocator, const property_list &propList = {}) {
AllocatorT allocator, const property_list &propList = {})
: Range(bufferRange), MemRange(bufferRange) {
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 = {}) {
const property_list &propList = {})
: Range(bufferRange), MemRange(bufferRange) {
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,
const property_list &propList = {})
: Range(bufferRange) {
: Range(bufferRange), MemRange(bufferRange) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
hostData, get_count() * sizeof(T), propList);
}
Expand All @@ -89,7 +93,8 @@ class buffer {
typename = EnableIfOneDimension<N>>
buffer(InputIterator first, InputIterator last, AllocatorT allocator,
const property_list &propList = {})
: Range(range<1>(std::distance(first, last))) {
: Range(range<1>(std::distance(first, last))),
MemRange(range<1>(std::distance(first, last))) {
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
first, last, get_count() * sizeof(T), propList, allocator);
}
Expand All @@ -98,15 +103,16 @@ class buffer {
typename = EnableIfOneDimension<N>>
buffer(InputIterator first, InputIterator last,
const property_list &propList = {})
: Range(range<1>(std::distance(first, last))) {
: Range(range<1>(std::distance(first, last))),
MemRange(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>
// &baseIndex, const range<dimensions> &subRange) {
// impl = std::make_shared<detail::buffer_impl>(b, baseIndex, subRange);
// }
buffer(buffer<T, dimensions, AllocatorT> &b, const id<dimensions> &baseIndex,
const range<dimensions> &subRange)
: impl(b.impl), Offset(baseIndex + b.Offset), Range(subRange), MemRange(b.MemRange),
IsSubBuffer(true) {}

template <int N = dimensions, typename = EnableIfOneDimension<N>>
buffer(cl_mem MemObject, const context &SyclContext,
Expand All @@ -116,6 +122,7 @@ class buffer {
CHECK_OCL_CODE(clGetMemObjectInfo(MemObject, CL_MEM_SIZE, sizeof(size_t),
&BufSize, nullptr));
Range[0] = BufSize / sizeof(T);
MemRange[0] = BufSize / sizeof(T);
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
MemObject, SyclContext, BufSize, AvailableEvent);
}
Expand Down Expand Up @@ -150,6 +157,9 @@ class buffer {
access::target target = access::target::global_buffer>
accessor<T, dimensions, mode, target, access::placeholder::false_t>
get_access(handler &commandGroupHandler) {
if (IsSubBuffer)
return impl->template get_access<T, dimensions, mode, target>(
*this, commandGroupHandler, Range, Offset);
return impl->template get_access<T, dimensions, mode, target>(
*this, commandGroupHandler);
}
Expand All @@ -158,6 +168,9 @@ class buffer {
accessor<T, dimensions, mode, access::target::host_buffer,
access::placeholder::false_t>
get_access() {
if (IsSubBuffer)
return impl->template get_access<T, dimensions, mode>(*this, Range,
Offset);
return impl->template get_access<T, dimensions, mode>(*this);
}

Expand Down Expand Up @@ -185,7 +198,7 @@ class buffer {

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

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

template <typename ReinterpretT, int ReinterpretDim>
buffer<ReinterpretT, ReinterpretDim, AllocatorT>
Expand All @@ -212,12 +225,22 @@ class buffer {
template <class Obj>
friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
template <typename A, int dims, typename C> friend class buffer;
template <typename DataT, int dims, access::mode mode,
access::target target, access::placeholder isPlaceholder>
friend class accessor;
// If this buffer is subbuffer - this range represents range of the parent
// buffer
range<dimensions> MemRange;
bool IsSubBuffer = false;
range<dimensions> Range;
// If this buffer is sub-buffer - offset field specifies the origin of the
// sub-buffer inside the parent buffer
id<dimensions> Offset;

// Reinterpret contructor
buffer(shared_ptr_class<detail::buffer_impl<AllocatorT>> Impl,
range<dimensions> reinterpretRange)
: impl(Impl), Range(reinterpretRange){};
: impl(Impl), Range(reinterpretRange), MemRange(reinterpretRange) {};
};
} // namespace sycl
} // namespace cl
Expand Down
89 changes: 89 additions & 0 deletions sycl/test/basic_tests/buffer/subbuffer.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,89 @@
// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out
//==---------- subbuffer.cpp --- sub-buffer basic test ---------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include <CL/sycl.hpp>

using namespace cl::sycl;

int main() {

bool Failed = false;
// Basic test case
{
const int M = 6;
const int N = 7;
int Result[M][N] = {0};
{
auto OrigRange = range<2>(M, N);
buffer<int, 2> Buffer(OrigRange);
Buffer.set_final_data((int *)Result);
auto Offset = id<2>(1, 1);
auto SubRange = range<2>(M - 2, N - 2);
queue MyQueue;
buffer<int, 2> SubBuffer(Buffer, Offset, SubRange);
MyQueue.submit([&](handler &cgh) {
auto B = SubBuffer.get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class Subbuf_test>(SubRange,
[=](id<2> Index) { B[Index] = 1; });
});
}

// Check that we filled correct subset of buffer:
// 0000000 0000000
// 0000000 0111110
// 0000000 --> 0111110
// 0000000 0111110
// 0000000 0111110
// 0000000 0000000

for (size_t i = 0; i < M; ++i) {
for (size_t j = 0; j < N; ++j) {
size_t Expected =
((i == 0) || (i == M - 1) || (j == 0) || (j == N - 1)) ? 0 : 1;
if (Result[i][j] != Expected) {
std::cout << "line: " << __LINE__ << " Result[" << i << "][" << j
<< "] is " << Result[i][j] << " expected " << Expected
<< std::endl;
Failed = true;
}
}
}
}
// Try to create subbuffer from subbuffer
{
const int M = 10;
int Data[M] = {0};
auto OrigRange = range<1>(M);
buffer<int, 1> Buffer(Data, OrigRange);
auto Offset = id<1>(1);
auto SubRange = range<1>(M - 2);
auto SubSubRange = range<1>(M - 4);
queue MyQueue;
buffer<int, 1> SubBuffer(Buffer, Offset, SubRange);
buffer<int, 1> SubSubBuffer(SubBuffer, Offset, SubSubRange);
MyQueue.submit([&](handler &cgh) {
auto B = SubSubBuffer.get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class Subsubbuf_test>(SubSubRange,
[=](id<1> Index) { B[Index] = 1; });
});
auto Acc = Buffer.get_access<cl::sycl::access::mode::read>();
for (size_t i = 0; i < M; ++i) {
size_t Expected = (i > 1 && i < M - 2) ? 1 : 0;
if (Acc[i] != Expected) {
std::cout << "line: " << __LINE__ << " Data[" << i << "] is " << Acc[i]
<< " expected " << Expected << std::endl;
Failed = true;
}
}
}
return Failed;
}

0 comments on commit 82fead6

Please sign in to comment.