diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 74bb361bb96f2..700c074340261 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -763,8 +763,7 @@ class accessor : : impl(id(), range<1>{1}, BufferRef.get_range()) { #else : AccessorBaseHost( - /*Offset=*/{0, 0, 0}, - detail::convertToArrayOfN<3, 1>(range<1>{1}), + /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}), detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { @@ -773,18 +772,19 @@ class accessor : #endif } - template - accessor(buffer &BufferRef, - detail::enable_if_t &CommandGroupHandler) + template + > + accessor(buffer &BufferRef, + handler &CommandGroupHandler) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), range<1>{1}, BufferRef.get_range()) { } #else : AccessorBaseHost( - /*Offset=*/{0, 0, 0}, - detail::convertToArrayOfN<3, 1>(range<1>{1}), + /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}), detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { @@ -793,11 +793,11 @@ class accessor : #endif template 0) && ((!IsPlaceH && IsHostBuf) || - (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> - * = nullptr> - accessor(buffer &BufferRef) + typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) && + ((!IsPlaceH && IsHostBuf) || + (IsPlaceH && + (IsGlobalBuf || IsConstantBuf)))>> + accessor(buffer &BufferRef) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.get_range()) { } @@ -814,9 +814,10 @@ class accessor : #endif template 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>> - accessor(buffer &BufferRef, + typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) && + (!IsPlaceH && + (IsGlobalBuf || IsConstantBuf))>> + accessor(buffer &BufferRef, handler &CommandGroupHandler) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.get_range()) { @@ -833,10 +834,11 @@ class accessor : #endif template 0) && ((!IsPlaceH && IsHostBuf) || - (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>> - accessor(buffer &BufferRef, + typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) && + ((!IsPlaceH && IsHostBuf) || + (IsPlaceH && + (IsGlobalBuf || IsConstantBuf)))>> + accessor(buffer &BufferRef, range AccessRange, id AccessOffset = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(AccessOffset, AccessRange, BufferRef.get_range()) { @@ -854,9 +856,10 @@ class accessor : #endif template 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>> - accessor(buffer &BufferRef, + typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) && + (!IsPlaceH && + (IsGlobalBuf || IsConstantBuf))>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, range AccessRange, id AccessOffset = {}) #ifdef __SYCL_DEVICE_ONLY__ @@ -932,17 +935,17 @@ class accessor : } template - operator typename std::enable_if>::type() const { + operator typename detail::enable_if_t< + Dims == 0 && AccessMode == access::mode::atomic, atomic>() + const { const size_t LinearIndex = getLinearIndex(id()); return atomic( multi_ptr(getQualifiedPtr() + LinearIndex)); } template - typename std::enable_if<(Dims > 0) && AccessMode == access::mode::atomic, - atomic>::type + typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic, + atomic> operator[](id Index) const { const size_t LinearIndex = getLinearIndex(Index); return atomic( @@ -951,7 +954,7 @@ class accessor : template typename detail::enable_if_t>::type + atomic> operator[](size_t Index) const { const size_t LinearIndex = getLinearIndex(id(Index)); return atomic( diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp index b894bfeda9961..6fe59191e059b 100644 --- a/sycl/include/CL/sycl/buffer.hpp +++ b/sycl/include/CL/sycl/buffer.hpp @@ -21,7 +21,9 @@ class queue; template class range; template + typename AllocatorT = cl::sycl::buffer_allocator, + typename = typename std::enable_if<(dimensions > 0) && + (dimensions <= 3)>::type> class buffer { public: using value_type = T; @@ -300,9 +302,10 @@ class buffer { shared_ptr_class impl; template friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); - template friend class buffer; - template + template + friend class buffer; + template friend class accessor; range Range; // Offset field specifies the origin of the sub buffer inside the parent diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index f9c06b96709eb..34e44369e00e0 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -30,7 +30,8 @@ namespace sycl { template class accessor; -template class buffer; +template +class buffer; class handler; using buffer_allocator = detail::sycl_memory_object_allocator; diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index dd464de3d851c..fccc7cff68348 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -60,7 +60,8 @@ namespace sycl { // Forward declaration class handler; -template class buffer; +template +class buffer; namespace detail { /// This class is the default KernelName template parameter type for kernel