Skip to content

Commit 82fead6

Browse files
Fznamznonvladimirlaz
Fznamznon
authored and
vladimirlaz
committed
[SYCL] Implement basic sub-buffers support
Signed-off-by: Mariya Podchishchaeva <[email protected]>
1 parent 7e5a7aa commit 82fead6

File tree

3 files changed

+136
-24
lines changed

3 files changed

+136
-24
lines changed

‎sycl/include/CL/sycl/accessor.hpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -810,11 +810,11 @@ class accessor
810810
buffer<DataT, Dimensions>>::type &bufferRef)
811811
#ifdef __SYCL_DEVICE_ONLY__
812812
: impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
813-
bufferRef.get_range(), bufferRef.get_range()) {
813+
bufferRef.MemRange, bufferRef.MemRange) {
814814
#else
815815
: impl(std::make_shared<_ImplT>(
816816
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
817-
bufferRef.get_range(), bufferRef.get_range())) {
817+
bufferRef.MemRange, bufferRef.MemRange)) {
818818
#endif
819819
auto BufImpl = detail::getSyclObjImpl(bufferRef);
820820
if (AccessTarget == access::target::host_buffer) {
@@ -858,12 +858,12 @@ class accessor
858858
// Pass nullptr as a pointer to mem and use buffers from the ctor
859859
// arguments to avoid the need in adding utility functions for
860860
// dummy/default initialization of range fields.
861-
: impl(nullptr, bufferRef.get_range(), bufferRef.get_range(),
861+
: impl(nullptr, bufferRef.MemRange, bufferRef.MemRange,
862862
&commandGroupHandlerRef) {}
863863
#else
864864
: impl(std::make_shared<_ImplT>(
865865
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
866-
bufferRef.get_range(), bufferRef.get_range(),
866+
bufferRef.MemRange, bufferRef.MemRange,
867867
&commandGroupHandlerRef)) {
868868
auto BufImpl = detail::getSyclObjImpl(bufferRef);
869869
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
@@ -906,11 +906,11 @@ class accessor
906906
// arguments to avoid the need in adding utility functions for
907907
// dummy/default initialization of range<Dimensions> and
908908
// id<Dimension> fields.
909-
: impl(nullptr, Range, bufferRef.get_range(), Offset) {}
909+
: impl(nullptr, Range, bufferRef.MemRange, Offset) {}
910910
#else // !__SYCL_DEVICE_ONLY__
911911
: impl(std::make_shared<_ImplT>(
912912
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
913-
bufferRef.get_range(), Offset)) {
913+
bufferRef.MemRange, Offset)) {
914914
auto BufImpl = detail::getSyclObjImpl(bufferRef);
915915
if (AccessTarget == access::target::host_buffer) {
916916
if (BufImpl->OpenCLInterop) {
@@ -956,12 +956,12 @@ class accessor
956956
// arguments to avoid the need in adding utility functions for
957957
// dummy/default initialization of range<Dimensions> and
958958
// id<Dimension> fields.
959-
: impl(nullptr, Range, bufferRef.get_range(),
959+
: impl(nullptr, Range, bufferRef.MemRange,
960960
&commandGroupHandlerRef, Offset) {}
961961
#else // !__SYCL_DEVICE_ONLY__
962962
: impl(std::make_shared<_ImplT>(
963963
(dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
964-
bufferRef.get_range(), &commandGroupHandlerRef, Offset)) {
964+
bufferRef.MemRange, &commandGroupHandlerRef, Offset)) {
965965
auto BufImpl = detail::getSyclObjImpl(bufferRef);
966966
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
967967
throw cl::sycl::runtime_error(

‎sycl/include/CL/sycl/buffer.hpp

Lines changed: 39 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -33,54 +33,58 @@ class buffer {
3333

3434
buffer(const range<dimensions> &bufferRange,
3535
const property_list &propList = {})
36-
: Range(bufferRange) {
36+
: Range(bufferRange), MemRange(bufferRange) {
3737
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
3838
get_count() * sizeof(T), propList);
3939
}
4040

4141
buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
42-
const property_list &propList = {}) {
42+
const property_list &propList = {})
43+
: Range(bufferRange), MemRange(bufferRange) {
4344
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
4445
get_count() * sizeof(T), propList, allocator);
4546
}
4647

4748
buffer(T *hostData, const range<dimensions> &bufferRange,
4849
const property_list &propList = {})
49-
: Range(bufferRange) {
50+
: Range(bufferRange), MemRange(bufferRange) {
5051
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
5152
hostData, get_count() * sizeof(T), propList);
5253
}
5354

5455
buffer(T *hostData, const range<dimensions> &bufferRange,
55-
AllocatorT allocator, const property_list &propList = {}) {
56+
AllocatorT allocator, const property_list &propList = {})
57+
: Range(bufferRange), MemRange(bufferRange) {
5658
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
5759
hostData, get_count() * sizeof(T), propList, allocator);
5860
}
5961

6062
buffer(const T *hostData, const range<dimensions> &bufferRange,
6163
const property_list &propList = {})
62-
: Range(bufferRange) {
64+
: Range(bufferRange), MemRange(bufferRange) {
6365
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
6466
hostData, get_count() * sizeof(T), propList);
6567
}
6668

6769
buffer(const T *hostData, const range<dimensions> &bufferRange,
68-
AllocatorT allocator, const property_list &propList = {}) {
70+
AllocatorT allocator, const property_list &propList = {})
71+
: Range(bufferRange), MemRange(bufferRange) {
6972
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
7073
hostData, get_count() * sizeof(T), propList, allocator);
7174
}
7275

7376
buffer(const shared_ptr_class<T> &hostData,
7477
const range<dimensions> &bufferRange, AllocatorT allocator,
75-
const property_list &propList = {}) {
78+
const property_list &propList = {})
79+
: Range(bufferRange), MemRange(bufferRange) {
7680
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
7781
hostData, get_count() * sizeof(T), propList, allocator);
7882
}
7983

8084
buffer(const shared_ptr_class<T> &hostData,
8185
const range<dimensions> &bufferRange,
8286
const property_list &propList = {})
83-
: Range(bufferRange) {
87+
: Range(bufferRange), MemRange(bufferRange) {
8488
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
8589
hostData, get_count() * sizeof(T), propList);
8690
}
@@ -89,7 +93,8 @@ class buffer {
8993
typename = EnableIfOneDimension<N>>
9094
buffer(InputIterator first, InputIterator last, AllocatorT allocator,
9195
const property_list &propList = {})
92-
: Range(range<1>(std::distance(first, last))) {
96+
: Range(range<1>(std::distance(first, last))),
97+
MemRange(range<1>(std::distance(first, last))) {
9398
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
9499
first, last, get_count() * sizeof(T), propList, allocator);
95100
}
@@ -98,15 +103,16 @@ class buffer {
98103
typename = EnableIfOneDimension<N>>
99104
buffer(InputIterator first, InputIterator last,
100105
const property_list &propList = {})
101-
: Range(range<1>(std::distance(first, last))) {
106+
: Range(range<1>(std::distance(first, last))),
107+
MemRange(range<1>(std::distance(first, last))) {
102108
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
103109
first, last, get_count() * sizeof(T), propList);
104110
}
105111

106-
// buffer(buffer<T, dimensions, AllocatorT> b, const id<dimensions>
107-
// &baseIndex, const range<dimensions> &subRange) {
108-
// impl = std::make_shared<detail::buffer_impl>(b, baseIndex, subRange);
109-
// }
112+
buffer(buffer<T, dimensions, AllocatorT> &b, const id<dimensions> &baseIndex,
113+
const range<dimensions> &subRange)
114+
: impl(b.impl), Offset(baseIndex + b.Offset), Range(subRange), MemRange(b.MemRange),
115+
IsSubBuffer(true) {}
110116

111117
template <int N = dimensions, typename = EnableIfOneDimension<N>>
112118
buffer(cl_mem MemObject, const context &SyclContext,
@@ -116,6 +122,7 @@ class buffer {
116122
CHECK_OCL_CODE(clGetMemObjectInfo(MemObject, CL_MEM_SIZE, sizeof(size_t),
117123
&BufSize, nullptr));
118124
Range[0] = BufSize / sizeof(T);
125+
MemRange[0] = BufSize / sizeof(T);
119126
impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
120127
MemObject, SyclContext, BufSize, AvailableEvent);
121128
}
@@ -150,6 +157,9 @@ class buffer {
150157
access::target target = access::target::global_buffer>
151158
accessor<T, dimensions, mode, target, access::placeholder::false_t>
152159
get_access(handler &commandGroupHandler) {
160+
if (IsSubBuffer)
161+
return impl->template get_access<T, dimensions, mode, target>(
162+
*this, commandGroupHandler, Range, Offset);
153163
return impl->template get_access<T, dimensions, mode, target>(
154164
*this, commandGroupHandler);
155165
}
@@ -158,6 +168,9 @@ class buffer {
158168
accessor<T, dimensions, mode, access::target::host_buffer,
159169
access::placeholder::false_t>
160170
get_access() {
171+
if (IsSubBuffer)
172+
return impl->template get_access<T, dimensions, mode>(*this, Range,
173+
Offset);
161174
return impl->template get_access<T, dimensions, mode>(*this);
162175
}
163176

@@ -185,7 +198,7 @@ class buffer {
185198

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

188-
// bool is_sub_buffer() const { return impl->is_sub_buffer(); }
201+
bool is_sub_buffer() const { return IsSubBuffer; }
189202

190203
template <typename ReinterpretT, int ReinterpretDim>
191204
buffer<ReinterpretT, ReinterpretDim, AllocatorT>
@@ -212,12 +225,22 @@ class buffer {
212225
template <class Obj>
213226
friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
214227
template <typename A, int dims, typename C> friend class buffer;
228+
template <typename DataT, int dims, access::mode mode,
229+
access::target target, access::placeholder isPlaceholder>
230+
friend class accessor;
231+
// If this buffer is subbuffer - this range represents range of the parent
232+
// buffer
233+
range<dimensions> MemRange;
234+
bool IsSubBuffer = false;
215235
range<dimensions> Range;
236+
// If this buffer is sub-buffer - offset field specifies the origin of the
237+
// sub-buffer inside the parent buffer
238+
id<dimensions> Offset;
216239

217240
// Reinterpret contructor
218241
buffer(shared_ptr_class<detail::buffer_impl<AllocatorT>> Impl,
219242
range<dimensions> reinterpretRange)
220-
: impl(Impl), Range(reinterpretRange){};
243+
: impl(Impl), Range(reinterpretRange), MemRange(reinterpretRange) {};
221244
};
222245
} // namespace sycl
223246
} // namespace cl

‎sycl/test/basic_tests/buffer/subbuffer.cpp

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL
2+
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
6+
//==---------- subbuffer.cpp --- sub-buffer basic test ---------------------==//
7+
//
8+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
9+
// See https://llvm.org/LICENSE.txt for license information.
10+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
11+
//
12+
//===----------------------------------------------------------------------===//
13+
#include <CL/sycl.hpp>
14+
15+
using namespace cl::sycl;
16+
17+
int main() {
18+
19+
bool Failed = false;
20+
// Basic test case
21+
{
22+
const int M = 6;
23+
const int N = 7;
24+
int Result[M][N] = {0};
25+
{
26+
auto OrigRange = range<2>(M, N);
27+
buffer<int, 2> Buffer(OrigRange);
28+
Buffer.set_final_data((int *)Result);
29+
auto Offset = id<2>(1, 1);
30+
auto SubRange = range<2>(M - 2, N - 2);
31+
queue MyQueue;
32+
buffer<int, 2> SubBuffer(Buffer, Offset, SubRange);
33+
MyQueue.submit([&](handler &cgh) {
34+
auto B = SubBuffer.get_access<access::mode::read_write>(cgh);
35+
cgh.parallel_for<class Subbuf_test>(SubRange,
36+
[=](id<2> Index) { B[Index] = 1; });
37+
});
38+
}
39+
40+
// Check that we filled correct subset of buffer:
41+
// 0000000 0000000
42+
// 0000000 0111110
43+
// 0000000 --> 0111110
44+
// 0000000 0111110
45+
// 0000000 0111110
46+
// 0000000 0000000
47+
48+
for (size_t i = 0; i < M; ++i) {
49+
for (size_t j = 0; j < N; ++j) {
50+
size_t Expected =
51+
((i == 0) || (i == M - 1) || (j == 0) || (j == N - 1)) ? 0 : 1;
52+
if (Result[i][j] != Expected) {
53+
std::cout << "line: " << __LINE__ << " Result[" << i << "][" << j
54+
<< "] is " << Result[i][j] << " expected " << Expected
55+
<< std::endl;
56+
Failed = true;
57+
}
58+
}
59+
}
60+
}
61+
// Try to create subbuffer from subbuffer
62+
{
63+
const int M = 10;
64+
int Data[M] = {0};
65+
auto OrigRange = range<1>(M);
66+
buffer<int, 1> Buffer(Data, OrigRange);
67+
auto Offset = id<1>(1);
68+
auto SubRange = range<1>(M - 2);
69+
auto SubSubRange = range<1>(M - 4);
70+
queue MyQueue;
71+
buffer<int, 1> SubBuffer(Buffer, Offset, SubRange);
72+
buffer<int, 1> SubSubBuffer(SubBuffer, Offset, SubSubRange);
73+
MyQueue.submit([&](handler &cgh) {
74+
auto B = SubSubBuffer.get_access<access::mode::read_write>(cgh);
75+
cgh.parallel_for<class Subsubbuf_test>(SubSubRange,
76+
[=](id<1> Index) { B[Index] = 1; });
77+
});
78+
auto Acc = Buffer.get_access<cl::sycl::access::mode::read>();
79+
for (size_t i = 0; i < M; ++i) {
80+
size_t Expected = (i > 1 && i < M - 2) ? 1 : 0;
81+
if (Acc[i] != Expected) {
82+
std::cout << "line: " << __LINE__ << " Data[" << i << "] is " << Acc[i]
83+
<< " expected " << Expected << std::endl;
84+
Failed = true;
85+
}
86+
}
87+
}
88+
return Failed;
89+
}

0 commit comments

Comments
 (0)