diff options
author | Pablo Tello <pablo.tello@arm.com> | 2018-06-05 13:37:36 +0100 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:52:54 +0000 |
commit | d2d5f759601bd91ffa0f9258dcf501c9a6955869 (patch) | |
tree | da746f2dc2274cc5ac901f9e0bf3088a7aa97669 /src/core/CL/cl_kernels/color_convert.cl | |
parent | 343722b441e2be9a156745db2daab305edae0db6 (diff) | |
download | ComputeLibrary-d2d5f759601bd91ffa0f9258dcf501c9a6955869.tar.gz |
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 <gianmarco.iodice@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/color_convert.cl')
-rw-r--r-- | src/core/CL/cl_kernels/color_convert.cl | 66 |
1 files changed, 53 insertions, 13 deletions
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); |