diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 9d1c61d474a6e..26f72b73ead8e 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -267,7 +267,6 @@ class image_accessor #ifndef __SYCL_DEVICE_ONLY__ : public detail::AccessorBaseHost { size_t MImageCount; - size_t MImageSize; image_channel_order MImgChannelOrder; image_channel_type MImgChannelType; #else @@ -277,9 +276,8 @@ class image_accessor AccessTarget>::type; OCLImageTy MImageObj; char MPadding[sizeof(detail::AccessorBaseHost) + - sizeof(size_t /*MImageSize*/) + sizeof(size_t /*MImageCount*/) + - sizeof(image_channel_order) + sizeof(image_channel_type) - - sizeof(OCLImageTy)]; + sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) + + sizeof(image_channel_type) - sizeof(OCLImageTy)]; protected: void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; } @@ -342,7 +340,7 @@ class image_accessor #ifdef __SYCL_DEVICE_ONLY__ - sycl::vec getCountInternal() const { + sycl::vec getRangeInternal() const { return __invoke_ImageQuerySize, OCLImageTy>( MImageObj); } @@ -356,10 +354,10 @@ class image_accessor #else - sycl::vec getCountInternal() const { + sycl::vec getRangeInternal() const { // TODO: Implement for host. throw runtime_error( - "image::getCountInternal() is not implemented for host"); + "image::getRangeInternal() is not implemented for host"); return sycl::vec{1}; } @@ -397,7 +395,6 @@ class image_accessor AccessMode, detail::getSyclObjImpl(ImageRef).get(), Dimensions, ImageElementSize), MImageCount(ImageRef.get_count()), - MImageSize(MImageCount * ImageElementSize), MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()), MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) { detail::EventImplPtr Event = @@ -429,7 +426,6 @@ class image_accessor AccessMode, detail::getSyclObjImpl(ImageRef).get(), Dimensions, ImageElementSize), MImageCount(ImageRef.get_count()), - MImageSize(MImageCount * ImageElementSize), MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()), MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) { checkDeviceFeatureSupported( @@ -455,32 +451,39 @@ class image_accessor // get_count() method : Returns the number of elements of the SYCL image this // SYCL accessor is accessing. // - // get_size() method : Returns the size in bytes of the SYCL image this SYCL - // accessor is accessing. Returns ElementSize*get_count(). + // get_range() method : Returns a range object which represents the number of + // elements of dataT per dimension that this accessor may access. + // The range object returned must equal to the range of the image this + // accessor is associated with. #ifdef __SYCL_DEVICE_ONLY__ - size_t get_size() const { - int ChannelType = __invoke_ImageQueryFormat(MImageObj); - int ChannelOrder = __invoke_ImageQueryOrder(MImageObj); - int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder); - return (ElementSize * get_count()); - } - template size_t get_count() const; + size_t get_count() const { return get_range().size(); } - template <> size_t get_count<1>() const { return getCountInternal(); } - template <> size_t get_count<2>() const { - cl_int2 Count = getCountInternal(); - return (Count.x() * Count.y()); - }; - template <> size_t get_count<3>() const { - cl_int3 Count = getCountInternal(); - return (Count.x() * Count.y() * Count.z()); - }; + template > + range<1> get_range() const { + cl_int Range = getRangeInternal(); + return range<1>(Range); + } + template > + range<2> get_range() const { + cl_int2 Range = getRangeInternal(); + return range<2>(Range[0], Range[1]); + } + template > + range<3> get_range() const { + cl_int3 Range = getRangeInternal(); + return range<3>(Range[0], Range[1], Range[3]); + } #else - size_t get_size() const { return MImageSize; }; size_t get_count() const { return MImageCount; }; + + template 0)>> + range get_range() const { + return detail::convertToArrayOfN(getAccessRange()); + } + #endif // Available only when: @@ -566,7 +569,7 @@ class __image_array_slice__ { CoordElemType LastCoord = 0; if (std::is_same::value) { - sycl::vec Size = MBaseAcc.getCountInternal(); + sycl::vec Size = MBaseAcc.getRangeInternal(); LastCoord = MIdx / static_cast(Size.template swizzle()); } else { @@ -608,27 +611,31 @@ class __image_array_slice__ { } #ifdef __SYCL_DEVICE_ONLY__ - size_t get_size() const { return MBaseAcc.getElementSize() * get_count(); } - - template size_t get_count() const; + size_t get_count() const { return get_range().size(); } - template <> size_t get_count<1>() const { - cl_int2 Count = MBaseAcc.getCountInternal(); - return Count.x(); + template > + range<1> get_range() const { + cl_int2 Count = MBaseAcc.getRangeInternal(); + return range<1>(Count.x()); + } + template > + range<2> get_range() const { + cl_int3 Count = MBaseAcc.getRangeInternal(); + return range<2>(Count.x(), Count.y()); } - template <> size_t get_count<2>() const { - cl_int3 Count = MBaseAcc.getCountInternal(); - return (Count.x() * Count.y()); - }; -#else - size_t get_size() const { - return MBaseAcc.MImageSize / MBaseAcc.getAccessRange()[Dimensions]; - }; +#else size_t get_count() const { return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions]; - }; + } + + template > + range get_range() const { + return detail::convertToArrayOfN(MBaseAcc.getAccessRange()); + } + #endif private: @@ -1099,6 +1106,11 @@ class accessor 0)>> + range get_range() const { + return detail::convertToArrayOfN(getSize()); + } + template > operator RefType() const { diff --git a/sycl/test/basic_tests/device_event.cpp b/sycl/test/basic_tests/device_event.cpp index 0fc5843432585..879ca90aa833b 100644 --- a/sycl/test/basic_tests/device_event.cpp +++ b/sycl/test/basic_tests/device_event.cpp @@ -95,7 +95,8 @@ int test_strideN(size_t stride) { // that are not supposed to happen, but who knows..., c) to see those // values at the end if something goes wrong during the ASYNC MEM COPY. out_ptr[item.get_global_id()[0]] = item.get_global_id()[0] + 700; - + // Just a check of get_range() API. + local_acc.get_range(); item.barrier(); // Copy from local memory to global memory. diff --git a/sycl/test/basic_tests/image.cpp b/sycl/test/basic_tests/image.cpp index 9e86268ae51c8..532731008252f 100644 --- a/sycl/test/basic_tests/image.cpp +++ b/sycl/test/basic_tests/image.cpp @@ -69,7 +69,7 @@ int main() { TestQueue Q{sycl::default_selector()}; Q.submit([&](sycl::handler &CGH) { auto ImgAcc = Img.get_access(CGH); - CGH.single_task([=]() { ImgAcc.get_size(); }); + CGH.single_task([=]() { ImgAcc.get_range(); }); }); } diff --git a/sycl/test/basic_tests/image_array.cpp b/sycl/test/basic_tests/image_array.cpp index bb34648f78d24..398cd07f34c24 100644 --- a/sycl/test/basic_tests/image_array.cpp +++ b/sycl/test/basic_tests/image_array.cpp @@ -32,7 +32,7 @@ int main() { READ_I = 0, READ_SAMPLER_F = 0, READ_SAMPLER_I = 0, - GET_SIZE, + GET_RANGE, GET_COUNT, WRITE1, WRITE2, @@ -96,11 +96,11 @@ int main() { ResAcc[READ_SAMPLER_F] |= sycl::any(sycl::isnotequal(Val, ValRef)); } - // Check that the size and count of 1D image in 1D image array == width + // Check that the range and count of 1D image in 1D image array == width // of 2d image. - ResAcc[GET_SIZE] |= (ImgAcc.get_size() / ImgSize[1]) != - ImgArrayAcc[CoordI.y()].get_size(); + ResAcc[GET_RANGE] |= sycl::range<1>(ImgAcc.get_range()[0]) != + ImgArrayAcc[CoordI.y()].get_range(); ResAcc[GET_COUNT] |= (ImgAcc.get_count() / ImgSize[1]) != ImgArrayAcc[CoordI.y()].get_count(); diff --git a/sycl/test/basic_tests/image_constructors.cpp b/sycl/test/basic_tests/image_constructors.cpp index de4bf3ee14ec2..7115f89992c30 100644 --- a/sycl/test/basic_tests/image_constructors.cpp +++ b/sycl/test/basic_tests/image_constructors.cpp @@ -5,12 +5,14 @@ // RUN: %CPU_RUN_PLACEHOLDER %t2.out // RUN: %GPU_RUN_PLACEHOLDER %t2.out // RUN: %ACC_RUN_PLACEHOLDER %t2.out +// //==-------image_constructors.cpp - SYCL image constructors basic test------==// // 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 // //===----------------------------------------------------------------------===// +// Tests the constructors, get_count and get_range APIs. #include #include @@ -35,7 +37,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { { cl::sycl::image img = cl::sycl::image(imageHostPtr, channelOrder, channelType, r); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -44,7 +47,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { { cl::sycl::image img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -55,7 +59,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -66,7 +71,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (const void*, image_channel_order, * image_channel_type, const range&, @@ -76,7 +82,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { const auto constHostPtr = imageHostPtr; cl::sycl::image img = cl::sycl::image(constHostPtr, channelOrder, channelType, r); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (const void*, image_channel_order, @@ -86,7 +93,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { const auto constHostPtr = imageHostPtr; cl::sycl::image img = cl::sycl::image( constHostPtr, channelOrder, channelType, r, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (const void*, image_channel_order, @@ -98,7 +106,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image( constHostPtr, channelOrder, channelType, r, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (const void*, image_channel_order, @@ -110,7 +119,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image( constHostPtr, channelOrder, channelType, r, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -122,7 +132,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::shared_ptr_class(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image(hostPointer, channelOrder, channelType, r); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -133,7 +144,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::shared_ptr_class(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image(hostPointer, channelOrder, channelType, r, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -146,7 +158,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::shared_ptr_class(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image(hostPointer, channelOrder, channelType, r, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -159,7 +172,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::shared_ptr_class(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image( hostPointer, channelOrder, channelType, r, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -168,7 +182,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { { cl::sycl::image img = cl::sycl::image(channelOrder, channelType, r); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -177,7 +192,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { { cl::sycl::image img = cl::sycl::image(channelOrder, channelType, r, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -187,7 +203,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image(channelOrder, channelType, r, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -197,7 +214,8 @@ void test_constructors(cl::sycl::range r, void *imageHostPtr) { cl::sycl::image_allocator imgAlloc; cl::sycl::image img = cl::sycl::image(channelOrder, channelType, r, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } } @@ -220,7 +238,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, pitch); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -230,7 +249,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, pitch, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -242,7 +262,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, pitch, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (void *, image_channel_order, @@ -253,7 +274,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image( imageHostPtr, channelOrder, channelType, r, pitch, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -265,7 +287,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image(hostPointer, channelOrder, channelType, r, pitch); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -277,7 +300,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image( hostPointer, channelOrder, channelType, r, pitch, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -291,7 +315,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image( hostPointer, channelOrder, channelType, r, pitch, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (shared_ptr_class&, image_channel_order, @@ -304,7 +329,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range(imageHostPtr, &no_delete); cl::sycl::image img = cl::sycl::image( hostPointer, channelOrder, channelType, r, pitch, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -314,7 +340,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image(channelOrder, channelType, r, pitch); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -324,7 +351,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image(channelOrder, channelType, r, pitch, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -335,7 +363,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image(channelOrder, channelType, r, pitch, imgAlloc); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } /* Constructor (image_channel_order, image_channel_type, @@ -346,7 +375,8 @@ void test_constructors_with_pitch(cl::sycl::range r, cl::sycl::range img = cl::sycl::image( channelOrder, channelType, r, pitch, imgAlloc, propList); - assert(img.get_size() == (numElems * elementSize)); + assert(img.get_count() == numElems); + assert(img.get_range() == r); } }