COMPMID-286: CL colour convert to U8

Change-Id: I62bbf510cc106a90ed2884be3c9c0c127da25898
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/150681
Reviewed-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Tested-by: bsgcomp <bsgcomp@arm.com>

7 files changed, 102 insertions, 4 deletions

diff --git a/arm_compute/core/CL/kernels/CLColorConvertKernel.h b/arm_compute/core/CL/kernels/CLColorConvertKernel.h
index edd05ef007..63e11bb738 100644
--- a/arm_compute/core/CL/kernels/CLColorConvertKernel.h
+++ b/arm_compute/core/CL/kernels/CLColorConvertKernel.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -55,7 +55,8 @@ public:
      *
      * @param[in]  input  Source tensor. Formats supported: RGBA8888/UYVY422/YUYV422/RGB888
      * @param[out] output Destination tensor. Formats supported: RGB888 (if the formats of @p input are RGBA8888/UYVY422/YUYV422),
-     *                                                          RGBA8888 (if the formats of @p input are UYVY422/YUYV422/RGB888/)
+     *                                                          RGBA8888 (if the formats of @p input are UYVY422/YUYV422/RGB888/),
+     *                                                          U8 (if the formats of @p input is RGB888)
      */
     void configure(const ICLTensor *input, ICLTensor *output);
     /** Set the input and output of the kernel
diff --git a/arm_compute/runtime/CL/functions/CLColorConvert.h b/arm_compute/runtime/CL/functions/CLColorConvert.h
index dd7de45478..8f4fa56684 100644
--- a/arm_compute/runtime/CL/functions/CLColorConvert.h
+++ b/arm_compute/runtime/CL/functions/CLColorConvert.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -43,7 +43,8 @@ public:
      *
      * @param[in]  input  Source tensor. Formats supported: RGBA8888/UYVY422/YUYV422/RGB888
      * @param[out] output Destination tensor. Formats supported: RGB888 (if the formats of @p input are RGBA8888/UYVY422/YUYV422),
-     *                                                          RGBA8888 (if the formats of @p input are UYVY422/YUYV422/RGB888/)
+     *                                                          RGBA8888 (if the formats of @p input are UYVY422/YUYV422/RGB888/),
+     *                                                          U8 (if the formats of @p input is RGB888)
      */
     void configure(const ICLTensor *input, ICLTensor *output);
     /** Initialize the function's source, destination
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 7e8ef6b22d..87b588e16c 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -352,6 +352,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
     { "RGB888_to_IYUV_bt709", "color_convert.cl" },
     { "RGB888_to_NV12_bt709", "color_convert.cl" },
     { "RGB888_to_RGBA8888_bt709", "color_convert.cl" },
+    { "RGB888_to_U8_bt709", "color_convert.cl" },
     { "RGB888_to_YUV444_bt709", "color_convert.cl" },
     { "RGBA8888_to_IYUV_bt709", "color_convert.cl" },
     { "RGBA8888_to_NV12_bt709", "color_convert.cl" },
diff --git a/src/core/CL/cl_kernels/color_convert.cl b/src/core/CL/cl_kernels/color_convert.cl
index 02a0c8ee2a..7a872b47b5 100644
--- a/src/core/CL/cl_kernels/color_convert.cl
+++ b/src/core/CL/cl_kernels/color_convert.cl
@@ -64,6 +64,54 @@ __kernel void RGB888_to_RGBA8888_bt709(
     vstore16(rgba_3, 0, out.ptr + 48);
 }
 
+/** Convert an RGB888 image to U8
+ *
+ * Global Workgroup Size [ DIV_CEIL(width, 16), height ]
+ * No offset.
+ *
+ * @param[in]  input_ptr                            Pointer to the source image. Supported Format: RGB888
+ * @param[in]  input_stride_x                       Stride of the source image in X dimension (in bytes)
+ * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  input_stride_y                       Stride of the source image in Y dimension (in bytes)
+ * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the source image
+ * @param[out] output_ptr                           Pointer to the destination image. Supported Format: U8
+ * @param[in]  output_stride_x                      Stride of the destination image in X dimension (in bytes)
+ * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  output_stride_y                      Stride of the destination image in Y dimension (in bytes)
+ * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination image
+ */
+__kernel void RGB888_to_U8_bt709(
+    IMAGE_DECLARATION(input),
+    IMAGE_DECLARATION(output))
+{
+    Image in  = CONVERT_TO_IMAGE_STRUCT(input);
+    Image out = CONVERT_TO_IMAGE_STRUCT(output);
+
+    // handle 16 pixels every time
+    const uchar16 rgb_0 = vload16(0, in.ptr);
+    const uchar16 rgb_1 = vload16(0, in.ptr + 16);
+    const uchar16 rgb_2 = vload16(0, in.ptr + 32);
+
+    //Resequence values from a sequence of 16 RGB values to sequence of 16 R, 16 G, 16 B values
+    const uchar16 rgb_r = (uchar16)(rgb_0.s0369, rgb_0.scf, rgb_1.s258b, rgb_1.se, rgb_2.s147a, rgb_2.sd);
+    const uchar16 rgb_g = (uchar16)(rgb_0.s147a, rgb_0.sd, rgb_1.s0369, rgb_1.scf, rgb_2.s258b, rgb_2.se);
+    const uchar16 rgb_b = (uchar16)(rgb_0.s258b, rgb_0.se, rgb_1.s147a, rgb_1.sd, rgb_2.s0369, rgb_2.scf);
+
+    const float16 rgb2u8_red_coef_bt709   = 0.2126f;
+    const float16 rgb2u8_green_coef_bt709 = 0.7152f;
+    const float16 rgb2u8_blue_coef_bt709  = 0.0722f;
+
+    //Computation of 16 greyscale values in float
+    const float16 greyscale_f_0 = rgb2u8_red_coef_bt709 * convert_float16(rgb_r) + rgb2u8_green_coef_bt709 * convert_float16(rgb_g) + rgb2u8_blue_coef_bt709 * convert_float16(rgb_b);
+
+    //Convert it to 16 grayscale uchar values
+    const uchar16 greyscale_u8_0 = convert_uchar16_sat_rtz(greyscale_f_0);
+
+    vstore16(greyscale_u8_0, 0, out.ptr);
+}
+
 /** Convert an RGB888 image to RGBX8888
  *
  * Global Workgroup Size [ DIV_CEIL(width, 16), height ]
diff --git a/src/core/CL/kernels/CLColorConvertKernel.cpp b/src/core/CL/kernels/CLColorConvertKernel.cpp
index e79019eab9..4f178c9d75 100644
--- a/src/core/CL/kernels/CLColorConvertKernel.cpp
+++ b/src/core/CL/kernels/CLColorConvertKernel.cpp
@@ -61,6 +61,7 @@ void CLColorConvertKernel::configure(const ICLTensor *input, ICLTensor *output)
                     num_elems_processed_per_iteration = 16;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -75,6 +76,7 @@ void CLColorConvertKernel::configure(const ICLTensor *input, ICLTensor *output)
                     num_elems_processed_per_iteration = 8;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -84,9 +86,11 @@ void CLColorConvertKernel::configure(const ICLTensor *input, ICLTensor *output)
             switch(output->info()->format())
             {
                 case Format::RGBA8888:
+                case Format::U8:
                     num_elems_processed_per_iteration = 16;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -143,6 +147,7 @@ void CLColorConvertKernel::configure(const ICLMultiImage *input, ICLImage *outpu
                     num_elems_processed_per_iteration = 4;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -220,6 +225,7 @@ void CLColorConvertKernel::configure(const ICLImage *input, ICLMultiImage *outpu
                     num_elems_read_per_iteration_x    = 16;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -235,6 +241,7 @@ void CLColorConvertKernel::configure(const ICLImage *input, ICLMultiImage *outpu
                     num_elems_read_per_iteration_x    = 8;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -303,6 +310,7 @@ void CLColorConvertKernel::configure(const ICLMultiImage *input, ICLMultiImage *
                     num_elems_processed_per_iteration = 16;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
@@ -316,6 +324,7 @@ void CLColorConvertKernel::configure(const ICLMultiImage *input, ICLMultiImage *
                     num_elems_processed_per_iteration = 16;
                     break;
                 default:
+                    ARM_COMPUTE_ERROR("Not supported");
                     break;
             }
             break;
diff --git a/tests/benchmark/CL/ColorConvert.cpp b/tests/benchmark/CL/ColorConvert.cpp
index a6337919ed..a89357b9f1 100644
--- a/tests/benchmark/CL/ColorConvert.cpp
+++ b/tests/benchmark/CL/ColorConvert.cpp
@@ -47,6 +47,9 @@ const auto ColorConvert_RGBA_to_RGB = combine(framework::dataset::make("FormatTy
 const auto ColorConvert_RGB_to_RGBA = combine(framework::dataset::make("FormatType", { Format::RGB888 }),
                                               framework::dataset::make("FormatType", { Format::RGBA8888 }));
 
+const auto ColorConvert_RGB_to_U8 = combine(framework::dataset::make("FormatType", { Format::RGB888 }),
+                                            framework::dataset::make("FormatType", { Format::U8 }));
+
 const auto ColorConvert_YUYV_to_RGBDataset = combine(YUYVDataset,
                                                      RGBDataset);
 
@@ -82,6 +85,13 @@ REGISTER_FIXTURE_DATA_TEST_CASE(RunSmall, CLColorConvertFixture, framework::Data
 REGISTER_FIXTURE_DATA_TEST_CASE(RunLarge, CLColorConvertFixture, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), ColorConvert_RGB_to_RGBA));
 TEST_SUITE_END()
 
+TEST_SUITE(RGBtoU8)
+// *INDENT-OFF*
+// clang-format off
+REGISTER_FIXTURE_DATA_TEST_CASE(RunSmall, CLColorConvertFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), ColorConvert_RGB_to_U8));
+REGISTER_FIXTURE_DATA_TEST_CASE(RunLarge, CLColorConvertFixture, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), ColorConvert_RGB_to_U8));
+TEST_SUITE_END()
+
 TEST_SUITE(YUV)
 // *INDENT-OFF*
 // clang-format off
diff --git a/tests/validation/CL/ColorConvert.cpp b/tests/validation/CL/ColorConvert.cpp
index 34b0e0db67..7210a7a891 100644
--- a/tests/validation/CL/ColorConvert.cpp
+++ b/tests/validation/CL/ColorConvert.cpp
@@ -52,6 +52,9 @@ const auto ColorConvert_RGBA_to_RGB = combine(framework::dataset::make("FormatTy
 const auto ColorConvert_RGB_to_RGBA = combine(framework::dataset::make("FormatType", { Format::RGB888 }),
                                               framework::dataset::make("FormatType", { Format::RGBA8888 }));
 
+const auto ColorConvert_RGB_to_U8 = combine(framework::dataset::make("FormatType", { Format::RGB888 }),
+                                            framework::dataset::make("FormatType", { Format::U8 }));
+
 const auto ColorConvert_YUYV_to_RGBDataset = combine(YUYVDataset,
                                                      RGBDataset);
 
@@ -143,6 +146,12 @@ DATA_TEST_CASE(RGB, framework::DatasetMode::ALL, combine(concat(datasets::Small2
     validate_configuration(shape, src_format, dst_format);
 }
 
+DATA_TEST_CASE(RGBtoU8, framework::DatasetMode::ALL, combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), ColorConvert_RGB_to_U8),
+               shape, src_format, dst_format)
+{
+    validate_configuration(shape, src_format, dst_format);
+}
+
 DATA_TEST_CASE(YUV, framework::DatasetMode::ALL, combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), ColorConvert_YUYV_to_RGBDataset),
                shape, src_format, dst_format)
 {
@@ -213,6 +222,25 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLColorConvertFixture<uint8_t>, framework::Data
 }
 TEST_SUITE_END()
 
+TEST_SUITE(RGBtoU8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLColorConvertFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), ColorConvert_RGB_to_U8))
+{
+    // Validate output
+    for(unsigned int plane_idx = 0; plane_idx < _dst_num_planes; ++plane_idx)
+    {
+        validate(CLAccessor(*_target.cl_plane(plane_idx)), _reference[plane_idx]);
+    }
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLColorConvertFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), ColorConvert_RGB_to_U8))
+{
+    // Validate output
+    for(unsigned int plane_idx = 0; plane_idx < _dst_num_planes; ++plane_idx)
+    {
+        validate(CLAccessor(*_target.cl_plane(plane_idx)), _reference[plane_idx]);
+    }
+}
+TEST_SUITE_END()
+
 TEST_SUITE(YUV)
 FIXTURE_DATA_TEST_CASE(RunSmall, CLColorConvertFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), ColorConvert_YUYV_to_RGBDataset))
 {