From d2d5f759601bd91ffa0f9258dcf501c9a6955869 Mon Sep 17 00:00:00 2001 From: Pablo Tello Date: Tue, 5 Jun 2018 13:37:36 +0100 Subject: COMPMID-1199: Fixed mismatches in CLColorConvert. Mismatches caused by the CL kernel computing the green value in a different way than in NEON and C++. Luminance values must be added after multiplying the input UV values with the coefficients and not before. Change-Id: I359573a98cf12f3be5c3437c28822175a5703dbb Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/134158 Reviewed-by: Gian Marco Iodice Tested-by: Jenkins Reviewed-by: Georgios Pinitas --- src/core/CL/cl_kernels/color_convert.cl | 66 ++++++++++++++++++++++++++------- 1 file changed, 53 insertions(+), 13 deletions(-) (limited to 'src/core/CL/cl_kernels/color_convert.cl') diff --git a/src/core/CL/cl_kernels/color_convert.cl b/src/core/CL/cl_kernels/color_convert.cl index 01d8b90be9..72c74930b0 100644 --- a/src/core/CL/cl_kernels/color_convert.cl +++ b/src/core/CL/cl_kernels/color_convert.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -135,9 +135,19 @@ __kernel void UYVY422_to_RGB888_bt709( char8 cb = (char8)(uyvy.s0, uyvy.s0, uyvy.s4, uyvy.s4, uyvy.s8, uyvy.s8, uyvy.sc, uyvy.sc) - (char8)(128); char8 cr = (char8)(uyvy.s2, uyvy.s2, uyvy.s6, uyvy.s6, uyvy.sa, uyvy.sa, uyvy.se, uyvy.se) - (char8)(128); - float8 f_r = convert_float8(luma) + (float8)(0.0000f) * convert_float8(cb) + (float8)(1.5748f) * convert_float8(cr); - float8 f_g = convert_float8(luma) - (float8)(0.1873f) * convert_float8(cb) - (float8)(0.4681f) * convert_float8(cr); - float8 f_b = convert_float8(luma) + (float8)(1.8556f) * convert_float8(cb) + (float8)(0.0000f) * convert_float8(cr); + float8 red_coef_bt709 = (float8)(1.5748f); + float8 green_coef_bt709 = (float8)(-0.1873f); + float8 green_coef2_bt709 = (float8)(-0.4681f); + float8 blue_coef_bt709 = (float8)(1.8556f); + float8 lumav = convert_float8(luma); + + float8 f_r = red_coef_bt709 * convert_float8(cr); + float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); + float8 f_b = blue_coef_bt709 * convert_float8(cb); + + f_r += lumav; + f_g += lumav; + f_b += lumav; uchar8 r_0 = convert_uchar8_rtz(f_r); uchar8 g_0 = convert_uchar8_rtz(f_g); @@ -183,9 +193,19 @@ __kernel void UYVY422_to_RGBA8888_bt709( char8 cb = (char8)(uyvy.s0, uyvy.s0, uyvy.s4, uyvy.s4, uyvy.s8, uyvy.s8, uyvy.sc, uyvy.sc) - (char8)(128); char8 cr = (char8)(uyvy.s2, uyvy.s2, uyvy.s6, uyvy.s6, uyvy.sa, uyvy.sa, uyvy.se, uyvy.se) - (char8)(128); - float8 f_r = convert_float8(luma) + (float8)(0.0000f) * convert_float8(cb) + (float8)(1.5748f) * convert_float8(cr); - float8 f_g = convert_float8(luma) - (float8)(0.1873f) * convert_float8(cb) - (float8)(0.4681f) * convert_float8(cr); - float8 f_b = convert_float8(luma) + (float8)(1.8556f) * convert_float8(cb) + (float8)(0.0000f) * convert_float8(cr); + float8 red_coef_bt709 = (float8)(1.5748f); + float8 green_coef_bt709 = (float8)(-0.1873f); + float8 green_coef2_bt709 = (float8)(-0.4681f); + float8 blue_coef_bt709 = (float8)(1.8556f); + float8 lumav = convert_float8(luma); + + float8 f_r = red_coef_bt709 * convert_float8(cr); + float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); + float8 f_b = blue_coef_bt709 * convert_float8(cb); + + f_r += lumav; + f_g += lumav; + f_b += lumav; uchar8 r_0 = convert_uchar8_rtz(f_r); uchar8 g_0 = convert_uchar8_rtz(f_g); @@ -232,9 +252,19 @@ __kernel void YUYV422_to_RGB888_bt709( char8 cb = (char8)(uyvy.s1, uyvy.s1, uyvy.s5, uyvy.s5, uyvy.s9, uyvy.s9, uyvy.sd, uyvy.sd) - (char8)(128); char8 cr = (char8)(uyvy.s3, uyvy.s3, uyvy.s7, uyvy.s7, uyvy.sb, uyvy.sb, uyvy.sf, uyvy.sf) - (char8)(128); - float8 f_r = convert_float8(luma) + (float8)(0.0000f) * convert_float8(cb) + (float8)(1.5748f) * convert_float8(cr); - float8 f_g = convert_float8(luma) - (float8)(0.1873f) * convert_float8(cb) - (float8)(0.4681f) * convert_float8(cr); - float8 f_b = convert_float8(luma) + (float8)(1.8556f) * convert_float8(cb) + (float8)(0.0000f) * convert_float8(cr); + float8 red_coef_bt709 = (float8)(1.5748f); + float8 green_coef_bt709 = (float8)(-0.1873f); + float8 green_coef2_bt709 = (float8)(-0.4681f); + float8 blue_coef_bt709 = (float8)(1.8556f); + float8 lumav = convert_float8(luma); + + float8 f_r = red_coef_bt709 * convert_float8(cr); + float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); + float8 f_b = blue_coef_bt709 * convert_float8(cb); + + f_r += lumav; + f_g += lumav; + f_b += lumav; uchar8 r_0 = convert_uchar8_rtz(f_r); uchar8 g_0 = convert_uchar8_rtz(f_g); @@ -280,9 +310,19 @@ __kernel void YUYV422_to_RGBA8888_bt709( char8 cb = (char8)(uyvy.s1, uyvy.s1, uyvy.s5, uyvy.s5, uyvy.s9, uyvy.s9, uyvy.sd, uyvy.sd) - (char8)(128); char8 cr = (char8)(uyvy.s3, uyvy.s3, uyvy.s7, uyvy.s7, uyvy.sb, uyvy.sb, uyvy.sf, uyvy.sf) - (char8)(128); - float8 f_r = convert_float8(luma) + (float8)(0.0000f) * convert_float8(cb) + (float8)(1.5748f) * convert_float8(cr); - float8 f_g = convert_float8(luma) - (float8)(0.1873f) * convert_float8(cb) - (float8)(0.4681f) * convert_float8(cr); - float8 f_b = convert_float8(luma) + (float8)(1.8556f) * convert_float8(cb) + (float8)(0.0000f) * convert_float8(cr); + float8 red_coef_bt709 = (float8)(1.5748f); + float8 green_coef_bt709 = (float8)(-0.1873f); + float8 green_coef2_bt709 = (float8)(-0.4681f); + float8 blue_coef_bt709 = (float8)(1.8556f); + float8 lumav = convert_float8(luma); + + float8 f_r = red_coef_bt709 * convert_float8(cr); + float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); + float8 f_b = blue_coef_bt709 * convert_float8(cb); + + f_r += lumav; + f_g += lumav; + f_b += lumav; uchar8 r_0 = convert_uchar8_rtz(f_r); uchar8 g_0 = convert_uchar8_rtz(f_g); -- cgit v1.2.1