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.

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

[SYCL] Implement basic sub-buffers support #64

Merged
merged 1 commit into from
Apr 5, 2019
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
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) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What about writing modern C++ with some { } instead of () for most of the instance constructions?

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;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think you can replace all the const something in this code by constexpr something

const int N = 7;
int Result[M][N] = {0};
{
auto OrigRange = range<2>(M, N);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What about using modern C++ { } everywhere instead of old troublesome ()?
Specially if people look at tests as good coding example and if we consider that SYCL is about modern C++ and heterogeneous computing... :-)

buffer<int, 2> Buffer(OrigRange);
Buffer.set_final_data((int *)Result);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Old C cast in modern C++ code...

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;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not clear why a size_t. Actually a bool looks fine to me.

auto Expected = !(i == 0) || (i == M - 1) || (j == 0) || (j == N - 1));

Replacing auto by int is also possible if you think it would cause the reader more time to understand the next line...

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);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Shorter:

range<1> OrigRange { M };

buffer<int, 1> Buffer(Data, OrigRange);
auto Offset = id<1>(1);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Idem for the 3 next ones

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;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

       int Expected = i > 1 && i < M - 2;

if (Acc[i] != Expected) {
std::cout << "line: " << __LINE__ << " Data[" << i << "] is " << Acc[i]
<< " expected " << Expected << std::endl;
Failed = true;
}
}
}
return Failed;
}