[SYCL] Add image accessor read with sampler for host device (#597)

garimagu · bader · commit 98d93827b878 · 2019-09-19T19:26:05.000+03:00
- Support added for sampler with NEAREST Filtering Mode.
- Check if the results are consistent with the values read by CPU/GPU
 device using a test case.
- Currently, the test case is enabled only for CPU. Seg faults on GPU.
- Added a small test case for as() function in vec class. There is no
 efficient test case in CTS or in lit-tests.

Signed-off-by: Garima Gupta &lt;garima.gupta@intel.com&gt;
diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp
@@ -97,6 +97,14 @@ getImageOffset(const vec<T, 4> &Coords, id<3> ImgPitch,
 // read from based on Addressing Mode for Nearest filter mode.
 cl_int4 getPixelCoordNearestFiltMode(cl_float4, addressing_mode, range<3>);
 
+// Check if PixelCoord are out of range for Sampler with clamp adressing mode.
+bool isOutOfRange(cl_int4 PixelCoord, addressing_mode SmplAddrMode,
+                  range<3> ImgRange);
+
+// Get Border Color for the image_channel_order, the border color values are
+// only used when the sampler has clamp addressing mode.
+cl_float4 getBorderColor(image_channel_order ImgChannelOrder);
+
 // Reads data from a pixel at Ptr location, based on the number of Channels in
 // Order and returns the data.
 // The datatype used to read from the Ptr is based on the T of the
@@ -957,12 +965,18 @@ DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,
     // Get Pixel Coordinates in integers that will be read from in the Image.
     PixelCoord =
         getPixelCoordNearestFiltMode(FloatCoorduvw, SmplAddrMode, ImgRange);
-    // TODO: Check Out-of-range coordinates. Need to use Addressing Mode Of
-    // Sampler to find the appropriate return value. Eg: clamp_to_edge returns
-    // edge values and clamp returns border color for out-of-range coordinates.
-    RetData = ReadPixelDataNearestFiltMode<DataT>(
-        PixelCoord, ImgPitch, ImgChannelType, ImgChannelOrder, BasePtr,
-        ElementSize);
+
+    // Return Border Color for out-of-range coordinates for Sampler with
+    // addressing_mode::clamp.
+
+    if (isOutOfRange(PixelCoord, SmplAddrMode, ImgRange)) {
+      cl_float4 BorderColor = (getBorderColor(ImgChannelOrder));
+      RetData = BorderColor.convert<typename TryToGetElementType<DataT>::type>();
+    } else {
+      RetData = ReadPixelDataNearestFiltMode<DataT>(
+          PixelCoord, ImgPitch, ImgChannelType, ImgChannelOrder, BasePtr,
+          ElementSize);
+    }
     break;
   }
   case filtering_mode::linear:
diff --git a/sycl/source/detail/image_accessor_util.cpp b/sycl/source/detail/image_accessor_util.cpp
@@ -15,43 +15,106 @@ namespace detail {
 
 // For Nearest Filtering mode, process cl_float4 Coordinates and return the
 // appropriate Pixel Coordinates based on Addressing Mode.
-cl_int4 getPixelCoordNearestFiltMode(cl_float4 Coord_uvw,
+cl_int4 getPixelCoordNearestFiltMode(cl_float4 Coorduvw,
                                      addressing_mode SmplAddrMode,
                                      range<3> ImgRange) {
-  cl_float u = Coord_uvw.x();
-  cl_float v = Coord_uvw.y();
-  cl_float w = Coord_uvw.z();
-
-  cl_int i = 0;
-  cl_int j = 0;
-  cl_int k = 0;
-  cl_int width = ImgRange[0];
-  cl_int height = ImgRange[1];
-  cl_int depth = ImgRange[2];
+  cl_int4 Coordijk(0);
+  cl_int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0);
   switch (SmplAddrMode) {
-  case addressing_mode::mirrored_repeat:
-    // TODO: Add the computations.
-    break;
-  case addressing_mode::repeat:
-    // TODO: Add the computations.
-    break;
+  case addressing_mode::mirrored_repeat: {
+    cl_float4 Tempuvw(0);
+    Tempuvw = 2.0f * cl::sycl::rint(0.5f * Coorduvw);
+    Tempuvw = cl::sycl::fabs(Coorduvw - Tempuvw);
+    Tempuvw = Tempuvw * (Rangewhd.convert<cl_float>());
+    Tempuvw = (cl::sycl::floor(Tempuvw));
+    Coordijk = Tempuvw.convert<cl_int>();
+    Coordijk = cl::sycl::min(Coordijk, (Rangewhd - 1));
+    // Eg:
+    // u,v,w = {2.3,1.7,0.5} // normalized coordinates.
+    // w,h,d = {9,9,9}
+    // u1=2*rint(1.15)=2
+    // v1=2*rint(0.85)=2
+    // w1=2*rint(0.5)=0
+    // u1=fabs(2.3-2)=.3
+    // v1=fabs(1.7-2)=.3
+    // w1=fabs(0.5-0)=.5
+    // u1=0.3*9=2.7
+    // v1=0.3*9=2.7
+    // w1=0.5*9=4.5
+    // i,j,k = {2,2,4}
+
+  } break;
+  case addressing_mode::repeat: {
+
+    cl_float4 Tempuvw(0);
+    Tempuvw =
+        (Coorduvw - cl::sycl::floor(Coorduvw)) * Rangewhd.convert<cl_float>();
+    Coordijk = (cl::sycl::floor(Tempuvw)).convert<cl_int>();
+    cl_int4 GreaterThanEqual = (Coordijk >= Rangewhd);
+    Coordijk =
+        cl::sycl::select(Coordijk, (Coordijk - Rangewhd), GreaterThanEqual);
+    // Eg:
+    // u = 2.3; v = 1.5; w = 0.5; // normalized coordinates.
+    // w,h,d  = {9,9,9};
+    // u1= 0.3*w;
+    // v1= 0.5*d;
+    // w1= 0.5*h;
+    // i = floor(2.7);
+    // j = floor(4.5);
+    // k = floor(4.5);
+    // if (i/j/k > w/h/d-1)
+    //      // Condition is not satisfied.
+    //      (This condition I think will only be satisfied if the floating point
+    //      arithmetic of  multiplication
+    //      gave a value in u1/v1/w1 as > w/h/d)
+    // i = 2; j = 4; k = 4;
+  } break;
   case addressing_mode::clamp_to_edge:
-    i = cl::sycl::clamp((int)cl::sycl::floor(u), 0, (width - 1));
-    j = cl::sycl::clamp((int)cl::sycl::floor(v), 0, (height - 1));
-    k = cl::sycl::clamp((int)cl::sycl::floor(w), 0, (depth - 1));
+    Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
+    Coordijk = cl::sycl::clamp(Coordijk, cl_int4(0), (Rangewhd - 1));
     break;
   case addressing_mode::clamp:
-    i = cl::sycl::clamp((int)cl::sycl::floor(u), -1, width);
-    j = cl::sycl::clamp((int)cl::sycl::floor(v), -1, height);
-    k = cl::sycl::clamp((int)cl::sycl::floor(w), -1, depth);
+    Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
+    Coordijk = cl::sycl::clamp(Coordijk, cl_int4(-1), Rangewhd);
     break;
   case addressing_mode::none:
-    i = (int)cl::sycl::floor(u);
-    j = (int)cl::sycl::floor(v);
-    k = (int)cl::sycl::floor(w);
+    Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
+    break;
+  }
+  return Coordijk;
+}
+
+bool isOutOfRange(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode,
+                  const range<3> ImgRange) {
+
+  if (SmplAddrMode != addressing_mode::clamp)
+    return false;
+
+  auto CheckOutOfRange = [](cl_int Coord, cl_int Range) {
+    return ((Coord < 0) || (Coord >= Range));
+  };
+
+  bool CheckWidth = CheckOutOfRange(PixelCoord.x(),ImgRange[0]);
+  bool CheckHeight = CheckOutOfRange(PixelCoord.y(),ImgRange[1]);
+  bool CheckDepth = CheckOutOfRange(PixelCoord.z(),ImgRange[2]);
+
+  return (CheckWidth || CheckHeight || CheckDepth);
+}
+
+cl_float4 getBorderColor(image_channel_order ImgChannelOrder) {
+
+  cl_float4 BorderColor(0.0f);
+  switch (ImgChannelOrder) {
+  case image_channel_order::r:
+  case image_channel_order::rg:
+  case image_channel_order::rgb:
+  case image_channel_order::luminance:
+    BorderColor.w() = 1.0f;
+    break;
+  default:
     break;
   }
-  return cl_int4{i, j, k, 0};
+  return BorderColor;
 }
 
 } // namespace detail
diff --git a/sycl/test/basic_tests/image_accessor_readsampler.cpp b/sycl/test/basic_tests/image_accessor_readsampler.cpp
@@ -0,0 +1,204 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==------------------- image_accessor_readsampler.cpp ---------------------==//
+//==-----------------image_accessor read API test with sampler--------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+#include <cassert>
+#include <iomanip>
+#include <iostream>
+
+namespace s = cl::sycl;
+
+template <int unique_number> class kernel_class;
+
+void validateReadData(s::cl_float4 ReadData, s::cl_float4 ExpectedColor) {
+  // Maximum difference of 1.5 ULP is allowed.
+  s::cl_int4 PixelDataInt = ReadData.template as<s::cl_int4>();
+  s::cl_int4 ExpectedDataInt = ExpectedColor.template as<s::cl_int4>();
+  s::cl_int4 Diff = ExpectedDataInt - PixelDataInt;
+#if DEBUG_OUTPUT
+  {
+    if (((s::cl_int)Diff.x() <= 1 && (s::cl_int)Diff.x() >= -1) &&
+        ((s::cl_int)Diff.y() <= 1 && (s::cl_int)Diff.y() >= -1) &&
+        ((s::cl_int)Diff.z() <= 1 && (s::cl_int)Diff.z() >= -1) &&
+        ((s::cl_int)Diff.w() <= 1 && (s::cl_int)Diff.w() >= -1)) {
+      std::cout << "Read Data is correct within precision: " << std::endl;
+    } else {
+      std::cout << "Read Data is WRONG/ outside precision: " << std::endl;
+    }
+    std::cout << "ReadData: \t"
+              << std::setprecision(std::numeric_limits<long double>::digits10 +
+                                   1)
+              << (float)ReadData.x() * 127 << "  " << (float)ReadData.y() * 127
+              << "  " << (float)ReadData.z() * 127 << "  "
+              << (float)ReadData.w() * 127 << std::endl;
+
+    std::cout << "ExpectedColor: \t" << (float)ExpectedColor.x() * 127 << "  "
+              << (float)ExpectedColor.y() * 127 << "  "
+              << (float)ExpectedColor.z() * 127 << "  "
+              << (float)ExpectedColor.w() * 127 << std::endl;
+  }
+#else
+  {
+    assert((s::cl_int)Diff.x() <= 1 && (s::cl_int)Diff.x() >= -1);
+    assert((s::cl_int)Diff.y() <= 1 && (s::cl_int)Diff.y() >= -1);
+    assert((s::cl_int)Diff.z() <= 1 && (s::cl_int)Diff.z() >= -1);
+    assert((s::cl_int)Diff.w() <= 1 && (s::cl_int)Diff.w() >= -1);
+  }
+#endif
+}
+
+template <int i>
+void checkReadSampler(char *host_ptr, s::sampler Sampler, s::cl_float4 Coord,
+                      s::cl_float4 ExpectedColor) {
+
+  s::cl_float4 ReadData;
+  {
+    // image with dim = 3
+    s::image<3> Img(host_ptr, s::image_channel_order::rgba,
+                    s::image_channel_type::snorm_int8, s::range<3>{2, 3, 4});
+    s::queue myQueue;
+    s::buffer<s::cl_float4, 1> ReadDataBuf(&ReadData, s::range<1>(1));
+    myQueue.submit([&](s::handler &cgh) {
+      auto ReadAcc = Img.get_access<s::cl_float4, s::access::mode::read>(cgh);
+      s::accessor<s::cl_float4, 1, s::access::mode::write> ReadDataBufAcc(
+          ReadDataBuf, cgh);
+
+    cgh.single_task<class kernel_class<i>>([=](){
+      s::cl_float4 RetColor = ReadAcc.read(Coord, Sampler);
+      ReadDataBufAcc[0] = RetColor;
+    });
+    });
+  }
+  validateReadData(ReadData, ExpectedColor);
+}
+
+void checkSamplerNearest() {
+
+  // create image:
+  char host_ptr[100];
+  for (int i = 0; i < 100; i++)
+    host_ptr[i] = i;
+
+  // Calling only valid configurations.
+  // A. coordinate normalization mode::normalized
+  // addressing_mode::mirrored_repeat
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f,
+                       0.0f); // Out-of-range mirrored_repeat mode
+    auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
+                              s::addressing_mode::mirrored_repeat,
+                              s::filtering_mode::nearest);
+    checkReadSampler<1>(host_ptr, Sampler, Coord,
+                        s::cl_float4((56.0f / 127.0f), (57.0f / 127.0f),
+                                     (58.0f / 127.0f),
+                                     (59.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // addressing_mode::repeat
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f); // Out-of-range repeat mode
+    auto Sampler =
+        s::sampler(s::coordinate_normalization_mode::normalized,
+                   s::addressing_mode::repeat, s::filtering_mode::nearest);
+    checkReadSampler<2>(host_ptr, Sampler, Coord,
+                        s::cl_float4((56.0f / 127.0f), (57.0f / 127.0f),
+                                     (58.0f / 127.0f),
+                                     (59.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // addressing_mode::clamp_to_edge
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f); // Out-of-range Edge Color
+    auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
+                              s::addressing_mode::clamp_to_edge,
+                              s::filtering_mode::nearest);
+    checkReadSampler<3>(host_ptr, Sampler, Coord,
+                        s::cl_float4((88.0f / 127.0f), (89.0f / 127.0f),
+                                     (90.0f / 127.0f),
+                                     (91.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // addressing_mode::clamp
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f); // Out-of-range Border Color
+    auto Sampler =
+        s::sampler(s::coordinate_normalization_mode::normalized,
+                   s::addressing_mode::clamp, s::filtering_mode::nearest);
+    checkReadSampler<4>(
+        host_ptr, Sampler, Coord,
+        s::cl_float4(0.0f, 0.0f, 0.0f, 0.0f) /*Expected Value*/);
+  }
+
+  // addressing_mode::none
+  {
+    s::cl_float4 Coord(0.0f, 0.5f, 0.75f,
+                       0.0f); // In-range for consistent return value.
+    auto Sampler =
+        s::sampler(s::coordinate_normalization_mode::normalized,
+                   s::addressing_mode::none, s::filtering_mode::nearest);
+    checkReadSampler<5>(host_ptr, Sampler, Coord,
+                        s::cl_float4((80.0f / 127.0f), (81.0f / 127.0f),
+                                     (82.0f / 127.0f),
+                                     (83.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // B. coordinate_normalization_mode::unnormalized
+  // addressing_mode::clamp_to_edge
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
+    auto Sampler = s::sampler(s::coordinate_normalization_mode::unnormalized,
+                              s::addressing_mode::clamp_to_edge,
+                              s::filtering_mode::nearest);
+    checkReadSampler<6>(host_ptr, Sampler, Coord,
+                        s::cl_float4((56.0f / 127.0f), (57.0f / 127.0f),
+                                     (58.0f / 127.0f),
+                                     (59.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // addressing_mode::clamp
+  {
+    s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
+    auto Sampler =
+        s::sampler(s::coordinate_normalization_mode::unnormalized,
+                   s::addressing_mode::clamp, s::filtering_mode::nearest);
+    checkReadSampler<7>(host_ptr, Sampler, Coord,
+                        s::cl_float4((56.0f / 127.0f), (57.0f / 127.0f),
+                                     (58.0f / 127.0f),
+                                     (59.0f / 127.0f)) /*Expected Value*/);
+  }
+
+  // addressing_mode::none
+  {
+    s::cl_float4 Coord(0.0f, 1.0f, 2.0f,
+                       0.0f); // In-range for consistent return value.
+    auto Sampler =
+        s::sampler(s::coordinate_normalization_mode::unnormalized,
+                   s::addressing_mode::none, s::filtering_mode::nearest);
+    checkReadSampler<8>(host_ptr, Sampler, Coord,
+                        s::cl_float4((56.0f / 127.0f), (57.0f / 127.0f),
+                                     (58.0f / 127.0f),
+                                     (59.0f / 127.0f)) /*Expected Value*/);
+  }
+}
+
+void checkSamplerLinear(){
+    // TODO. Implement this code.
+};
+
+int main() {
+
+  checkSamplerNearest();
+  // checkSamplerLinear();
+}
diff --git a/sycl/test/basic_tests/vectors/vector_operators.cpp b/sycl/test/basic_tests/vectors/vector_operators.cpp
