From 369d9dd45a4b9fb6284ac6125cca9e857efa50e5 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 07:53:32 -0700 Subject: [PATCH 01/12] [SYCL] Align sub-group implementation and docs - Remove previously deprecated functions - Deprecate two-input shuffles - Fix linear_id type (size_t => uint32_t) Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 86 ++---------------------- 1 file changed, 6 insertions(+), 80 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index a8de9028ad557..777b2ee000b6f 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -126,7 +126,7 @@ struct sub_group { using id_type = id<1>; using range_type = range<1>; - using linear_id_type = size_t; + using linear_id_type = uint32_t; static constexpr int dimensions = 1; /* --- common interface members --- */ @@ -144,89 +144,11 @@ struct sub_group { unsigned int get_group_range() const { return __spirv_BuiltInNumSubgroups; } - unsigned int get_uniform_group_range() const { - return __spirv_BuiltInNumEnqueuedSubgroups; - } - - /* --- vote / ballot functions --- */ - - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::any_of instead.") - bool any(bool predicate) const { - return __spirv_GroupAny(__spv::Scope::Subgroup, predicate); - } - - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::all_of instead.") - bool all(bool predicate) const { - return __spirv_GroupAll(__spv::Scope::Subgroup, predicate); - } - template using EnableIfIsScalarArithmetic = sycl::detail::enable_if_t::value, T>; - /* --- collectives --- */ - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::broadcast instead.") - EnableIfIsScalarArithmetic broadcast(T x, id<1> local_id) const { - return sycl::detail::spirv::GroupBroadcast(x, local_id); - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::reduce instead.") - EnableIfIsScalarArithmetic reduce(T x, BinaryOperation op) const { - return sycl::detail::calc( - typename sycl::detail::GroupOpTag::type(), x, op); - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::reduce instead.") - EnableIfIsScalarArithmetic reduce(T x, T init, BinaryOperation op) const { - return op(init, reduce(x, op)); - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::exclusive_scan instead.") - EnableIfIsScalarArithmetic exclusive_scan(T x, BinaryOperation op) const { - return sycl::detail::calc( - typename sycl::detail::GroupOpTag::type(), x, op); - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::exclusive_scan instead.") - EnableIfIsScalarArithmetic exclusive_scan(T x, T init, - BinaryOperation op) const { - if (get_local_id().get(0) == 0) { - x = op(init, x); - } - T scan = exclusive_scan(x, op); - if (get_local_id().get(0) == 0) { - scan = init; - } - return scan; - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::inclusive_scan instead.") - EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op) const { - return sycl::detail::calc( - typename sycl::detail::GroupOpTag::type(), x, op); - } - - template - __SYCL_EXPORT_DEPRECATED("Use sycl::intel::inclusive_scan instead.") - EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op, - T init) const { - if (get_local_id().get(0) == 0) { - x = op(init, x); - } - return inclusive_scan(x, op); - } - /* --- one-input shuffles --- */ /* indices in [0 , sub_group size) */ @@ -249,17 +171,21 @@ struct sub_group { /* --- two-input shuffles --- */ /* indices in [0 , 2 * sub_group size) */ - template T shuffle(T x, T y, id<1> local_id) const { + template + __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") + T shuffle(T x, T y, id<1> local_id) const { return sycl::detail::sub_group::shuffle_down( x, y, (local_id - get_local_id()).get(0)); } template + __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_down(T current, T next, uint32_t delta) const { return sycl::detail::sub_group::shuffle_down(current, next, delta); } template + __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_up(T previous, T current, uint32_t delta) const { return sycl::detail::sub_group::shuffle_up(previous, current, delta); } From 9788887724aca20c7695071f2e5915f77c89496e Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 08:11:21 -0700 Subject: [PATCH 02/12] [SYCL] Combine device and host sub-group headers Preparation for updated sub-group features. Combining files to avoid accidental interface differences. Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 184 +++++++++++++++- sycl/include/CL/sycl/intel/sub_group_host.hpp | 196 ------------------ 2 files changed, 176 insertions(+), 204 deletions(-) delete mode 100644 sycl/include/CL/sycl/intel/sub_group_host.hpp diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 777b2ee000b6f..939bf3dc5b740 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -23,8 +23,6 @@ #include -#ifdef __SYCL_DEVICE_ONLY__ - __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { template class multi_ptr; @@ -33,6 +31,23 @@ namespace detail { namespace sub_group { +// Selects 8-bit, 16-bit or 32-bit type depending on size of T. If T doesn't +// maps to mentioned types, then void is returned +template +using SelectBlockT = + select_apply_cl_scalar_t; + +template +using AcceptableForGlobalLoadStore = + bool_constant>::value && + Space == access::address_space::global_space>; + +template +using AcceptableForLocalLoadStore = + bool_constant>::value && + Space == access::address_space::local_space>; + +#ifdef __SYCL_DEVICE_ONLY__ #define __SYCL_SG_GENERATE_BODY_1ARG(name, SPIRVOperation) \ template T name(T x, id<1> local_id) { \ using OCLT = sycl::detail::ConvertToOpenCLType_t; \ @@ -115,6 +130,7 @@ void store(multi_ptr dst, const vec &x) { __spirv_SubgroupBlockWriteINTEL(reinterpret_cast(dst.get()), sycl::detail::bit_cast(x)); } +#endif // __SYCL_DEVICE_ONLY__ } // namespace sub_group @@ -132,17 +148,49 @@ struct sub_group { /* --- common interface members --- */ id<1> get_local_id() const { +#ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupLocalInvocationId; +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range<1> get_local_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_BuiltInSubgroupSize; +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } - range<1> get_local_range() const { return __spirv_BuiltInSubgroupSize; } range<1> get_max_local_range() const { +#ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupMaxSize; +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } - id<1> get_group_id() const { return __spirv_BuiltInSubgroupId; } + id<1> get_group_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_BuiltInSubgroupId; +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } - unsigned int get_group_range() const { return __spirv_BuiltInNumSubgroups; } + unsigned int get_group_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_BuiltInNumSubgroups; +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } template using EnableIfIsScalarArithmetic = @@ -153,19 +201,47 @@ struct sub_group { /* indices in [0 , sub_group size) */ template T shuffle(T x, id<1> local_id) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle(x, local_id); +#else + (void)x; + (void)local_id; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template T shuffle_down(T x, uint32_t delta) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down(x, x, delta); +#else + (void)x; + (void)delta; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template T shuffle_up(T x, uint32_t delta) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_up(x, x, delta); +#else + (void)x; + (void)delta; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template T shuffle_xor(T x, id<1> value) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_xor(x, value); +#else + (void)x; + (void)value; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } /* --- two-input shuffles --- */ @@ -174,20 +250,44 @@ struct sub_group { template __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle(T x, T y, id<1> local_id) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down( x, y, (local_id - get_local_id()).get(0)); +#else + (void)x; + (void)y; + (void)local_id; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_down(T current, T next, uint32_t delta) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down(current, next, delta); +#else + (void)current; + (void)next; + (void)delta; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_up(T previous, T current, uint32_t delta) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_up(previous, current, delta); +#else + (void)previous; + (void)current; + (void)delta; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } /* --- sub_group load/stores --- */ @@ -197,14 +297,26 @@ struct sub_group { sycl::detail::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore::value, T> load(const multi_ptr src) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::load(src); +#else + (void)src; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template sycl::detail::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore::value, T> load(const multi_ptr src) const { +#ifdef __SYCL_DEVICE_ONLY__ return src.get()[get_local_id()[0]]; +#else + (void)src; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -213,7 +325,13 @@ struct sub_group { N != 1, vec> load(const multi_ptr src) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::load(src); +#else + (void)src; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -221,11 +339,17 @@ struct sub_group { sycl::detail::sub_group::AcceptableForLocalLoadStore::value, vec> load(const multi_ptr src) const { +#ifdef __SYCL_DEVICE_ONLY__ vec res; for (int i = 0; i < N; ++i) { res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]); } return res; +#else + (void)src; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -234,21 +358,41 @@ struct sub_group { N == 1, vec> load(const multi_ptr src) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::load(src); +#else + (void)src; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template sycl::detail::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore::value> store(multi_ptr dst, const T &x) const { +#ifdef __SYCL_DEVICE_ONLY__ sycl::detail::sub_group::store(dst, x); +#else + (void)dst; + (void)x; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template sycl::detail::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore::value> store(multi_ptr dst, const T &x) const { +#ifdef __SYCL_DEVICE_ONLY__ dst.get()[get_local_id()[0]] = x; +#else + (void)dst; + (void)x; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -256,7 +400,14 @@ struct sub_group { sycl::detail::sub_group::AcceptableForGlobalLoadStore::value && N == 1> store(multi_ptr dst, const vec &x) const { +#ifdef __SYCL_DEVICE_ONLY__ store(dst, x); +#else + (void)dst; + (void)x; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -264,24 +415,44 @@ struct sub_group { sycl::detail::sub_group::AcceptableForGlobalLoadStore::value && N != 1> store(multi_ptr dst, const vec &x) const { +#ifdef __SYCL_DEVICE_ONLY__ sycl::detail::sub_group::store(dst, x); +#else + (void)dst; + (void)x; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template sycl::detail::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore::value> store(multi_ptr dst, const vec &x) const { +#ifdef __SYCL_DEVICE_ONLY__ for (int i = 0; i < N; ++i) { *(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i]; } +#else + (void)dst; + (void)x; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } /* --- synchronization functions --- */ void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const { +#ifdef __SYCL_DEVICE_ONLY__ uint32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, flags); +#else + (void)accessSpace; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } protected: @@ -291,6 +462,3 @@ struct sub_group { } // namespace intel } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) -#else -#include -#endif diff --git a/sycl/include/CL/sycl/intel/sub_group_host.hpp b/sycl/include/CL/sycl/intel/sub_group_host.hpp deleted file mode 100644 index 0c5762462e1ff..0000000000000 --- a/sycl/include/CL/sycl/intel/sub_group_host.hpp +++ /dev/null @@ -1,196 +0,0 @@ -//==- sub_group_host.hpp --- SYCL sub-group for host device ---------------==// -// -// 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 -// -//===----------------------------------------------------------------------===// - -#pragma once - -#include -#include -#include -#include -#include -#ifndef __SYCL_DEVICE_ONLY__ - -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -template class multi_ptr; -namespace intel { -struct sub_group { - - using id_type = id<1>; - using range_type = range<1>; - using linear_id_type = size_t; - static constexpr int dimensions = 1; - - /* --- common interface members --- */ - - id<1> get_local_id() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - range<1> get_local_range() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - range<1> get_max_local_range() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - id<1> get_group_id() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - size_t get_group_range() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - size_t get_uniform_group_range() const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- vote / ballot functions --- */ - - bool any(bool) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - bool all(bool) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- collectives --- */ - - template T broadcast(T, id<1>) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T reduce(T, BinaryOperation) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T reduce(T, T, BinaryOperation) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T exclusive_scan(T, BinaryOperation) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T exclusive_scan(T, T, BinaryOperation) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T inclusive_scan(T, BinaryOperation) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - T inclusive_scan(T, BinaryOperation, T) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- one - input shuffles --- */ - /* indices in [0 , sub - group size ) */ - - template T shuffle(T, id<1>) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template T shuffle_down(T, uint32_t) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - template T shuffle_up(T, uint32_t) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template T shuffle_xor(T, id<1>) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- two - input shuffles --- */ - /* indices in [0 , 2* sub - group size ) */ - template T shuffle(T, T, id<1>) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - template T shuffle_down(T, T, uint32_t) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - template T shuffle_up(T, T, uint32_t) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- sub - group load / stores --- */ - /* these can map to SIMD or block read / write hardware where available */ - template - T load(const multi_ptr) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - vec load(const multi_ptr) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - void store(multi_ptr, const T &) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - template - void store(multi_ptr, const vec &) const { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - - /* --- synchronization functions --- */ - void barrier(access::fence_space accessSpace = - access::fence_space::global_and_local) const { - (void)accessSpace; - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } - -protected: - template friend class cl::sycl::nd_item; - sub_group() { - throw runtime_error("Subgroups are not supported on host device. ", - PI_INVALID_DEVICE); - } -}; -} // namespace intel -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) -#endif From 8ddc993f8d83bc1cb5d1260a65d652d0730aa5ac Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 08:33:04 -0700 Subject: [PATCH 03/12] [SYCL] Add linear ID queries to sub-group class Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 34 ++++++++++++++++++------ 1 file changed, 26 insertions(+), 8 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 939bf3dc5b740..d0a93e8e30544 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -147,7 +147,7 @@ struct sub_group { /* --- common interface members --- */ - id<1> get_local_id() const { + id_type get_local_id() const { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupLocalInvocationId; #else @@ -156,7 +156,16 @@ struct sub_group { #endif } - range<1> get_local_range() const { + linear_id_type get_local_linear_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(get_local_id()); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range_type get_local_range() const { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupSize; #else @@ -165,7 +174,7 @@ struct sub_group { #endif } - range<1> get_max_local_range() const { + range_type get_max_local_range() const { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupMaxSize; #else @@ -174,7 +183,7 @@ struct sub_group { #endif } - id<1> get_group_id() const { + id_type get_group_id() const { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupId; #else @@ -183,7 +192,16 @@ struct sub_group { #endif } - unsigned int get_group_range() const { + linear_id_type get_group_linear_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(get_group_id()); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range_type get_group_range() const { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInNumSubgroups; #else @@ -200,7 +218,7 @@ struct sub_group { /* --- one-input shuffles --- */ /* indices in [0 , sub_group size) */ - template T shuffle(T x, id<1> local_id) const { + template T shuffle(T x, id_type local_id) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle(x, local_id); #else @@ -233,7 +251,7 @@ struct sub_group { #endif } - template T shuffle_xor(T x, id<1> value) const { + template T shuffle_xor(T x, id_type value) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_xor(x, value); #else @@ -249,7 +267,7 @@ struct sub_group { template __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") - T shuffle(T x, T y, id<1> local_id) const { + T shuffle(T x, T y, id_type local_id) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down( x, y, (local_id - get_local_id()).get(0)); From 86d0210357e42255cf67d7b30de0e7fedc6a1820 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 08:40:03 -0700 Subject: [PATCH 04/12] [SYCL] Remove fence_space from sub-group barrier Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 21 ++++++++++++++++++--- 1 file changed, 18 insertions(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index d0a93e8e30544..72c0f652b8576 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -460,12 +460,27 @@ struct sub_group { } /* --- synchronization functions --- */ + void barrier() const { +#ifdef __SYCL_DEVICE_ONLY__ + __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, + __spv::MemorySemanticsMask::AcquireRelease | + __spv::MemorySemanticsMask::SubgroupMemory | + __spv::MemorySemanticsMask::WorkgroupMemory | + __spv::MemorySemanticsMask::CrossWorkgroupMemory); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + __SYCL_EXPORT_DEPRECATED("Sub-group barrier accepting fence_space is deprecated." + "Use barrier() without a fence_space instead.") void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const { #ifdef __SYCL_DEVICE_ONLY__ - uint32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); - __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, - flags); + int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); + __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, + flags); #else (void)accessSpace; throw runtime_error("Sub-groups are not supported on host device.", From a8e34177106b7c682fca7c199d93575d9d5a637e Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 10:14:05 -0700 Subject: [PATCH 05/12] [SYCL] Fix syntax in sub-group tests - get_group_range() changed from uint32_t to range<1> - get_uniform_group_range() no longer supported Signed-off-by: John Pennycook --- sycl/test/sub_group/common.cpp | 7 +------ sycl/test/sub_group/common_ocl.cpp | 7 +------ sycl/test/sub_group/sg.cl | 2 -- 3 files changed, 2 insertions(+), 14 deletions(-) diff --git a/sycl/test/sub_group/common.cpp b/sycl/test/sub_group/common.cpp index 9d2fb5f1e0314..be4f9c8cd7ac0 100644 --- a/sycl/test/sub_group/common.cpp +++ b/sycl/test/sub_group/common.cpp @@ -24,7 +24,6 @@ struct Data { unsigned int max_local_range; unsigned int group_id; unsigned int group_range; - unsigned int uniform_group_range; }; void check(queue &Queue, unsigned int G, unsigned int L) { @@ -43,9 +42,7 @@ void check(queue &Queue, unsigned int G, unsigned int L) { syclacc[NdItem.get_global_id()].max_local_range = SG.get_max_local_range().get(0); syclacc[NdItem.get_global_id()].group_id = SG.get_group_id().get(0); - syclacc[NdItem.get_global_id()].group_range = SG.get_group_range(); - syclacc[NdItem.get_global_id()].uniform_group_range = - SG.get_uniform_group_range(); + syclacc[NdItem.get_global_id()].group_range = SG.get_group_range().get(0); }); }); auto syclacc = syclbuf.get_access(); @@ -69,8 +66,6 @@ void check(queue &Queue, unsigned int G, unsigned int L) { } exit_if_not_equal(syclacc[j].group_id, group_id, "group_id"); exit_if_not_equal(syclacc[j].group_range, num_sg, "group_range"); - exit_if_not_equal(syclacc[j].uniform_group_range, num_sg, - "uniform_group_range"); } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); diff --git a/sycl/test/sub_group/common_ocl.cpp b/sycl/test/sub_group/common_ocl.cpp index 3e4cb3a7d664b..232e6c6c11acc 100644 --- a/sycl/test/sub_group/common_ocl.cpp +++ b/sycl/test/sub_group/common_ocl.cpp @@ -29,7 +29,6 @@ struct Data { unsigned int max_local_range; unsigned int group_id; unsigned int group_range; - unsigned int uniform_group_range; }; void check(queue &Queue, const int G, const int L, const char *SpvFile) { @@ -72,9 +71,7 @@ void check(queue &Queue, const int G, const int L, const char *SpvFile) { syclacc[NdItem.get_global_id()].max_local_range = SG.get_max_local_range().get(0); syclacc[NdItem.get_global_id()].group_id = SG.get_group_id().get(0); - syclacc[NdItem.get_global_id()].group_range = SG.get_group_range(); - syclacc[NdItem.get_global_id()].uniform_group_range = - SG.get_uniform_group_range(); + syclacc[NdItem.get_global_id()].group_range = SG.get_group_range().get(0); }); }); auto syclacc = syclbuf.get_access(); @@ -87,8 +84,6 @@ void check(queue &Queue, const int G, const int L, const char *SpvFile) { exit_if_not_equal(syclacc[j].group_id, oclacc[j].group_id, "group_id"); exit_if_not_equal(syclacc[j].group_range, oclacc[j].group_range, "group_range"); - exit_if_not_equal(syclacc[j].uniform_group_range, - oclacc[j].uniform_group_range, "uniform_group_range"); } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); diff --git a/sycl/test/sub_group/sg.cl b/sycl/test/sub_group/sg.cl index 0dcee4129807e..1c91667300632 100644 --- a/sycl/test/sub_group/sg.cl +++ b/sycl/test/sub_group/sg.cl @@ -12,7 +12,6 @@ struct Data { uint max_local_range; uint group_id; uint group_range; - uint uniform_group_range; }; __kernel void ocl_subgr(__global struct Data *a) { uint id = get_global_id(0); @@ -21,5 +20,4 @@ __kernel void ocl_subgr(__global struct Data *a) { a[id].max_local_range = get_max_sub_group_size(); a[id].group_id = get_sub_group_id(); a[id].group_range = get_num_sub_groups(); - a[id].uniform_group_range = get_num_sub_groups(); } From 9abb244767ba3cd24bf9893f1ce7bf7aaaaf6435 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 10:17:22 -0700 Subject: [PATCH 06/12] [SYCL] Run clang-format-9 on sub_group.hpp Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 20 +++++++++++--------- 1 file changed, 11 insertions(+), 9 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 72c0f652b8576..f67efd70f2798 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -462,25 +462,27 @@ struct sub_group { /* --- synchronization functions --- */ void barrier() const { #ifdef __SYCL_DEVICE_ONLY__ - __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, - __spv::MemorySemanticsMask::AcquireRelease | - __spv::MemorySemanticsMask::SubgroupMemory | - __spv::MemorySemanticsMask::WorkgroupMemory | - __spv::MemorySemanticsMask::CrossWorkgroupMemory); + __spirv_ControlBarrier( + __spv::Scope::Subgroup, __spv::Scope::Subgroup, + __spv::MemorySemanticsMask::AcquireRelease | + __spv::MemorySemanticsMask::SubgroupMemory | + __spv::MemorySemanticsMask::WorkgroupMemory | + __spv::MemorySemanticsMask::CrossWorkgroupMemory); #else throw runtime_error("Sub-groups are not supported on host device.", PI_INVALID_DEVICE); #endif } - __SYCL_EXPORT_DEPRECATED("Sub-group barrier accepting fence_space is deprecated." - "Use barrier() without a fence_space instead.") + __SYCL_EXPORT_DEPRECATED( + "Sub-group barrier accepting fence_space is deprecated." + "Use barrier() without a fence_space instead.") void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const { #ifdef __SYCL_DEVICE_ONLY__ int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); - __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, - flags); + __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, + flags); #else (void)accessSpace; throw runtime_error("Sub-groups are not supported on host device.", From 838001e42ecdbfbb2b7e5d176fac62eee2a6a8f6 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 18 Jun 2020 13:23:58 -0700 Subject: [PATCH 07/12] [SYCL] Extract id component 0 before cast Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index f67efd70f2798..5c3a792482976 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -158,7 +158,7 @@ struct sub_group { linear_id_type get_local_linear_id() const { #ifdef __SYCL_DEVICE_ONLY__ - return static_cast(get_local_id()); + return static_cast(get_local_id()[0]); #else throw runtime_error("Sub-groups are not supported on host device.", PI_INVALID_DEVICE); @@ -194,7 +194,7 @@ struct sub_group { linear_id_type get_group_linear_id() const { #ifdef __SYCL_DEVICE_ONLY__ - return static_cast(get_group_id()); + return static_cast(get_group_id()[0]); #else throw runtime_error("Sub-groups are not supported on host device.", PI_INVALID_DEVICE); From b1b18ad4d85da8a8ff25f670de31a7b2d4e9a487 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 19 Jun 2020 06:26:55 -0700 Subject: [PATCH 08/12] [SYCL] Remove duplicated using statements Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 16 ---------------- 1 file changed, 16 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 5c3a792482976..dc3ab74be9533 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -70,22 +70,6 @@ __SYCL_SG_GENERATE_BODY_2ARG(shuffle_up, SubgroupShuffleUpINTEL) #undef __SYCL_SG_GENERATE_BODY_2ARG -// Selects 8-bit, 16-bit or 32-bit type depending on size of T. If T doesn't -// maps to mentioned types, then void is returned -template -using SelectBlockT = - select_apply_cl_scalar_t; - -template -using AcceptableForGlobalLoadStore = - bool_constant>::value && - Space == access::address_space::global_space>; - -template -using AcceptableForLocalLoadStore = - bool_constant>::value && - Space == access::address_space::local_space>; - template T load(const multi_ptr src) { using BlockT = SelectBlockT; From 90d5156536967707902e39c9b6a7edb70444f49f Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 19 Jun 2020 10:45:56 -0700 Subject: [PATCH 09/12] [SYCL] Reintroduce __SYCL_DEPRECATED macro __SYCL_EXPORT_DEPRECATED() should not apply to header-only functions. Signed-off-by: John Pennycook --- sycl/include/CL/sycl/detail/defines.hpp | 6 ++++++ sycl/include/CL/sycl/intel/sub_group.hpp | 11 +++++------ 2 files changed, 11 insertions(+), 6 deletions(-) diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp index 7c4381fae189e..eb723a411439c 100644 --- a/sycl/include/CL/sycl/detail/defines.hpp +++ b/sycl/include/CL/sycl/detail/defines.hpp @@ -46,3 +46,9 @@ #warning "No assumptions will be emitted due to no __builtin_assume available" #endif #endif + +#ifdef _WIN32 +#define __SYCL_DEPRECATED(message) __declspec(deprecated(message)) +#else +#define __SYCL_DEPRECATED(message) __attribute__((deprecated(message))) +#endif diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index dc3ab74be9533..d1e6d9762b8e1 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -250,7 +250,7 @@ struct sub_group { /* indices in [0 , 2 * sub_group size) */ template - __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") + __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle(T x, T y, id_type local_id) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down( @@ -265,7 +265,7 @@ struct sub_group { } template - __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") + __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_down(T current, T next, uint32_t delta) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_down(current, next, delta); @@ -279,7 +279,7 @@ struct sub_group { } template - __SYCL_EXPORT_DEPRECATED("Two-input sub-group shuffles are deprecated.") + __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.") T shuffle_up(T previous, T current, uint32_t delta) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::sub_group::shuffle_up(previous, current, delta); @@ -458,9 +458,8 @@ struct sub_group { #endif } - __SYCL_EXPORT_DEPRECATED( - "Sub-group barrier accepting fence_space is deprecated." - "Use barrier() without a fence_space instead.") + __SYCL_DEPRECATED("Sub-group barrier accepting fence_space is deprecated." + "Use barrier() without a fence_space instead.") void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const { #ifdef __SYCL_DEVICE_ONLY__ From 048dca0c7a8b9e73fd3daa219c5f8335c11d64cf Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 23 Jun 2020 06:40:06 -0700 Subject: [PATCH 10/12] [SYCL] Remove default argument from barrier Meaning of barrier() was ambiguous, and could have meant: - New barrier() with no fence_space argument - Old barrier() with default fence_space argument Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index d1e6d9762b8e1..814d5731cc777 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -460,8 +460,7 @@ struct sub_group { __SYCL_DEPRECATED("Sub-group barrier accepting fence_space is deprecated." "Use barrier() without a fence_space instead.") - void barrier(access::fence_space accessSpace = - access::fence_space::global_and_local) const { + void barrier(access::fence_space accessSpace) const { #ifdef __SYCL_DEVICE_ONLY__ int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, From b382ba67936c3e8608077b558ee237b65b318498 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Mon, 29 Jun 2020 07:26:32 -0700 Subject: [PATCH 11/12] [SYCL] Restore sub-group collectives Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 68 ++++++++++++++++++++++++ 1 file changed, 68 insertions(+) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 40ed51dc715d7..cddc1174ea59c 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -470,6 +470,74 @@ struct sub_group { #endif } + /* --- deprecated collective functions --- */ + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::broadcast instead.") + EnableIfIsScalarArithmetic broadcast(T x, id<1> local_id) const { + return sycl::detail::spirv::GroupBroadcast(x, local_id); + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::reduce instead.") + EnableIfIsScalarArithmetic reduce(T x, BinaryOperation op) const { + return sycl::detail::calc( + typename sycl::detail::GroupOpTag::type(), x, op); + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::reduce instead.") + EnableIfIsScalarArithmetic reduce(T x, T init, BinaryOperation op) const { + return op(init, reduce(x, op)); + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::exclusive_scan instead.") + EnableIfIsScalarArithmetic exclusive_scan(T x, BinaryOperation op) const { + return sycl::detail::calc( + typename sycl::detail::GroupOpTag::type(), x, op); + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::exclusive_scan instead.") + EnableIfIsScalarArithmetic exclusive_scan(T x, T init, + BinaryOperation op) const { + if (get_local_id().get(0) == 0) { + x = op(init, x); + } + T scan = exclusive_scan(x, op); + if (get_local_id().get(0) == 0) { + scan = init; + } + return scan; + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::inclusive_scan instead.") + EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op) const { + return sycl::detail::calc( + typename sycl::detail::GroupOpTag::type(), x, op); + } + + template + __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " + "sycl::intel::inclusive_scan instead.") + EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op, + T init) const { + if (get_local_id().get(0) == 0) { + x = op(init, x); + } + return inclusive_scan(x, op); + } + protected: template friend class cl::sycl::nd_item; sub_group() = default; From dad7d1af63f4c4c74a6c4dd9d5c56eb591e38fcc Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Mon, 29 Jun 2020 10:33:27 -0700 Subject: [PATCH 12/12] [SYCL] Add #ifdef __SYCL_DEVICE_ONLY__ guards Signed-off-by: John Pennycook --- sycl/include/CL/sycl/intel/sub_group.hpp | 52 ++++++++++++++++++++++++ 1 file changed, 52 insertions(+) diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index cddc1174ea59c..db24282239dc8 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -475,32 +475,61 @@ struct sub_group { __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " "sycl::intel::broadcast instead.") EnableIfIsScalarArithmetic broadcast(T x, id<1> local_id) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::GroupBroadcast(x, local_id); +#else + (void)x; + (void)local_id; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " "sycl::intel::reduce instead.") EnableIfIsScalarArithmetic reduce(T x, BinaryOperation op) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::calc( typename sycl::detail::GroupOpTag::type(), x, op); +#else + (void)x; + (void)op; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " "sycl::intel::reduce instead.") EnableIfIsScalarArithmetic reduce(T x, T init, BinaryOperation op) const { +#ifdef __SYCL_DEVICE_ONLY__ return op(init, reduce(x, op)); +#else + (void)x; + (void)init; + (void)op; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " "sycl::intel::exclusive_scan instead.") EnableIfIsScalarArithmetic exclusive_scan(T x, BinaryOperation op) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::calc( typename sycl::detail::GroupOpTag::type(), x, op); +#else + (void)x; + (void)op; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -508,6 +537,7 @@ struct sub_group { "sycl::intel::exclusive_scan instead.") EnableIfIsScalarArithmetic exclusive_scan(T x, T init, BinaryOperation op) const { +#ifdef __SYCL_DEVICE_ONLY__ if (get_local_id().get(0) == 0) { x = op(init, x); } @@ -516,15 +546,29 @@ struct sub_group { scan = init; } return scan; +#else + (void)x; + (void)init; + (void)op; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " "sycl::intel::inclusive_scan instead.") EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op) const { +#ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::calc( typename sycl::detail::GroupOpTag::type(), x, op); +#else + (void)x; + (void)op; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } template @@ -532,10 +576,18 @@ struct sub_group { "sycl::intel::inclusive_scan instead.") EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op, T init) const { +#ifdef __SYCL_DEVICE_ONLY__ if (get_local_id().get(0) == 0) { x = op(init, x); } return inclusive_scan(x, op); +#else + (void)x; + (void)op; + (void)init; + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif } protected: