[SYCL] Rewrite image_api LIT as Google Unit test

sycl/include/CL/sycl/detail/cg.hpp

@@ -208,7 +208,7 @@ class HostKernel : public HostKernelBase {
     // If local size for host is not set explicitly, let's adjust it to 1,
     // so nd_range_error for zero local size is not thrown.
     if (AdjustedRange.LocalSize[0] == 0)
-      for (int I = 0; I < AdjustedRange.Dims; ++I)
+      for (size_t I = 0; I < AdjustedRange.Dims; ++I)
         AdjustedRange.LocalSize[I] = 1;
     if (HPI)
       HPI->start();

sycl/include/CL/sycl/detail/image_accessor_util.hpp

@@ -530,13 +530,20 @@ convertWriteData(const vec<cl_int, 4> WriteData,
   }
 }
 
+template <typename T> struct pixel_type_helper { using type = T; };
+
+template <> struct pixel_type_helper<float> { using type = int32_t; };
+
+template <> struct pixel_type_helper<double> { using type = int64_t; };
+
 template <typename ChannelType>
 vec<ChannelType, 4> processFloatDataToPixel(vec<cl_float, 4> WriteData,
                                             float MulFactor) {
   vec<cl_float, 4> Temp = WriteData * MulFactor;
   vec<cl_int, 4> TempInInt = Temp.convert<int, rounding_mode::rte>();
-  vec<cl_int, 4> TempInIntSaturated =
-      cl::sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
+  vec<cl_int, 4> TempInIntSaturated = cl::sycl::clamp(
+      TempInInt, min_v<typename pixel_type_helper<ChannelType>::type>(),
+      max_v<typename pixel_type_helper<ChannelType>::type>());
   return TempInIntSaturated.convert<ChannelType>();
 }
 
@@ -545,8 +552,6 @@ vec<ChannelType, 4>
 convertWriteData(const vec<cl_float, 4> WriteData,
                  const image_channel_type ImageChannelType) {
 
-  vec<ChannelType, 4> PixelData;
-
   switch (ImageChannelType) {
   case image_channel_type::snorm_int8:
     // convert_char_sat_rte(f * 127.0f)
@@ -574,7 +579,8 @@ convertWriteData(const vec<cl_float, 4> WriteData,
     {
       vec<cl_ushort, 4> PixelData =
           processFloatDataToPixel<cl_ushort>(WriteData, 32.0f);
-      PixelData = cl::sycl::min(PixelData, static_cast<ChannelType>(0x1f));
+      PixelData =
+          cl::sycl::min(PixelData, static_cast<const ChannelType>(0x1f));
       // Compressing the data into the first element of PixelData.
       // This is needed so that the data can be directly stored into the pixel
       // location from the first element.
@@ -616,7 +622,8 @@ convertWriteData(const vec<cl_float, 4> WriteData,
   case image_channel_type::fp32:
     return WriteData.convert<ChannelType>();
   default:
-    break;
+    throw cl::sycl::invalid_parameter_error("Unsupported data type",
+                                            PI_INVALID_VALUE);
   }
 }
 
@@ -659,7 +666,8 @@ convertWriteData(const vec<cl_half, 4> WriteData,
         "image_channel_type of the image.",
         PI_INVALID_VALUE);
   default:
-    break;
+    throw cl::sycl::invalid_parameter_error("Unsupported data type",
+                                            PI_INVALID_VALUE);
   }
 }
 
@@ -782,8 +790,8 @@ void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color,
 template <typename DataT>
 DataT ReadPixelData(const cl_int4 PixelCoord, const id<3> ImgPitch,
                     const image_channel_type ImageChannelType,
-                    const image_channel_order ImageChannelOrder,
-                    void *BasePtr, const uint8_t ElementSize) {
+                    const image_channel_order ImageChannelOrder, void *BasePtr,
+                    const uint8_t ElementSize) {
   DataT Color(0);
   auto Ptr = static_cast<unsigned char *>(BasePtr) +
              getImageOffset(PixelCoord, ImgPitch,
@@ -908,29 +916,28 @@ DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues,
   cl_int i0 = CoordValues.s0(), j0 = CoordValues.s1(), k0 = CoordValues.s2(),
          i1 = CoordValues.s4(), j1 = CoordValues.s5(), k1 = CoordValues.s6();
 
-  auto getColorInFloat =
-      [&](cl_int4 V) {
-        DataT Res = getColor<DataT>(V, SmplAddrMode,
-                                    ImgRange, ImgPitch, ImgChannelType,
-                                    ImgChannelOrder, BasePtr, ElementSize);
-        return Res.template convert<cl_float>();
-      };
+  auto getColorInFloat = [&](cl_int4 V) {
+    DataT Res =
+        getColor<DataT>(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
+                        ImgChannelOrder, BasePtr, ElementSize);
+    return Res.template convert<cl_float>();
+  };
 
   // Get Color Values at each Coordinate.
   cl_float4 Ci0j0k0 = getColorInFloat(cl_int4{i0, j0, k0, 0});
-  
+
   cl_float4 Ci1j0k0 = getColorInFloat(cl_int4{i1, j0, k0, 0});
-  
+
   cl_float4 Ci0j1k0 = getColorInFloat(cl_int4{i0, j1, k0, 0});
-  
+
   cl_float4 Ci1j1k0 = getColorInFloat(cl_int4{i1, j1, k0, 0});
-  
+
   cl_float4 Ci0j0k1 = getColorInFloat(cl_int4{i0, j0, k1, 0});
-  
+
   cl_float4 Ci1j0k1 = getColorInFloat(cl_int4{i1, j0, k1, 0});
-  
+
   cl_float4 Ci0j1k1 = getColorInFloat(cl_int4{i0, j1, k1, 0});
-  
+
   cl_float4 Ci1j1k1 = getColorInFloat(cl_int4{i1, j1, k1, 0});
 
   cl_float a = abc.x();

sycl/include/CL/sycl/handler.hpp

@@ -23,6 +23,7 @@
 #include <CL/sycl/stl.hpp>
 
 #include <algorithm>
+#include <cstring>
 #include <functional>
 #include <memory>
 #include <type_traits>
@@ -135,7 +136,7 @@ device getDeviceFromHandler(handler &);
 /// "finalization" it constructs CG object, that represents specific operation,
 /// passing fields that are required only.
 class handler {
-private:
+protected:
   /// Constructs SYCL handler from queue.
   ///
   /// \param Queue is a SYCL queue.
@@ -337,7 +338,7 @@ class handler {
     using KI = sycl::detail::KernelInfo<KernelName>;
     // Empty name indicates that the compilation happens without integration
     // header, so don't perform things that require it.
-    if (KI::getName() != "") {
+    if (std::strlen(KI::getName()) > 0) {
       MArgs.clear();
       extractArgsAndReqsFromLambda(MHostKernel->getPtr(), KI::getNumParams(),
                                    &KI::getParamDesc(0));
@@ -1241,7 +1242,7 @@ class handler {
     MCGType = detail::CG::PREFETCH_USM;
   }
 
-private:
+protected:
   shared_ptr_class<detail::queue_impl> MQueue;
   /// The storage for the arguments passed.
   /// We need to store a copy of values that are passed explicitly through

sycl/unittests/scheduler/CMakeLists.txt

@@ -11,6 +11,7 @@ add_sycl_unittest(SchedulerTests
     FinishedCmdCleanup.cpp
     LeafLimit.cpp
     MemObjCommandCleanup.cpp
+    ImageApi.cpp
     utils.cpp
 )