diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index afdcfa1d6d60..454e13a00dec 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -336,11 +336,11 @@ class accessor : using reference = DataT &; using const_reference = const DataT &; - template + template accessor( enable_if_t> &BufferRef) + buffer> &BufferRef) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.MemRange) { #else @@ -357,9 +357,9 @@ class accessor : #endif } - template + template accessor( - buffer &BufferRef, + buffer &BufferRef, enable_if_t &CommandGroupHandler) #ifdef __SYCL_DEVICE_ONLY__ @@ -376,11 +376,11 @@ class accessor : } #endif - template 0) && ((!IsPlaceH && IsHostBuf) || (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>> - accessor(buffer &BufferRef) + accessor(buffer &BufferRef) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.MemRange) { } @@ -398,11 +398,10 @@ class accessor : } #endif - template 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>> - accessor(buffer &BufferRef, - handler &CommandGroupHandler) + accessor(buffer &BufferRef, handler &CommandGroupHandler) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.MemRange) { } @@ -417,12 +416,12 @@ class accessor : } #endif - template 0) && ((!IsPlaceH && IsHostBuf) || (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>> - accessor(buffer &BufferRef, - range AccessRange, id AccessOffset = {}) + accessor(buffer &BufferRef, range AccessRange, + id AccessOffset = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(AccessOffset, AccessRange, BufferRef.MemRange) { } @@ -439,12 +438,11 @@ class accessor : } #endif - template 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>> - accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - id AccessOffset = {}) + accessor(buffer &BufferRef, handler &CommandGroupHandler, + range AccessRange, id AccessOffset = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(AccessOffset, AccessRange, BufferRef.MemRange) { } diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp index f150a60a341f..d9f75d520d23 100644 --- a/sycl/include/CL/sycl/buffer.hpp +++ b/sycl/include/CL/sycl/buffer.hpp @@ -21,7 +21,7 @@ class queue; template class range; template > + typename AllocatorT = cl::sycl::buffer_allocator> class buffer { public: using value_type = T; diff --git a/sycl/include/CL/sycl/detail/aligned_allocator.hpp b/sycl/include/CL/sycl/detail/aligned_allocator.hpp index a7925aac6333..d99dd9dba166 100644 --- a/sycl/include/CL/sycl/detail/aligned_allocator.hpp +++ b/sycl/include/CL/sycl/detail/aligned_allocator.hpp @@ -10,11 +10,9 @@ #include #include -#include #include #include -#include #include #include #include @@ -22,8 +20,8 @@ namespace cl { namespace sycl { -namespace detail { -template class aligned_allocator { +template +class aligned_allocator { public: using value_type = T; using pointer = T*; @@ -32,7 +30,10 @@ template class aligned_allocator { using const_reference = const T&; public: - template struct rebind { typedef aligned_allocator other; }; + template + struct rebind { + typedef aligned_allocator other; + }; // Construct an object void construct(pointer Ptr, const_reference Val) { @@ -45,15 +46,11 @@ template class aligned_allocator { pointer address(reference Val) const { return &Val; } const_pointer address(const_reference Val) { return &Val; } - // Allocate sufficiently aligned memory + // Allocate aligned (to Alignment) memory pointer allocate(size_t Size) { - size_t NumBytes = Size * sizeof(value_type); - const size_t Alignment = - std::max(getNextPowerOfTwo(sizeof(value_type)), 64); - NumBytes = ((NumBytes - 1) | (Alignment - 1)) + 1; - + Size += Alignment - Size % Alignment; pointer Result = reinterpret_cast( - detail::OSUtil::alignedAlloc(Alignment, NumBytes)); + detail::OSUtil::alignedAlloc(Alignment, Size * sizeof(value_type))); if (!Result) throw std::bad_alloc(); return Result; @@ -68,6 +65,5 @@ template class aligned_allocator { bool operator==(const aligned_allocator&) { return true; } bool operator!=(const aligned_allocator& rhs) { return false; } }; -} // namespace detail } // namespace sycl } // namespace cl diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index c0765fa00d48..bd7ad79bd012 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -22,7 +22,6 @@ #include #include -#include #include #include #include @@ -36,7 +35,7 @@ class accessor; template class buffer; class handler; -using buffer_allocator = detail::aligned_allocator; +using buffer_allocator = aligned_allocator; namespace detail { using EventImplPtr = std::shared_ptr; @@ -60,10 +59,7 @@ template class buffer_impl : public SYCLMemObjT { return; set_final_data(reinterpret_cast(HostData)); - size_t RequiredAlignment = - getNextPowerOfTwo(sizeof(typename AllocatorT::value_type)); - if (reinterpret_cast(HostData) % RequiredAlignment == 0 || - MProps.has_property()) { + if (MProps.has_property()) { MUserPtr = HostData; return; } diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index 42a522d59c9e..4031e138f8d4 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -103,9 +103,6 @@ template T createSyclObjFromImpl(decltype(T::impl) ImplObj) { return T(ImplObj); } -// Returns the smallest power of two not less than Var -size_t getNextPowerOfTwo(size_t Var); - } // namespace detail } // namespace sycl } // namespace cl diff --git a/sycl/include/CL/sycl/detail/image_impl.hpp b/sycl/include/CL/sycl/detail/image_impl.hpp index 65617cd4e8e2..51a11639081f 100644 --- a/sycl/include/CL/sycl/detail/image_impl.hpp +++ b/sycl/include/CL/sycl/detail/image_impl.hpp @@ -26,7 +26,7 @@ enum class image_channel_type : unsigned int; namespace detail { // utility functions and typedefs for image_impl -using image_allocator = aligned_allocator; +using image_allocator = aligned_allocator; // utility function: Returns the Number of Channels for a given Order. uint8_t getImageNumberChannels(image_channel_order Order); diff --git a/sycl/source/detail/common.cpp b/sycl/source/detail/common.cpp index 25379b367945..d9841251aa19 100644 --- a/sycl/source/detail/common.cpp +++ b/sycl/source/detail/common.cpp @@ -237,17 +237,6 @@ vector_class split_string(const string_class &str, return result; } -size_t getNextPowerOfTwo(size_t Var) { - --Var; - Var |= Var >> 1; - Var |= Var >> 2; - Var |= Var >> 4; - Var |= Var >> 8; - Var |= Var >> 16; - Var |= Var >> 32; - return ++Var; -} - } // namespace detail } // namespace sycl } // namespace cl