[SYCL] Improve image accessors support on a host device (#1502)

- Support more image_channel_types in image_accessors read/write API for
half4 datatype on host device.
- Enable image_accessor_readwrite test on GPU.

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

Author: garimagu

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

@@ -423,12 +423,26 @@ template <typename ChannelType>
 void convertReadData(const vec<ChannelType, 4> PixelData,
                      const image_channel_type ImageChannelType,
                      vec<cl_half, 4> &RetData) {
-
+  vec<cl_float, 4> RetDataFloat;
   switch (ImageChannelType) {
   case image_channel_type::snorm_int8:
+    //  max(-1.0f, (half)c / 127.0f)
+    RetDataFloat = (PixelData.template convert<cl_float>()) / 127.0f;
+    RetDataFloat = cl::sycl::fmax(RetDataFloat, -1);
+    break;
   case image_channel_type::snorm_int16:
+    // max(-1.0f, (half)c / 32767.0f)
+    RetDataFloat = (PixelData.template convert<cl_float>()) / 32767.0f;
+    RetDataFloat = cl::sycl::fmax(RetDataFloat, -1);
+    break;
   case image_channel_type::unorm_int8:
+    // (half)c / 255.0f
+    RetDataFloat = (PixelData.template convert<cl_float>()) / 255.0f;
+    break;
   case image_channel_type::unorm_int16:
+    // (half)c / 65535.0f
+    RetDataFloat = (PixelData.template convert<cl_float>()) / 65535.0f;
+    break;
   case image_channel_type::unorm_short_565:
   case image_channel_type::unorm_short_555:
   case image_channel_type::unorm_int_101010:
@@ -452,7 +466,7 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
         PI_INVALID_VALUE);
   case image_channel_type::fp16:
     RetData = PixelData.template convert<cl_half>();
-    break;
+    return;
   case image_channel_type::fp32:
     throw cl::sycl::invalid_parameter_error(
         "Datatype to read - cl_half4 is incompatible with the "
@@ -461,6 +475,7 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
   default:
     break;
   }
+  RetData = RetDataFloat.template convert<cl_half>();
 }
 
 // Converts data to write into appropriate datatype based on the channel of the
@@ -629,12 +644,20 @@ template <typename ChannelType>
 vec<ChannelType, 4>
 convertWriteData(const vec<cl_half, 4> WriteData,
                  const image_channel_type ImageChannelType) {
-
+  vec<cl_float, 4> WriteDataFloat = WriteData.convert<cl_float>();
   switch (ImageChannelType) {
   case image_channel_type::snorm_int8:
+    // convert_char_sat_rte(h * 127.0f)
+    return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
   case image_channel_type::snorm_int16:
+    // convert_short_sat_rte(h * 32767.0f)
+    return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
   case image_channel_type::unorm_int8:
+    // convert_uchar_sat_rte(h * 255.0f)
+    return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
   case image_channel_type::unorm_int16:
+    // convert_ushort_sat_rte(h * 65535.0f)
+    return processFloatDataToPixel<ChannelType>(WriteDataFloat, 65535.0f);
   case image_channel_type::unorm_short_565:
   case image_channel_type::unorm_short_555:
   case image_channel_type::unorm_int_101010:
@@ -994,10 +1017,13 @@ DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues,
 //     ImgChannelType.
 //   Convert to DataT as per conversion rules in section 8.3 in OpenCL Spec.
 //
-// TODO:
-// Extend support for Step2 and Step3 for Linear Filtering Mode.
-// Extend support to find out of bounds Coordinates and return appropriate
-// value based on Addressing Mode.
+// TODO: Add additional check for half datatype read.
+// Based on OpenCL spec 2.0:
+// "The read_imageh calls that take integer coordinates must use a sampler with
+// filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to
+// CLK_NORMALIZED_COORDS_FALSE and addressing mode set to
+// CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise
+// the values returned are undefined."
 
 template <typename CoordT, typename DataT>
 DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,