34 #include <type_traits>
37 inline namespace _V1 {
43 #ifdef __SYCL_DEVICE_ONLY__
44 constexpr uint32_t flags =
45 static_cast<uint32_t
>(
50 #endif // __SYCL_DEVICE_ONLY__
61 template <
typename T,
int Dimensions = 1>
62 class __SYCL_TYPE(private_memory) private_memory {
66 #ifndef __SYCL_DEVICE_ONLY__
69 Val.reset(
new T[G.get_local_range().size()]);
70 #endif // __SYCL_DEVICE_ONLY__
76 #ifndef __SYCL_DEVICE_ONLY__
80 return Val.get()[Ind];
84 #endif // __SYCL_DEVICE_ONLY__
88 #ifdef __SYCL_DEVICE_ONLY__
95 std::unique_ptr<T[]> Val;
96 #endif // #ifdef __SYCL_DEVICE_ONLY__
103 template <
int Dimensions = 1>
class __SYCL_TYPE(group)
group {
105 #ifndef __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
108 using linear_id_type = size_t;
110 #endif // __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
113 sycl::memory_scope::work_group;
121 size_t get_id(
int dimension)
const {
return index[dimension]; }
125 size_t get_group_id(
int dimension)
const {
return index[dimension]; }
128 "sycl::group::get_max_local_range() instead")
131 size_t get_global_range(
int dimension)
const {
132 return globalRange[dimension];
136 #ifdef __SYCL_DEVICE_ONLY__
137 return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
140 "get_local_id() is not implemented on host");
144 size_t get_local_id(
int dimention)
const {
return get_local_id()[dimention]; }
147 return get_local_linear_id_impl<Dimensions>();
152 size_t get_local_range(
int dimension)
const {
return localRange[dimension]; }
155 return get_local_linear_range_impl();
160 size_t get_group_range(
int dimension)
const {
161 return get_group_range()[dimension];
164 size_t get_group_linear_range()
const {
165 return get_group_linear_range_impl();
170 size_t operator[](
int dimension)
const {
return index[dimension]; }
173 size_t get_linear_id()
const {
return get_group_linear_id(); }
175 size_t get_group_linear_id()
const {
return get_group_linear_id_impl(); }
179 template <
typename WorkItemFunctionT>
180 void parallel_for_work_item(WorkItemFunctionT Func)
const {
184 #ifdef __SYCL_DEVICE_ONLY__
186 __spirv::initGlobalSize<Dimensions, range<Dimensions>>()};
188 __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>()};
190 __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>()};
192 __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>()};
199 detail::Builder::createItem<Dimensions, false>(GlobalSize, GlobalId);
201 detail::Builder::createItem<Dimensions, false>(LocalSize, LocalId);
203 detail::Builder::createHItem<Dimensions>(GlobalItem, LocalItem);
213 detail::Builder::createItem<Dimensions, false>(
214 globalRange, GroupStartID + LocalID);
216 detail::Builder::createItem<Dimensions, false>(localRange,
219 detail::Builder::createHItem<Dimensions>(GlobalItem, LocalItem);
222 #endif // __SYCL_DEVICE_ONLY__
230 template <
typename WorkItemFunctionT>
232 WorkItemFunctionT Func)
const {
234 #ifdef __SYCL_DEVICE_ONLY__
236 __spirv::initGlobalSize<Dimensions, range<Dimensions>>()};
238 __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>()};
240 __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>()};
242 __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>()};
245 detail::Builder::createItem<Dimensions, false>(GlobalSize, GlobalId);
247 detail::Builder::createItem<Dimensions, false>(LocalSize, LocalId);
249 GlobalItem, LocalItem, flexibleRange);
255 LocalId, LocalSize, flexibleRange,
266 detail::Builder::createItem<Dimensions, false>(
267 globalRange, GroupStartID + LocalID);
269 detail::Builder::createItem<Dimensions, false>(localRange,
272 GlobalItem, LocalItem, flexibleRange);
275 LocalID, localRange, flexibleRange,
281 #endif // __SYCL_DEVICE_ONLY__
287 template <access::mode accessMode = access::mode::read_write>
311 template <
typename dataT>
313 std::enable_if_t<!detail::is_bool<dataT>::value,
317 size_t srcStride)
const {
323 numElements, srcStride, 0);
332 template <
typename dataT>
334 std::enable_if_t<!detail::is_bool<dataT>::value,
345 numElements, destStride, 0);
355 template <
typename DestDataT,
typename SrcDataT>
356 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
357 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
361 size_t srcStride)
const {
367 numElements, srcStride, 0);
377 template <
typename DestDataT,
typename SrcDataT>
378 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
379 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
383 size_t destStride)
const {
389 numElements, destStride, 0);
398 template <
typename T, access::address_space DestS, access::address_space SrcS>
403 access::decorated::legacy>
406 access::decorated::legacy>
409 size_t Stride)
const {
410 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
411 "Async copy to/from bool memory is not supported.");
413 reinterpret_cast<uint8_t *
>(Dest.get()));
415 reinterpret_cast<uint8_t *
>(Src.get()));
416 return async_work_group_copy(DestP, SrcP, NumElements, Stride);
424 template <
typename T, access::address_space DestS, access::address_space SrcS>
429 access::decorated::legacy>
432 access::decorated::legacy>
435 size_t Stride)
const {
436 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
437 "Async copy to/from bool memory is not supported.");
439 auto DestP = address_space_cast<DestS, access::decorated::legacy>(
440 reinterpret_cast<VecT *
>(Dest.get()));
441 auto SrcP = address_space_cast<SrcS, access::decorated::legacy>(
442 reinterpret_cast<VecT *
>(Src.get()));
443 return async_work_group_copy(DestP, SrcP, NumElements, Stride);
453 std::enable_if_t<detail::is_scalar_bool<DestT>::value &&
454 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
458 size_t NumElements,
size_t Stride)
const {
459 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
460 "Async copy to/from bool memory is not supported.");
462 std::conditional_t<std::is_const_v<SrcT>,
const uint8_t, uint8_t>;
471 return async_work_group_copy(DestP, SrcP, NumElements, Stride);
481 std::enable_if_t<detail::is_vector_bool<DestT>::value &&
482 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
486 size_t NumElements,
size_t Stride)
const {
487 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
488 "Async copy to/from bool memory is not supported.");
491 std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
500 return async_work_group_copy(DestP, SrcP, NumElements, Stride);
508 template <
typename dataT>
512 size_t numElements)
const {
513 return async_work_group_copy(dest, src, numElements, 1);
521 template <
typename dataT>
525 size_t numElements)
const {
526 return async_work_group_copy(dest, src, numElements, 1);
535 template <
typename DestDataT,
typename SrcDataT>
536 typename std::enable_if_t<
537 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
540 size_t numElements)
const {
541 return async_work_group_copy(dest, src, numElements, 1);
550 template <
typename DestDataT,
typename SrcDataT>
551 typename std::enable_if_t<
552 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
555 size_t numElements)
const {
556 return async_work_group_copy(dest, src, numElements, 1);
559 template <
typename... eventTN>
void wait_for(eventTN... Events)
const {
560 waitForHelper(Events...);
564 bool Result = (rhs.globalRange == globalRange) &&
565 (rhs.localRange == localRange) && (rhs.index == index);
567 "inconsistent group class fields");
572 return !((*this) == rhs);
581 template <
int dims = Dimensions>
582 typename std::enable_if_t<(dims == 1),
size_t>
583 get_local_linear_id_impl()
const {
588 template <
int dims = Dimensions>
589 typename std::enable_if_t<(dims == 2),
size_t>
590 get_local_linear_id_impl()
const {
592 return localId[0] * localRange[1] + localId[1];
595 template <
int dims = Dimensions>
596 typename std::enable_if_t<(dims == 3),
size_t>
597 get_local_linear_id_impl()
const {
599 return (localId[0] * localRange[1] * localRange[2]) +
600 (localId[1] * localRange[2]) + localId[2];
603 template <
int dims = Dimensions>
604 typename std::enable_if_t<(dims == 1),
size_t>
605 get_local_linear_range_impl()
const {
606 auto localRange = get_local_range();
607 return localRange[0];
610 template <
int dims = Dimensions>
611 typename std::enable_if_t<(dims == 2),
size_t>
612 get_local_linear_range_impl()
const {
613 auto localRange = get_local_range();
614 return localRange[0] * localRange[1];
617 template <
int dims = Dimensions>
618 typename std::enable_if_t<(dims == 3),
size_t>
619 get_local_linear_range_impl()
const {
620 auto localRange = get_local_range();
621 return localRange[0] * localRange[1] * localRange[2];
624 template <
int dims = Dimensions>
625 typename std::enable_if_t<(dims == 1),
size_t>
626 get_group_linear_range_impl()
const {
627 auto groupRange = get_group_range();
628 return groupRange[0];
631 template <
int dims = Dimensions>
632 typename std::enable_if_t<(dims == 2),
size_t>
633 get_group_linear_range_impl()
const {
634 auto groupRange = get_group_range();
635 return groupRange[0] * groupRange[1];
638 template <
int dims = Dimensions>
639 typename std::enable_if_t<(dims == 3),
size_t>
640 get_group_linear_range_impl()
const {
641 auto groupRange = get_group_range();
642 return groupRange[0] * groupRange[1] * groupRange[2];
645 template <
int dims = Dimensions>
646 typename std::enable_if_t<(dims == 1),
size_t>
647 get_group_linear_id_impl()
const {
651 template <
int dims = Dimensions>
652 typename std::enable_if_t<(dims == 2),
size_t>
653 get_group_linear_id_impl()
const {
654 return index[0] * groupRange[1] + index[1];
667 template <
int dims = Dimensions>
668 typename std::enable_if_t<(dims == 3),
size_t>
669 get_group_linear_id_impl()
const {
670 return (index[0] * groupRange[1] * groupRange[2]) +
671 (index[1] * groupRange[2]) + index[2];
674 void waitForHelper()
const {}
678 template <
typename T,
typename... Ts>
679 void waitForHelper(T E, Ts... Es)
const {
681 waitForHelper(Es...);
688 : globalRange(G), localRange(L), groupRange(GroupRange), index(I) {
691 "global range is not multiple of local");
693 "inconsistent group constructor arguments");
700 #ifdef __SYCL_DEVICE_ONLY__
705 "Free function calls are not supported on host");
709 namespace ext::oneapi::experimental {
711 #ifdef __SYCL_DEVICE_ONLY__
712 return sycl::detail::Builder::getElement(
717 "Free function calls are not supported on host");