diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 48936bfe84c19..ddd783aa25d02 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -13,10 +13,13 @@ #include #include #include +#include #include #include #include +#include #include +#include // The file contains implementations of accessor class. Objects of accessor // class define a requirement to access some SYCL memory object or local memory @@ -234,6 +237,194 @@ class accessor_common { }; }; +// Image accessor +template +class image_accessor +#ifndef __SYCL_DEVICE_ONLY__ + : public detail::AccessorBaseHost { + size_t MImageSize; + size_t MImageCount; +#else +{ + /* + // TODO: Define the datatype here based on Dimensions, AccessMode and + AccessTarget. + __ocl_image[Dim]d[array/non_array]_[AM]_t MImage; + __init(__ocl_imagexx_t Image) { MImage = Image; } + */ +#endif + constexpr static bool IsHostImageAcc = + (AccessTarget == access::target::host_image); + + constexpr static bool IsImageAcc = (AccessTarget == access::target::image); + + constexpr static bool IsImageArrayAcc = + (AccessTarget == access::target::image_array); + + constexpr static bool IsImageAccessAnyWrite = + (AccessMode == access::mode::write || + AccessMode == access::mode::discard_write); + + constexpr static bool IsImageAccessRead = (AccessMode == access::mode::read); + + static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc, + "Expected image type"); + + static_assert(IsPlaceholder == access::placeholder::false_t, + "Expected false as Placeholder value for image accessor."); + + static_assert( + AccessMode == access::mode::read || AccessMode == access::mode::write || + AccessMode == access::mode::discard_write, + "Access modes can be only read/write/discard_write for image accessor."); + + static_assert(Dimensions > 0 && Dimensions <= 3, + "Dimensions can be 1/2/3 for image accessor."); + +public: + using value_type = DataT; + using reference = DataT &; + using const_reference = const DataT &; + + // image_accessor Constructors. + + // Available only when: accessTarget == access::target::host_image + // template + // accessor(image &imageRef); + template < + typename AllocatorT, int Dims = Dimensions, + typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>> + image_accessor(image &ImageRef, int ImageElementSize) +#ifdef __SYCL_DEVICE_ONLY__ + { + // TODO: Implement this function. + } +#else + : AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/, + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + AccessMode, detail::getSyclObjImpl(ImageRef).get(), + Dimensions, ImageElementSize), + MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) { + detail::EventImplPtr Event = + detail::Scheduler::getInstance().addHostAccessor( + AccessorBaseHost::impl.get()); + Event->wait(Event); + } +#endif + + // Available only when: accessTarget == access::target::image + // template + // accessor(image &imageRef, + // handler &commandGroupHandlerRef); + template < + typename AllocatorT, int Dims = Dimensions, + typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>> + image_accessor(image &ImageRef, + handler &CommandGroupHandlerRef, int ImageElementSize) +#ifdef __SYCL_DEVICE_ONLY__ + { + // TODO: Implement this function. + } +#else + : AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/, + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + AccessMode, detail::getSyclObjImpl(ImageRef).get(), + Dimensions, ImageElementSize), + MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) { + } +#endif + + template 0) && (Dims < 3) && + IsImageArrayAcc>> + image_accessor(image &ImageRef, + handler &CommandGroupHandlerRef, int ImageElementSize) +#ifdef __SYCL_DEVICE_ONLY__ + { + // TODO: Implement this function. + } +#else + : AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/, + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + detail::convertToArrayOfN<3, 1>(ImageRef.get_range()), + AccessMode, detail::getSyclObjImpl(ImageRef).get(), + Dimensions, ImageElementSize), + MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) { + // TODO: Implement this function. + } +#endif + + /* TODO -- common interface members -- */ + +#ifdef __SYCL_DEVICE_ONLY__ + // TODO: Define the get_size(), get_count() methods. +#else + size_t get_size() const { return MImageSize; }; + size_t get_count() const { return MImageCount; }; +#endif + + template struct IsValidCoordDataT; + template struct IsValidCoordDataT<1, T> { + constexpr static bool value = + detail::is_contained>::type::value; + }; + template struct IsValidCoordDataT<2, T> { + constexpr static bool value = detail::is_contained< + T, detail::type_list>::type::value; + }; + template struct IsValidCoordDataT<3, T> { + constexpr static bool value = detail::is_contained< + T, detail::type_list>::type::value; + }; + + // Available only when: (accessTarget == access::target::image || + // accessTarget == access::target::host_image) && accessMode == + // access::mode::read + template 0) && (IsValidCoordDataT::value) && + (IsImageAcc || IsHostImageAcc) && IsImageAccessRead>> + DataT read(const CoordT &Coords) const { + // TODO: To be implemented. + throw cl::sycl::feature_not_supported("Read API is not implemented."); + return; + }; + + // Available only when: (accessTarget == access::target::image || accessTarget + // == access::target::host_image) && accessMode == access::mode::read + template 0) && (IsValidCoordDataT::value) && + (IsImageAcc || IsHostImageAcc) && IsImageAccessRead>> + DataT read(const CoordT &Coords, const sampler &Smpl) const { + // TODO: To be implemented. + throw cl::sycl::feature_not_supported("Read API is not implemented."); + return; + }; + + // Available only when: (accessTarget == access::target::image || accessTarget + // == access::target::host_image) && accessMode == access::mode::write || + // accessMode == access::mode::discard_write + template 0) && (detail::is_intn::value) && + (IsValidCoordDataT::value) && + (IsImageAcc || IsHostImageAcc) && IsImageAccessAnyWrite>> + void write(const CoordT &Coords, const DataT &Color) const { + // TODO: To be implemented. + throw cl::sycl::feature_not_supported("Write API is not implemented."); + return; + }; + + // Available only when: accessTarget == access::target::image_array && + // dimensions < 3 + //__image_array_slice__ operator[](size_t index) const; +}; + } // namespace detail template ) + sizeof(PtrType) - sizeof(detail::AccessorBaseHost)]; @@ -720,86 +911,62 @@ class accessor -class image_accessor { - static_assert(AccessTarget == access::target::image || - AccessTarget == access::target::host_image || - AccessTarget == access::target::image_array, - "Expected image type"); - // TODO: Check if placeholder is applicable here. -public: - using value_type = DataT; - using reference = DataT &; - using const_reference = const DataT &; - - /* Available only when: accessTarget == access::target::host_image */ - // template - // accessor(image &imageRef); - /* Available only when: accessTarget == access::target::image */ - // template - // accessor(image &imageRef, - // handler &commandGroupHandlerRef); - - /* Available only when: accessTarget == access::target::image_array && - dimensions < 3 */ - // template - // accessor(image &imageRef, - // handler &commandGroupHandlerRef); - - /* TODO -- common interface members -- */ - // size_t get_size() const; - - // size_t get_count() const; - - /* Available only when: (accessTarget == access::target::image || accessTarget - == access::target::host_image) && accessMode == access::mode::read */ - // template dataT read(const coordT &coords) const; - - /* Available only when: (accessTarget == access::target::image || accessTarget - == access::target::host_image) && accessMode == access::mode::read */ - // template - // dataT read(const coordT &coords, const sampler &smpl) const; - - /* Available only when: (accessTarget == access::target::image || accessTarget - == access::target::host_image) && accessMode == access::mode::write || - accessMode == access::mode::discard_write */ - // template - // void write(const coordT &coords, const dataT &color) const; - - /* Available only when: accessTarget == access::target::image_array && - dimensions < 3 */ - //__image_array_slice__ operator[](size_t index) const; -}; - // Image accessors +// Available only when: accessTarget == access::target::host_image +// template +// accessor(image &imageRef); template class accessor - : public image_accessor {}; + : public detail::image_accessor { +public: + template + accessor(cl::sycl::image &Image, + handler &CommandGroupHandler) + : detail::image_accessor( + Image, CommandGroupHandler, + (detail::getSyclObjImpl(Image))->getElementSize()) { + CommandGroupHandler.associateWithHandler(*this); + } +}; +// Available only when: accessTarget == access::target::image +// template +// accessor(image &imageRef, +// handler &commandGroupHandlerRef); template class accessor - : public image_accessor {}; + : public detail::image_accessor { +public: + template + accessor(cl::sycl::image &Image) + : detail::image_accessor( + Image, (detail::getSyclObjImpl(Image))->getElementSize()) {} +}; +// Available only when: accessTarget == access::target::image_array && +// dimensions < 3 +// template accessor(image &imageRef, handler &commandGroupHandlerRef); template class accessor - : public image_accessor {}; + : public detail::image_accessor { + // TODO: To be Implemented. +}; } // namespace sycl } // namespace cl - - namespace std { template #include +#include #include #include #include @@ -20,9 +21,16 @@ namespace cl { namespace sycl { +// forward declarations enum class image_channel_order : unsigned int; enum class image_channel_type : unsigned int; +template class image; +template +class accessor; +class handler; + namespace detail { // utility functions and typedefs for image_impl @@ -35,10 +43,13 @@ uint8_t getImageNumberChannels(image_channel_order Order); uint8_t getImageElementSize(uint8_t NumChannels, image_channel_type Type); // validImageDataT: cl_int4, cl_uint4, cl_float4, cl_half4 -// To be used in get_access method. Uncomment after get_access is implemented. -// template -// using is_validImageDataT = typename detail::is_contained< -// T, type_list>::type; +template +using is_validImageDataT = typename detail::is_contained< + T, type_list>::type; + +template +using EnableIfImgAccDataT = + typename std::enable_if::value, DataT>::type; template class image_impl : public SYCLMemObjT { @@ -136,7 +147,8 @@ class image_impl : public SYCLMemObjT { image_impl(void *HData, image_channel_order Order, image_channel_type Type, const range &ImageRange, AllocatorT Allocator, const property_list &PropList = {}) - : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), MOrder(Order), MType(Type), + : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), + MOrder(Order), MType(Type), MNumChannels(getImageNumberChannels(MOrder)), MElementSize(getImageElementSize(MNumChannels, MType)) { setPitches(); @@ -156,7 +168,8 @@ class image_impl : public SYCLMemObjT { image_impl(const void *HData, image_channel_order Order, image_channel_type Type, const range &ImageRange, AllocatorT Allocator, const property_list &PropList = {}) - : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), MOrder(Order), MType(Type), + : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), + MOrder(Order), MType(Type), MNumChannels(getImageNumberChannels(MOrder)), MElementSize(getImageElementSize(MNumChannels, MType)) { setPitches(); @@ -179,7 +192,8 @@ class image_impl : public SYCLMemObjT { const range &ImageRange, const EnableIfPitchT &Pitch, AllocatorT Allocator, const property_list &PropList = {}) - : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), MOrder(Order), MType(Type), + : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), + MOrder(Order), MType(Type), MNumChannels(getImageNumberChannels(MOrder)), MElementSize(getImageElementSize(MNumChannels, MType)) { setPitches(Pitch); @@ -199,7 +213,8 @@ class image_impl : public SYCLMemObjT { image_impl(shared_ptr_class &HData, image_channel_order Order, image_channel_type Type, const range &ImageRange, AllocatorT Allocator, const property_list &PropList = {}) - : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), MOrder(Order), MType(Type), + : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), + MOrder(Order), MType(Type), MNumChannels(getImageNumberChannels(MOrder)), MElementSize(getImageElementSize(MNumChannels, MType)) { setPitches(); @@ -224,7 +239,8 @@ class image_impl : public SYCLMemObjT { image_channel_type Type, const range &ImageRange, const EnableIfPitchT &Pitch, AllocatorT Allocator, const property_list &PropList = {}) - : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), MOrder(Order), MType(Type), + : MAllocator(Allocator), MProps(PropList), MRange(ImageRange), + MOrder(Order), MType(Type), MNumChannels(getImageNumberChannels(MOrder)), MElementSize(getImageElementSize(MNumChannels, MType)) { setPitches(Pitch); @@ -268,6 +284,31 @@ class image_impl : public SYCLMemObjT { return MProps.get_property(); } + // Returns a valid accessor to the image with the specified access mode and + // target. The only valid types for dataT are cl_int4, cl_uint4, cl_float4 and + // cl_half4. + template > + accessor + get_access(image &Image, + handler &CommandGroupHandler) { + return accessor(Image, CommandGroupHandler); + } + + // Returns a valid accessor to the image immediately on the host with the + // specified access mode and target. The only valid types for dataT are + // cl_int4, cl_uint4, cl_float4 and cl_half4. + template //, typename = EnableIfImgAccDataT> + accessor + get_access(image &Image) { + return accessor(Image); + } + // TODO: Implement this function. void *allocateHostMem() override { if (true) @@ -307,6 +348,11 @@ class image_impl : public SYCLMemObjT { // Implementation of the pure virtual function. } + // This utility api is currently used by accessor to get the element size of + // the image. Element size is dependent on num of channels and channel type. + // This information is not accessible from the image using any public API. + uint8_t getElementSize() const { return MElementSize; }; + private: bool MHostPtrReadOnly = false; AllocatorT MAllocator; diff --git a/sycl/include/CL/sycl/image.hpp b/sycl/include/CL/sycl/image.hpp index 93fd7c1ce96c1..99519dc04c983 100644 --- a/sycl/include/CL/sycl/image.hpp +++ b/sycl/include/CL/sycl/image.hpp @@ -219,6 +219,21 @@ class image { // Returns the allocator provided to the image AllocatorT get_allocator() const { return impl->get_allocator(); } + template + accessor, Dimensions, AccessMode, + access::target::image, access::placeholder::false_t> + get_access(handler &commandGroupHandler) { + return impl->template get_access(*this, + commandGroupHandler); + } + + template + accessor, Dimensions, AccessMode, + access::target::host_image, access::placeholder::false_t> + get_access() { + return impl->template get_access(*this); + } + template void set_final_data(Destination FinalData = nullptr) { if (true) @@ -234,6 +249,7 @@ class image { private: shared_ptr_class> impl; + template friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); };