[SYCL] Add get_range accessor method (#1050)

Align with the change in the spec KhronosGroup/SYCL-Docs#31.

Image Accessors: Removed get_size() method and added get_range() method.
Local Accessors: Added get_range() method.

Signed-off-by: Garima Gupta <[email protected]>

Author: garimagu

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<int, Dimensions> getCountInternal() const {
+  sycl::vec<int, Dimensions> getRangeInternal() const {
     return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
         MImageObj);
   }
@@ -356,10 +354,10 @@ class image_accessor
 
 #else
 
-  sycl::vec<int, Dimensions> getCountInternal() const {
+  sycl::vec<int, Dimensions> 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<int, Dimensions>{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<info::device::image_support>(
@@ -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<int, OCLImageTy>(MImageObj);
-    int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
-    int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
-    return (ElementSize * get_count());
-  }
 
-  template <int Dims = Dimensions> size_t get_count() const;
+  size_t get_count() const { return get_range<Dimensions>().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 <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
+  range<1> get_range() const {
+    cl_int Range = getRangeInternal();
+    return range<1>(Range);
+  }
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
+  range<2> get_range() const {
+    cl_int2 Range = getRangeInternal();
+    return range<2>(Range[0], Range[1]);
+  }
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 3>>
+  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 <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
+  range<Dims> get_range() const {
+    return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
+  }
+
 #endif
 
   // Available only when:
@@ -566,7 +569,7 @@ class __image_array_slice__ {
     CoordElemType LastCoord = 0;
 
     if (std::is_same<float, CoordElemType>::value) {
-      sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getCountInternal();
+      sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
       LastCoord =
           MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
     } else {
@@ -608,27 +611,31 @@ class __image_array_slice__ {
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
-  size_t get_size() const { return MBaseAcc.getElementSize() * get_count(); }
-
-  template <int Dims = Dimensions> size_t get_count() const;
+  size_t get_count() const { return get_range<Dimensions>().size(); }
 
-  template <> size_t get_count<1>() const {
-    cl_int2 Count = MBaseAcc.getCountInternal();
-    return Count.x();
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
+  range<1> get_range() const {
+    cl_int2 Count = MBaseAcc.getRangeInternal();
+    return range<1>(Count.x());
+  }
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
+  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 <int Dims = Dimensions,
+            typename = detail::enable_if_t<(Dims == 1 || Dims == 2)>>
+  range<Dims> get_range() const {
+    return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
+  }
+
 #endif
 
 private:
@@ -1099,6 +1106,11 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
 
   size_t get_count() const { return getSize().size(); }
 
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
+  range<Dims> get_range() const {
+    return detail::convertToArrayOfN<Dims, 1>(getSize());
+  }
+
   template <int Dims = Dimensions,
             typename = detail::enable_if_t<Dims == 0 && IsAccessAnyWrite>>
   operator RefType() const {

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.

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<sycl::float4, SYCLRead>(CGH);
-      CGH.single_task<class EmptyKernel>([=]() { ImgAcc.get_size(); });
+      CGH.single_task<class EmptyKernel>([=]() { ImgAcc.get_range(); });
     });
   }
 

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();