24 template <
int dimensions>
class range;
28 #ifdef SYCL2020_CONFORMANT_APIS
29 template <
typename DataT>
40 #ifdef SYCL2020_CONFORMANT_APIS
41 template <
typename DataT>
47 template <
typename T,
int Dimensions,
typename AllocatorT>
50 bool OwnNativeHandle =
true) {
59 buffer<DataT, Dimensions, Allocator, void>>;
62 typename AllocatorT = detail::default_buffer_allocator<DataT>>
74 template <
typename T,
int dimensions = 1,
75 typename AllocatorT = detail::default_buffer_allocator<T>,
81 static_assert(!std::is_same<T, std::string>::value,
82 "'std::string' is not a device copyable type");
92 template <
class Container>
96 decltype(std::declval<Container>().data())> (*)[],
97 const T (*)[]>::value>,
98 decltype(std::declval<Container>().size())>;
101 std::is_convertible<typename std::iterator_traits<It>::iterator_category,
102 std::input_iterator_tag>::value>;
103 template <
typename ItA,
typename ItB>
105 std::is_same<ItA, ItB>::value && !std::is_const<ItA>::value, ItA>;
116 : Range(bufferRange) {
117 impl = std::make_shared<detail::buffer_impl>(
120 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
nullptr,
121 (
const void *)
typeid(T).name(), dimensions,
122 sizeof(T), rangeToArray(Range).data());
128 : Range(bufferRange) {
129 impl = std::make_shared<detail::buffer_impl>(
133 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
nullptr,
134 (
const void *)
typeid(T).name(), dimensions,
135 sizeof(T), rangeToArray(Range).data());
141 : Range(bufferRange) {
142 impl = std::make_shared<detail::buffer_impl>(
146 impl->constructorNotification(CodeLoc, (
void *)impl.get(), hostData,
147 (
const void *)
typeid(T).name(), dimensions,
148 sizeof(T), rangeToArray(Range).data());
154 : Range(bufferRange) {
155 impl = std::make_shared<detail::buffer_impl>(
160 impl->constructorNotification(CodeLoc, (
void *)impl.get(), hostData,
161 (
const void *)
typeid(T).name(), dimensions,
162 sizeof(T), rangeToArray(Range).data());
165 template <
typename _T = T>
170 : Range(bufferRange) {
171 impl = std::make_shared<detail::buffer_impl>(
175 impl->constructorNotification(CodeLoc, (
void *)impl.get(), hostData,
176 (
const void *)
typeid(T).name(), dimensions,
177 sizeof(T), rangeToArray(Range).data());
180 template <
typename _T = T>
185 : Range(bufferRange) {
186 impl = std::make_shared<detail::buffer_impl>(
191 impl->constructorNotification(CodeLoc, (
void *)impl.get(), hostData,
192 (
const void *)
typeid(T).name(), dimensions,
193 sizeof(T), rangeToArray(Range).data());
196 buffer(
const std::shared_ptr<T> &hostData,
200 : Range(bufferRange) {
201 impl = std::make_shared<detail::buffer_impl>(
206 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
207 (
void *)hostData.get(),
208 (
const void *)
typeid(T).name(), dimensions,
209 sizeof(T), rangeToArray(Range).data());
212 buffer(
const std::shared_ptr<T[]> &hostData,
216 : Range(bufferRange) {
217 impl = std::make_shared<detail::buffer_impl>(
222 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
223 (
void *)hostData.get(),
224 (
const void *)
typeid(T).name(), dimensions,
225 sizeof(T), rangeToArray(Range).data());
228 buffer(
const std::shared_ptr<T> &hostData,
232 : Range(bufferRange) {
233 impl = std::make_shared<detail::buffer_impl>(
237 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
238 (
void *)hostData.get(),
239 (
const void *)
typeid(T).name(), dimensions,
240 sizeof(T), rangeToArray(Range).data());
243 buffer(
const std::shared_ptr<T[]> &hostData,
247 : Range(bufferRange) {
248 impl = std::make_shared<detail::buffer_impl>(
252 impl->constructorNotification(CodeLoc, (
void *)impl.get(),
253 (
void *)hostData.get(),
254 (
const void *)
typeid(T).name(), dimensions,
255 sizeof(T), rangeToArray(Range).data());
258 template <
class InputIterator,
int N = dimensions,
259 typename = EnableIfOneDimension<N>,
260 typename = EnableIfItInputIterator<InputIterator>>
261 buffer(InputIterator first, InputIterator last, AllocatorT allocator,
265 impl = std::make_shared<detail::buffer_impl>(
270 size_t r[3] = {Range[0], 0, 0};
271 impl->constructorNotification(CodeLoc, (
void *)impl.get(), &first,
272 (
const void *)
typeid(T).name(), dimensions,
276 template <
class InputIterator,
int N = dimensions,
277 typename = EnableIfOneDimension<N>,
278 typename = EnableIfItInputIterator<InputIterator>>
279 buffer(InputIterator first, InputIterator last,
283 impl = std::make_shared<detail::buffer_impl>(
287 size_t r[3] = {Range[0], 0, 0};
288 impl->constructorNotification(CodeLoc, (
void *)impl.get(), &first,
289 (
const void *)
typeid(T).name(), dimensions,
294 template <
class Container,
int N = dimensions,
295 typename = EnableIfOneDimension<N>,
296 typename = EnableIfContiguous<Container>>
297 buffer(Container &container, AllocatorT allocator,
300 : Range(
range<1>(container.size())) {
301 impl = std::make_shared<detail::buffer_impl>(
302 container.data(), size() *
sizeof(T),
306 size_t r[3] = {Range[0], 0, 0};
307 impl->constructorNotification(CodeLoc, (
void *)impl.get(), container.data(),
308 (
const void *)
typeid(T).name(), dimensions,
313 template <
class Container,
int N = dimensions,
314 typename = EnableIfOneDimension<N>,
315 typename = EnableIfContiguous<Container>>
318 :
buffer(container, {}, propList, CodeLoc) {}
323 : impl(b.impl), Range(subRange),
324 OffsetInBytes(getOffsetInBytes<T>(baseIndex, b.Range)),
326 impl->constructorNotification(CodeLoc, (
void *)impl.get(), impl.get(),
327 (
const void *)
typeid(T).name(), dimensions,
328 sizeof(T), rangeToArray(Range).data());
330 if (b.is_sub_buffer())
331 throw cl::sycl::invalid_object_error(
332 "Cannot create sub buffer from sub buffer.", PI_ERROR_INVALID_VALUE);
333 if (isOutOfBounds(baseIndex, subRange, b.Range))
334 throw cl::sycl::invalid_object_error(
335 "Requested sub-buffer size exceeds the size of the parent buffer",
336 PI_ERROR_INVALID_VALUE);
337 if (!isContiguousRegion(baseIndex, subRange, b.Range))
338 throw cl::sycl::invalid_object_error(
339 "Requested sub-buffer region is not contiguous",
340 PI_ERROR_INVALID_VALUE);
343 #ifdef __SYCL_INTERNAL_API
344 template <
int N = dimensions,
typename = EnableIfOneDimension<N>>
346 event AvailableEvent = {},
347 const detail::code_location CodeLoc = detail::code_location::current())
350 impl = std::make_shared<detail::buffer_impl>(
351 detail::pi::cast<pi_native_handle>(MemObject), SyclContext,
353 true, AvailableEvent);
354 Range[0] = impl->getSize() /
sizeof(
T);
355 impl->constructorNotification(CodeLoc, (
void *)impl.get(), &MemObject,
356 (
const void *)
typeid(
T).name(), dimensions,
357 sizeof(
T), rangeToArray(Range).data());
363 : impl(rhs.impl), Range(rhs.Range), OffsetInBytes(rhs.OffsetInBytes),
364 IsSubBuffer(rhs.IsSubBuffer) {
365 impl->constructorNotification(CodeLoc, (
void *)impl.get(), impl.get(),
366 (
const void *)
typeid(T).name(), dimensions,
367 sizeof(T), rangeToArray(Range).data());
372 : impl(
std::move(rhs.impl)), Range(rhs.Range),
373 OffsetInBytes(rhs.OffsetInBytes), IsSubBuffer(rhs.IsSubBuffer) {
374 impl->constructorNotification(CodeLoc, (
void *)impl.get(), impl.get(),
375 (
const void *)
typeid(T).name(), dimensions,
376 sizeof(T), rangeToArray(Range).data());
396 size_t get_count()
const {
return size(); }
397 size_t size() const noexcept {
return Range.size(); }
400 "get_size() is deprecated, please use byte_size() instead")
401 size_t get_size()
const {
return byte_size(); }
402 size_t byte_size() const noexcept {
return size() *
sizeof(T); }
405 return impl->template get_allocator<AllocatorT>();
408 template <access::mode Mode, access::target Target = access::target::device>
409 accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
414 return accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
416 *
this, CommandGroupHandler, {}, CodeLoc);
419 template <access::mode mode>
420 accessor<T, dimensions,
mode, access::target::host_buffer,
424 return accessor<T, dimensions,
mode, access::target::host_buffer,
425 access::placeholder::false_t,
429 template <access::mode mode, access::target target = access::target::device>
436 if (isOutOfBounds(accessOffset, accessRange, this->Range))
437 throw cl::sycl::invalid_object_error(
438 "Requested accessor would exceed the bounds of the buffer",
439 PI_ERROR_INVALID_VALUE);
443 *
this, commandGroupHandler, accessRange, accessOffset, {}, CodeLoc);
446 template <access::mode mode>
447 accessor<
T, dimensions,
mode, access::target::host_buffer,
448 access::placeholder::false_t, ext::oneapi::accessor_property_list<>>
452 if (isOutOfBounds(accessOffset, accessRange, this->Range))
453 throw cl::sycl::invalid_object_error(
454 "Requested accessor would exceed the bounds of the buffer",
455 PI_ERROR_INVALID_VALUE);
457 return accessor<T, dimensions,
mode, access::target::host_buffer,
458 access::placeholder::false_t,
460 *
this, accessRange, accessOffset, {}, CodeLoc);
463 #if __cplusplus >= 201703L
465 template <
typename... Ts>
auto get_access(Ts... args) {
466 return accessor{*
this, args...};
469 template <
typename... Ts>
470 auto get_access(handler &commandGroupHandler, Ts... args) {
471 return accessor{*
this, commandGroupHandler, args...};
474 template <
typename... Ts>
auto get_host_access(Ts... args) {
475 return host_accessor{*
this, args...};
478 template <
typename... Ts>
479 auto get_host_access(handler &commandGroupHandler, Ts... args) {
480 return host_accessor{*
this, commandGroupHandler, args...};
485 template <
typename Destination = std::
nullptr_t>
487 impl->set_final_data(finalData);
494 template <
typename Re
interpretT,
int Re
interpretDim>
497 if (
sizeof(ReinterpretT) * reinterpretRange.
size() != byte_size())
498 throw cl::sycl::invalid_object_error(
499 "Total size in bytes represented by the type and range of the "
500 "reinterpreted SYCL buffer does not equal the total size in bytes "
501 "represented by the type and range of this SYCL buffer",
502 PI_ERROR_INVALID_VALUE);
505 impl, reinterpretRange, OffsetInBytes, IsSubBuffer);
508 template <
typename Re
interpretT,
int Re
interpretDim = dimensions>
509 typename std::enable_if<
510 (
sizeof(ReinterpretT) ==
sizeof(T)) && (dimensions == ReinterpretDim),
514 impl, get_range(), OffsetInBytes, IsSubBuffer);
517 template <
typename Re
interpretT,
int Re
interpretDim = dimensions>
518 typename std::enable_if<
519 (ReinterpretDim == 1) && ((dimensions != ReinterpretDim) ||
520 (
sizeof(ReinterpretT) !=
sizeof(T))),
523 long sz = byte_size();
524 if (sz %
sizeof(ReinterpretT) != 0)
525 throw cl::sycl::invalid_object_error(
526 "Total byte size of buffer is not evenly divisible by the size of "
527 "the reinterpreted type",
528 PI_ERROR_INVALID_VALUE);
531 impl,
range<1>{sz /
sizeof(ReinterpretT)}, OffsetInBytes, IsSubBuffer);
535 return impl->template has_property<propertyT>();
539 return impl->template get_property<propertyT>();
546 bool outOfBounds =
false;
547 for (
int i = 0; i < dimensions; ++i)
548 outOfBounds |= newRange[i] + offset[i] > parentRange[i];
554 std::shared_ptr<detail::buffer_impl> impl;
557 template <
typename A,
int dims,
typename C,
typename Enable>
562 template <
typename HT,
int HDims,
typename HAllocT>
568 size_t OffsetInBytes = 0;
569 bool IsSubBuffer =
false;
572 template <
int N = dimensions,
typename = EnableIfOneDimension<N>>
574 bool OwnNativeHandle,
event AvailableEvent = {},
578 impl = std::make_shared<detail::buffer_impl>(
579 MemObject, SyclContext,
581 OwnNativeHandle, AvailableEvent);
582 Range[0] = impl->getSize() /
sizeof(T);
583 impl->constructorNotification(CodeLoc, (
void *)impl.get(), &MemObject,
584 (
const void *)
typeid(T).name(), dimensions,
585 sizeof(T), rangeToArray(Range).data());
589 buffer(std::shared_ptr<detail::buffer_impl> Impl,
590 range<dimensions> reinterpretRange,
size_t reinterpretOffset,
592 const detail::code_location CodeLoc = detail::code_location::current())
593 : impl(Impl), Range(reinterpretRange), OffsetInBytes(reinterpretOffset),
594 IsSubBuffer(isSubBuffer) {
595 impl->constructorNotification(CodeLoc, (
void *)impl.get(), Impl.get(),
596 (
const void *)
typeid(
T).name(), dimensions,
597 sizeof(
T), rangeToArray(Range).data());
600 template <
typename Type,
int N>
601 size_t getOffsetInBytes(
const id<N> &offset,
const range<N> &range) {
605 bool isContiguousRegion(
const id<1> &,
const range<1> &,
const range<1> &) {
610 bool isContiguousRegion(
const id<2> &offset,
const range<2> &newRange,
611 const range<2> &parentRange) {
619 return newRange[0] == 1;
620 return newRange[1] == parentRange[1];
623 bool isContiguousRegion(
const id<3> &offset,
const range<3> &newRange,
624 const range<3> &parentRange) {
634 return newRange[0] == 1 && newRange[1] == 1;
636 return newRange[0] == 1 && newRange[2] == parentRange[2];
637 return newRange[1] == parentRange[1] && newRange[2] == parentRange[2];
643 const buffer<DataT, Dimensions, Allocator, void> &Obj)
645 buffer<DataT, Dimensions, Allocator, void>>;
647 template <backend BackendName>
648 backend_return_t<BackendName, buffer<T, dimensions, AllocatorT>>
650 auto NativeHandles = impl->getNativeVector(BackendName);
651 return detail::BufferInterop<BackendName,
T, dimensions,
652 AllocatorT>::GetNativeObjs(NativeHandles);
656 #ifdef __cpp_deduction_guides
657 template <
class InputIterator,
class AllocatorT>
658 buffer(InputIterator, InputIterator, AllocatorT,
const property_list & = {})
659 -> buffer<
typename std::iterator_traits<InputIterator>::value_type, 1,
661 template <
class InputIterator>
662 buffer(InputIterator, InputIterator,
const property_list & = {})
663 -> buffer<
typename std::iterator_traits<InputIterator>::value_type, 1>;
664 template <
class Container,
class AllocatorT>
665 buffer(Container &, AllocatorT,
const property_list & = {})
666 -> buffer<typename Container::value_type, 1, AllocatorT>;
667 template <
class Container>
668 buffer(Container &,
const property_list & = {})
669 -> buffer<typename Container::value_type, 1>;
670 template <
class T,
int dimensions,
class AllocatorT>
671 buffer(
const T *,
const range<dimensions> &, AllocatorT,
672 const property_list & = {}) -> buffer<T, dimensions, AllocatorT>;
673 template <
class T,
int dimensions>
674 buffer(
const T *,
const range<dimensions> &,
const property_list & = {})
675 -> buffer<T, dimensions>;
676 #endif // __cpp_deduction_guides
682 template <
typename T,
int dimensions,
typename AllocatorT>
683 struct hash<
cl::sycl::buffer<T, dimensions, AllocatorT>> {
686 return hash<std::shared_ptr<cl::sycl::detail::buffer_impl>>()(