From 6ff3b19ee6120edf015fad8caab2991faa3070af Mon Sep 17 00:00:00 2001 From: Anthony Barbier Date: Mon, 4 Sep 2017 18:44:23 +0100 Subject: COMPMID-344 Updated doxygen Change-Id: I32f7b84daa560e460b77216add529c8fa8b327ae --- src/core/CL/cl_kernels/color_convert.cl | 1823 +++++++++++++++++++++++++++++++ 1 file changed, 1823 insertions(+) create mode 100644 src/core/CL/cl_kernels/color_convert.cl (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 new file mode 100644 index 0000000000..f5ec85ae76 --- /dev/null +++ b/src/core/CL/cl_kernels/color_convert.cl @@ -0,0 +1,1823 @@ +/* + * Copyright (c) 2016, 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +/** Convert an RGB888 image to RGBX8888 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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_RGBA8888_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 + uchar16 rgb_0 = vload16(0, in.ptr); + uchar16 rgb_1 = vload16(0, in.ptr + 16); + uchar16 rgb_2 = vload16(0, in.ptr + 32); + + uchar16 rgba_0 = (uchar16)(rgb_0.s012, 255, rgb_0.s345, 255, rgb_0.s678, 255, rgb_0.s9ab, 255); + uchar16 rgba_1 = (uchar16)(rgb_0.scde, 255, rgb_0.f, rgb_1.s01, 255, rgb_1.s234, 255, rgb_1.s567, 255); + uchar16 rgba_2 = (uchar16)(rgb_1.s89a, 255, rgb_1.sbcd, 255, rgb_1.sef, rgb_2.s0, 255, rgb_2.s123, 255); + uchar16 rgba_3 = (uchar16)(rgb_2.s456, 255, rgb_2.s789, 255, rgb_2.sabc, 255, rgb_2.sdef, 255); + + vstore16(rgba_0, 0, out.ptr); + vstore16(rgba_1, 0, out.ptr + 16); + vstore16(rgba_2, 0, out.ptr + 32); + vstore16(rgba_3, 0, out.ptr + 48); +} + +/** Convert an RGB888 image to RGBX8888 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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 RGBA8888_to_RGB888_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 + uchar16 rgba_0 = vload16(0, in.ptr); + uchar16 rgba_1 = vload16(0, in.ptr + 16); + uchar16 rgba_2 = vload16(0, in.ptr + 32); + uchar16 rgba_3 = vload16(0, in.ptr + 48); + + uchar16 rgb_0 = (uchar16)(rgba_0.s01245689, rgba_0.sacde, rgba_1.s0124); + uchar16 rgb_1 = (uchar16)(rgba_1.s5689acde, rgba_2.s01245689); + uchar16 rgb_2 = (uchar16)(rgba_2.sacde, rgba_3.s01245689, rgba_3.sacde); + + vstore16(rgb_0, 0, out.ptr); + vstore16(rgb_1, 0, out.ptr + 16); + vstore16(rgb_2, 0, out.ptr + 32); +} + +/** Convert a UYVY422 image to RGB888 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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 UYVY422_to_RGB888_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out = CONVERT_TO_IMAGE_STRUCT(output); + + // handle 8 pixels every time + uchar16 uyvy = vload16(0, in.ptr); + + uchar8 luma = (uchar8)(uyvy.s1, uyvy.s3, uyvy.s5, uyvy.s7, uyvy.s9, uyvy.sb, uyvy.sd, uyvy.sf); + 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); + + uchar8 r_0 = convert_uchar8_rtz(f_r); + uchar8 g_0 = convert_uchar8_rtz(f_g); + uchar8 b_0 = convert_uchar8_rtz(f_b); + + uchar16 rgb_0 = (uchar16)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2, b_0.s2, + r_0.s3, g_0.s3, b_0.s3, r_0.s4, g_0.s4, b_0.s4, r_0.s5); + uchar8 rgb_1 = (uchar8)(g_0.s5, b_0.s5, r_0.s6, g_0.s6, b_0.s6, r_0.s7, g_0.s7, b_0.s7); + + vstore16(rgb_0, 0, out.ptr); + vstore8(rgb_1, 0, out.ptr + 16); +} + +/** Convert a UYVY422 image to RGBX8888 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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 UYVY422_to_RGBA8888_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out = CONVERT_TO_IMAGE_STRUCT(output); + + // handle 8 pixels every time + uchar16 uyvy = vload16(0, in.ptr); + + uchar8 luma = (uchar8)(uyvy.s1, uyvy.s3, uyvy.s5, uyvy.s7, uyvy.s9, uyvy.sb, uyvy.sd, uyvy.sf); + 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); + + uchar8 r_0 = convert_uchar8_rtz(f_r); + uchar8 g_0 = convert_uchar8_rtz(f_g); + uchar8 b_0 = convert_uchar8_rtz(f_b); + + uchar16 rgba_0 = (uchar16)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255, + r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + uchar16 rgba_1 = (uchar16)(r_0.s4, g_0.s4, b_0.s4, 255, r_0.s5, g_0.s5, b_0.s5, 255, + r_0.s6, g_0.s6, b_0.s6, 255, r_0.s7, g_0.s7, b_0.s7, 255); + + vstore16(rgba_0, 0, out.ptr); + vstore16(rgba_1, 0, out.ptr + 16); +} + +/** Convert a YUYV422 image to RGB888 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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 YUYV422_to_RGB888_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out = CONVERT_TO_IMAGE_STRUCT(output); + + // handle 8 pixels every time + uchar16 uyvy = vload16(0, in.ptr); + + uchar8 luma = (uchar8)(uyvy.s0, uyvy.s2, uyvy.s4, uyvy.s6, uyvy.s8, uyvy.sa, uyvy.sc, uyvy.se); + 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); + + uchar8 r_0 = convert_uchar8_rtz(f_r); + uchar8 g_0 = convert_uchar8_rtz(f_g); + uchar8 b_0 = convert_uchar8_rtz(f_b); + + uchar16 rgb_0 = (uchar16)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2, b_0.s2, + r_0.s3, g_0.s3, b_0.s3, r_0.s4, g_0.s4, b_0.s4, r_0.s5); + uchar8 rgb_1 = (uchar8)(g_0.s5, b_0.s5, r_0.s6, g_0.s6, b_0.s6, r_0.s7, g_0.s7, b_0.s7); + + vstore16(rgb_0, 0, out.ptr); + vstore8(rgb_1, 0, out.ptr + 16); +} + +/** Convert a YUYV422 image to RGBX8888 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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 YUYV422_to_RGBA8888_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out = CONVERT_TO_IMAGE_STRUCT(output); + + // handle 8 pixels every time + uchar16 uyvy = vload16(0, in.ptr); + + uchar8 luma = (uchar8)(uyvy.s0, uyvy.s2, uyvy.s4, uyvy.s6, uyvy.s8, uyvy.sa, uyvy.sc, uyvy.se); + 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); + + uchar8 r_0 = convert_uchar8_rtz(f_r); + uchar8 g_0 = convert_uchar8_rtz(f_g); + uchar8 b_0 = convert_uchar8_rtz(f_b); + + uchar16 rgba_0 = (uchar16)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255, + r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + uchar16 rgba_1 = (uchar16)(r_0.s4, g_0.s4, b_0.s4, 255, r_0.s5, g_0.s5, b_0.s5, 255, + r_0.s6, g_0.s6, b_0.s6, 255, r_0.s7, g_0.s7, b_0.s7, 255); + + vstore16(rgba_0, 0, out.ptr); + vstore16(rgba_1, 0, out.ptr + 16); +} + +/** Convert a RGB image to NV12 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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] luma_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_step_x luma_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_step_y luma_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_offset_first_element_in_bytes The offset of the first element in the destination image luma channel + * @param[out] uv_ptr Pointer to the destination uv channel. Supported Format: U8 + * @param[in] uv_stride_x Stride of the destination uv channel in X dimension (in bytes) + * @param[in] uv_step_x uv_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] uv_step_y uv_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_offset_first_element_in_bytes The offset of the first element in the destination image uv channel + * + */ +__kernel void RGB888_to_NV12_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(luma), + IMAGE_DECLARATION(uv)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma); + Image out_uv = CONVERT_TO_IMAGE_STRUCT(uv); + + // handle 4 pixels every time, two lines, each line for 2 pixels + // Read 2 pixel of the first line + uchar8 rgb_0 = vload8(0, in.ptr); + uchar2 r_0 = (uchar2)(rgb_0.s0, rgb_0.s3); + uchar2 g_0 = (uchar2)(rgb_0.s1, rgb_0.s4); + uchar2 b_0 = (uchar2)(rgb_0.s2, rgb_0.s5); + + float2 f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_0) + (float2)(0.7152f) * convert_float2(g_0) + (float2)(0.0722f) * convert_float2(b_0); + float2 f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_0) - (float2)(0.3854f) * convert_float2(g_0) + (float2)(0.5000f) * convert_float2(b_0); + float2 f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_0) - (float2)(0.4542f) * convert_float2(g_0) - (float2)(0.0458f) * convert_float2(b_0); + + short2 i_y = convert_short2_rtz(f_y); + short2 i_u = convert_short2_rtz(f_u) + (short2)(128); + short2 i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_0 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_0, 0, out_y.ptr); + + uchar2 cb_0 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_0 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + + // Read 2 pixel of the second line + uchar8 rgb_1 = vload8(0, in.ptr + input_stride_y); + uchar2 r_1 = (uchar2)(rgb_1.s0, rgb_1.s3); + uchar2 g_1 = (uchar2)(rgb_1.s1, rgb_1.s4); + uchar2 b_1 = (uchar2)(rgb_1.s2, rgb_1.s5); + + f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_1) + (float2)(0.7152f) * convert_float2(g_1) + (float2)(0.0722f) * convert_float2(b_1); + f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_1) - (float2)(0.3854f) * convert_float2(g_1) + (float2)(0.5000f) * convert_float2(b_1); + f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_1) - (float2)(0.4542f) * convert_float2(g_1) - (float2)(0.0458f) * convert_float2(b_1); + + i_y = convert_short2_rtz(f_y); + i_u = convert_short2_rtz(f_u) + (short2)(128); + i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_1 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_1, 0, out_y.ptr + luma_stride_y); + + uchar2 cb_1 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_1 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + uchar2 cbcr = (uchar2)(((cb_0.s0 + cb_0.s1 + cb_1.s0 + cb_1.s1) / 4), + ((cr_0.s0 + cr_0.s1 + cr_1.s0 + cr_1.s1) / 4)); + + vstore2(cbcr, 0, out_uv.ptr); +} + +/* + R'= Y' + 0.0000*U + 1.5748*V + G'= Y' - 0.1873*U - 0.4681*V + B'= Y' + 1.8556*U + 0.0000*V +*/ + +/** Convert an NV12 image to RGB888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] rgb_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgb_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgb_output_step_x rgb_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgb_output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void NV12_to_RGB888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(rgb_output)) +{ + Image in_luma = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_output); + + // handle 8 pixels every time, two lines, each line for 4 pixels + uchar4 luma_0 = vload4(0, in_luma.ptr); + uchar4 luma_1 = vload4(0, in_luma.ptr + luma_input_stride_y); + uchar4 cbcr = vload4(0, in_uv.ptr); + char4 cb = (char4)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2) - (char4)(128); + char4 cr = (char4)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + uchar4 rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore4(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr + rgb_output_stride_y); + vstore4(rgb_1, 0, out_rgb.ptr + rgb_output_stride_y + 8); +} + +/** Convert a RGB image to YUV444 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] rgb_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] rgb_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] rgb_input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] rgb_input_step_y rgb_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination image V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void RGB888_to_YUV444_bt709( + IMAGE_DECLARATION(rgb_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + // handle 4 pixels every time + Image in_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // Read 4 pixel + uchar16 rgb_0 = vload16(0, in_rgb.ptr); + uchar4 r_0 = (uchar4)(rgb_0.s0, rgb_0.s3, rgb_0.s6, rgb_0.s9); + uchar4 g_0 = (uchar4)(rgb_0.s1, rgb_0.s4, rgb_0.s7, rgb_0.sa); + uchar4 b_0 = (uchar4)(rgb_0.s2, rgb_0.s5, rgb_0.s8, rgb_0.sb); + + float4 f_y = (float4)(0.0000f) + (float4)(0.2126f) * convert_float4(r_0) + (float4)(0.7152f) * convert_float4(g_0) + (float4)(0.0722f) * convert_float4(b_0); + float4 f_u = (float4)(0.0000f) - (float4)(0.1146f) * convert_float4(r_0) - (float4)(0.3854f) * convert_float4(g_0) + (float4)(0.5000f) * convert_float4(b_0); + float4 f_v = (float4)(0.0000f) + (float4)(0.5000f) * convert_float4(r_0) - (float4)(0.4542f) * convert_float4(g_0) - (float4)(0.0458f) * convert_float4(b_0); + + short4 i_y = convert_short4_rtz(f_y); + short4 i_u = convert_short4_rtz(f_u) + (short4)(128); + short4 i_v = convert_short4_rtz(f_v) + (short4)(128); + + uchar4 luma_0 = convert_uchar4(max((short4)(0), min(i_y, (short4)(255)))); + vstore4(luma_0, 0, out_y.ptr); + + uchar4 cb_0 = convert_uchar4(max((short4)(0), min(i_u, (short4)(255)))); + uchar4 cr_0 = convert_uchar4(max((short4)(0), min(i_v, (short4)(255)))); + vstore4(cb_0, 0, out_u.ptr); + vstore4(cr_0, 0, out_v.ptr); +} + +/** Convert a RGB image to IYUV using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 2), height ] + * No offset. + * + * @param[in] rgb_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] rgb_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] rgb_input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] rgb_input_step_y rgb_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void RGB888_to_IYUV_bt709( + IMAGE_DECLARATION(rgb_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + // handle 4 pixels every time, two lines, each line for 2 pixels + Image in_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // Read 2 pixel of the first line + uchar8 rgb_0 = vload8(0, in_rgb.ptr); + uchar2 r_0 = (uchar2)(rgb_0.s0, rgb_0.s3); + uchar2 g_0 = (uchar2)(rgb_0.s1, rgb_0.s4); + uchar2 b_0 = (uchar2)(rgb_0.s2, rgb_0.s5); + + float2 f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_0) + (float2)(0.7152f) * convert_float2(g_0) + (float2)(0.0722f) * convert_float2(b_0); + float2 f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_0) - (float2)(0.3854f) * convert_float2(g_0) + (float2)(0.5000f) * convert_float2(b_0); + float2 f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_0) - (float2)(0.4542f) * convert_float2(g_0) - (float2)(0.0458f) * convert_float2(b_0); + + short2 i_y = convert_short2_rtz(f_y); + short2 i_u = convert_short2_rtz(f_u) + (short2)(128); + short2 i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_0 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_0, 0, out_y.ptr); + + uchar2 cb_0 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_0 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + + // Read 2 pixel of the second line + uchar8 rgb_1 = vload8(0, in_rgb.ptr + rgb_input_stride_y); + uchar2 r_1 = (uchar2)(rgb_1.s0, rgb_1.s3); + uchar2 g_1 = (uchar2)(rgb_1.s1, rgb_1.s4); + uchar2 b_1 = (uchar2)(rgb_1.s2, rgb_1.s5); + + f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_1) + (float2)(0.7152f) * convert_float2(g_1) + (float2)(0.0722f) * convert_float2(b_1); + f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_1) - (float2)(0.3854f) * convert_float2(g_1) + (float2)(0.5000f) * convert_float2(b_1); + f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_1) - (float2)(0.4542f) * convert_float2(g_1) - (float2)(0.0458f) * convert_float2(b_1); + + i_y = convert_short2_rtz(f_y); + i_u = convert_short2_rtz(f_u) + (short2)(128); + i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_1 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_1, 0, out_y.ptr + luma_output_stride_y); + + uchar2 cb_1 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_1 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + uchar2 cbcr = (uchar2)(((cb_0.s0 + cb_0.s1 + cb_1.s0 + cb_1.s1) / 4), + ((cr_0.s0 + cr_0.s1 + cr_1.s0 + cr_1.s1) / 4)); + *out_u.ptr = cbcr.x; + *out_v.ptr = cbcr.y; +} + +/** Convert a RGBA image to YUV444 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] rgba_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] rgba_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] rgba_input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgba_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] rgba_input_step_y rgb_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgba_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination image V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void RGBA8888_to_YUV444_bt709( + IMAGE_DECLARATION(rgba_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + // handle 4 pixels every time + Image in_rgba = CONVERT_TO_IMAGE_STRUCT(rgba_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // Read 4 pixel + uchar16 rgb_0 = vload16(0, in_rgba.ptr); + uchar4 r_0 = (uchar4)(rgb_0.s0, rgb_0.s4, rgb_0.s8, rgb_0.sc); + uchar4 g_0 = (uchar4)(rgb_0.s1, rgb_0.s5, rgb_0.s9, rgb_0.sd); + uchar4 b_0 = (uchar4)(rgb_0.s2, rgb_0.s6, rgb_0.sa, rgb_0.se); + + float4 f_y = (float4)(0.0000f) + (float4)(0.2126f) * convert_float4(r_0) + (float4)(0.7152f) * convert_float4(g_0) + (float4)(0.0722f) * convert_float4(b_0); + float4 f_u = (float4)(0.0000f) - (float4)(0.1146f) * convert_float4(r_0) - (float4)(0.3854f) * convert_float4(g_0) + (float4)(0.5000f) * convert_float4(b_0); + float4 f_v = (float4)(0.0000f) + (float4)(0.5000f) * convert_float4(r_0) - (float4)(0.4542f) * convert_float4(g_0) - (float4)(0.0458f) * convert_float4(b_0); + + short4 i_y = convert_short4(f_y); + short4 i_u = convert_short4(f_u) + (short4)(128); + short4 i_v = convert_short4(f_v) + (short4)(128); + + uchar4 luma_0 = convert_uchar4_sat(max((short4)(0), min(i_y, (short4)(255)))); + vstore4(luma_0, 0, out_y.ptr); + + uchar4 cb_0 = convert_uchar4_sat(max((short4)(0), min(i_u, (short4)(255)))); + uchar4 cr_0 = convert_uchar4_sat(max((short4)(0), min(i_v, (short4)(255)))); + vstore4(cb_0, 0, out_u.ptr); + vstore4(cr_0, 0, out_v.ptr); +} + +/** Convert a RGBA image to NV12 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 2), height ] + * No offset. + * + * @param[in] input_ptr Pointer to the source image. Supported Format: U8 + * @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] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination image luma channel + * @param[out] uv_output_ptr Pointer to the destination uv channel. Supported Format: U8 + * @param[in] uv_output_stride_x Stride of the destination uv channel in X dimension (in bytes) + * @param[in] uv_output_step_x uv_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_output_stride_y Stride of the destination image uv channel in Y dimension (in bytes) + * @param[in] uv_output_step_y uv_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_output_offset_first_element_in_bytes The offset of the first element in the destination image uv channel + * + */ +__kernel void RGBA8888_to_NV12_bt709( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(uv_output)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_uv = CONVERT_TO_IMAGE_STRUCT(uv_output); + + // Read 2 pixel of the first line + uchar8 rgb_0 = vload8(0, in.ptr); + uchar2 r_0 = (uchar2)(rgb_0.s0, rgb_0.s4); + uchar2 g_0 = (uchar2)(rgb_0.s1, rgb_0.s5); + uchar2 b_0 = (uchar2)(rgb_0.s2, rgb_0.s6); + + float2 f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_0) + (float2)(0.7152f) * convert_float2(g_0) + (float2)(0.0722f) * convert_float2(b_0); + float2 f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_0) - (float2)(0.3854f) * convert_float2(g_0) + (float2)(0.5000f) * convert_float2(b_0); + float2 f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_0) - (float2)(0.4542f) * convert_float2(g_0) - (float2)(0.0458f) * convert_float2(b_0); + + short2 i_y = convert_short2_rtz(f_y); + short2 i_u = convert_short2_rtz(f_u) + (short2)(128); + short2 i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_0 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_0, 0, out_y.ptr); + + uchar2 cb_0 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_0 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + + // Read 2 pixel of the second line + uchar8 rgb_1 = vload8(0, in.ptr + input_stride_y); + uchar2 r_1 = (uchar2)(rgb_1.s0, rgb_1.s4); + uchar2 g_1 = (uchar2)(rgb_1.s1, rgb_1.s5); + uchar2 b_1 = (uchar2)(rgb_1.s2, rgb_1.s6); + + f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_1) + (float2)(0.7152f) * convert_float2(g_1) + (float2)(0.0722f) * convert_float2(b_1); + f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_1) - (float2)(0.3854f) * convert_float2(g_1) + (float2)(0.5000f) * convert_float2(b_1); + f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_1) - (float2)(0.4542f) * convert_float2(g_1) - (float2)(0.0458f) * convert_float2(b_1); + + i_y = convert_short2_rtz(f_y); + i_u = convert_short2_rtz(f_u) + (short2)(128); + i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_1 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_1, 0, out_y.ptr + luma_output_stride_y); + + uchar2 cb_1 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_1 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + uchar2 cbcr = (uchar2)(((cb_0.s0 + cb_0.s1 + cb_1.s0 + cb_1.s1) / 4), + ((cr_0.s0 + cr_0.s1 + cr_1.s0 + cr_1.s1) / 4)); + vstore2(cbcr, 0, out_uv.ptr); +} + +/** Convert a RGBA image to IYUV using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 2), height ] + * No offset. + * + * @param[in] rgba_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] rgba_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] rgba_input_step_x rgba_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgba_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] rgba_input_step_y rgba_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgba_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void RGBA8888_to_IYUV_bt709( + IMAGE_DECLARATION(rgba_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + // handle 4 pixels every time, two lines, each line for 2 pixels + Image in_rgb = CONVERT_TO_IMAGE_STRUCT(rgba_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // Read 2 pixel of the first line + uchar8 rgb_0 = vload8(0, in_rgb.ptr); + uchar2 r_0 = (uchar2)(rgb_0.s0, rgb_0.s4); + uchar2 g_0 = (uchar2)(rgb_0.s1, rgb_0.s5); + uchar2 b_0 = (uchar2)(rgb_0.s2, rgb_0.s6); + + float2 f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_0) + (float2)(0.7152f) * convert_float2(g_0) + (float2)(0.0722f) * convert_float2(b_0); + float2 f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_0) - (float2)(0.3854f) * convert_float2(g_0) + (float2)(0.5000f) * convert_float2(b_0); + float2 f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_0) - (float2)(0.4542f) * convert_float2(g_0) - (float2)(0.0458f) * convert_float2(b_0); + + short2 i_y = convert_short2_rtz(f_y); + short2 i_u = convert_short2_rtz(f_u) + (short2)(128); + short2 i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_0 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_0, 0, out_y.ptr); + + uchar2 cb_0 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_0 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + + // Read 2 pixel of the second line + uchar8 rgb_1 = vload8(0, in_rgb.ptr + rgba_input_stride_y); + uchar2 r_1 = (uchar2)(rgb_1.s0, rgb_1.s4); + uchar2 g_1 = (uchar2)(rgb_1.s1, rgb_1.s5); + uchar2 b_1 = (uchar2)(rgb_1.s2, rgb_1.s6); + + f_y = (float2)(0.0000f) + (float2)(0.2126f) * convert_float2(r_1) + (float2)(0.7152f) * convert_float2(g_1) + (float2)(0.0722f) * convert_float2(b_1); + f_u = (float2)(0.0000f) - (float2)(0.1146f) * convert_float2(r_1) - (float2)(0.3854f) * convert_float2(g_1) + (float2)(0.5000f) * convert_float2(b_1); + f_v = (float2)(0.0000f) + (float2)(0.5000f) * convert_float2(r_1) - (float2)(0.4542f) * convert_float2(g_1) - (float2)(0.0458f) * convert_float2(b_1); + + i_y = convert_short2_rtz(f_y); + i_u = convert_short2_rtz(f_u) + (short2)(128); + i_v = convert_short2_rtz(f_v) + (short2)(128); + + uchar2 luma_1 = convert_uchar2(max((short2)(0), min(i_y, (short2)(255)))); + vstore2(luma_1, 0, out_y.ptr + luma_output_stride_y); + + uchar2 cb_1 = convert_uchar2(max((short2)(0), min(i_u, (short2)(255)))); + uchar2 cr_1 = convert_uchar2(max((short2)(0), min(i_v, (short2)(255)))); + uchar2 cbcr = (uchar2)(((cb_0.s0 + cb_0.s1 + cb_1.s0 + cb_1.s1) / 4), + ((cr_0.s0 + cr_0.s1 + cr_1.s0 + cr_1.s1) / 4)); + *out_u.ptr = cbcr.x; + *out_v.ptr = cbcr.y; +} + +/** Convert an NV12 image to RGB8888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] rgb_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgb_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgb_output_step_x rgb_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgb_output_step_y rgb_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void NV12_to_RGBA8888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(rgb_output)) +{ + Image in_luma = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_output); + + uchar4 luma_0 = vload4(0, in_luma.ptr); + uchar4 luma_1 = vload4(0, in_luma.ptr + luma_input_stride_y); + uchar4 cbcr = vload4(0, in_uv.ptr); + char4 cb = (char4)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2) - (char4)(128); + char4 cr = (char4)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + uchar8 rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore8(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr + rgb_output_stride_y); + vstore8(rgb_1, 0, out_rgb.ptr + rgb_output_stride_y + 8); +} + +/** Convert an NV12 image to IYUV + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + */ +__kernel void NV12_to_IYUV_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 32 pixels every time, two lines, each line for 16 pixels + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar16 cbcr = vload16(0, in_uv.ptr); + uchar8 cb = (uchar8)(cbcr.s0, cbcr.s2, cbcr.s4, cbcr.s6, cbcr.s8, cbcr.sa, cbcr.sc, cbcr.se); + uchar8 cr = (uchar8)(cbcr.s1, cbcr.s3, cbcr.s5, cbcr.s7, cbcr.s9, cbcr.sb, cbcr.sd, cbcr.sf); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore8(cb, 0, out_u.ptr); + vstore8(cr, 0, out_v.ptr); +} + +/** Convert an NV12 image to YUV444 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + */ +__kernel void NV12_to_YUV444_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 32 pixels every time, two lines, each line for 16 pixels + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar16 cbcr = vload16(0, in_uv.ptr); + uchar16 cb = (uchar16)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2, cbcr.s4, cbcr.s4, cbcr.s6, cbcr.s6, cbcr.s8, cbcr.s8, + cbcr.sa, cbcr.sa, cbcr.sc, cbcr.sc, cbcr.se, cbcr.se); + uchar16 cr = (uchar16)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3, cbcr.s5, cbcr.s5, cbcr.s7, cbcr.s7, cbcr.s9, cbcr.s9, + cbcr.sb, cbcr.sb, cbcr.sd, cbcr.sd, cbcr.sf, cbcr.sf); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore16(cb, 0, out_u.ptr); + vstore16(cb, 0, out_u.ptr + u_output_stride_y); + vstore16(cr, 0, out_v.ptr); + vstore16(cr, 0, out_v.ptr + v_output_stride_y); +} + +/** Convert an NV21 image to RGB888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] rgb_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgb_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgb_output_step_x rgb_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgb_output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void NV21_to_RGB888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(rgb_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_output); + + // handle 8 pixels every time, two lines, each line for 4 pixels + uchar4 luma_0 = vload4(0, in_y.ptr); + uchar4 luma_1 = vload4(0, in_y.ptr + luma_input_stride_y); + uchar4 cbcr = vload4(0, in_uv.ptr); + char4 cr = (char4)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2) - (char4)(128); + char4 cb = (char4)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + uchar4 rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore4(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr + rgb_output_stride_y); + vstore4(rgb_1, 0, out_rgb.ptr + rgb_output_stride_y + 8); +} + +/** Convert an NV12 image to RGB8888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] rgba_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgba_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgba_output_step_x rgba_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgba_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgba_output_step_y rgba_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgba_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void NV21_to_RGBA8888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(rgba_output)) +{ + Image in_luma = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgba_output); + + // handle 8 pixels every time, two lines, each line for 4 pixels + uchar4 luma_0 = vload4(0, in_luma.ptr); + uchar4 luma_1 = vload4(0, in_luma.ptr + luma_input_stride_y); + uchar4 cbcr = vload4(0, in_uv.ptr); + char4 cr = (char4)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2) - (char4)(128); + char4 cb = (char4)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + uchar8 rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore8(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr + rgba_output_stride_y); + vstore8(rgb_1, 0, out_rgb.ptr + rgba_output_stride_y + 8); +} + +/** Convert an NV21 image to YUV444 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + */ +__kernel void NV21_to_YUV444_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 32 pixels every time, two lines, each line for 16 pixels + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar16 cbcr = vload16(0, in_uv.ptr); + uchar16 cr = (uchar16)(cbcr.s0, cbcr.s0, cbcr.s2, cbcr.s2, cbcr.s4, cbcr.s4, cbcr.s6, cbcr.s6, cbcr.s8, cbcr.s8, + cbcr.sa, cbcr.sa, cbcr.sc, cbcr.sc, cbcr.se, cbcr.se); + uchar16 cb = (uchar16)(cbcr.s1, cbcr.s1, cbcr.s3, cbcr.s3, cbcr.s5, cbcr.s5, cbcr.s7, cbcr.s7, cbcr.s9, cbcr.s9, + cbcr.sb, cbcr.sb, cbcr.sd, cbcr.sd, cbcr.sf, cbcr.sf); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore16(cb, 0, out_u.ptr); + vstore16(cb, 0, out_u.ptr + u_output_stride_y); + vstore16(cr, 0, out_v.ptr); + vstore16(cr, 0, out_v.ptr + v_output_stride_y); +} + +/** Convert an NV21 image to IYUV + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] uv_input_ptr Pointer to the source uv channel. Supported Format: U8 + * @param[in] uv_input_stride_x Stride of the source image uv channel in X dimension (in bytes) + * @param[in] uv_input_step_x uv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uv_input_step_y uv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + */ +__kernel void NV21_to_IYUV_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(uv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_uv = CONVERT_TO_IMAGE_STRUCT(uv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar16 cbcr = vload16(0, in_uv.ptr); + uchar8 cr = (uchar8)(cbcr.s0, cbcr.s2, cbcr.s4, cbcr.s6, cbcr.s8, cbcr.sa, cbcr.sc, cbcr.se); + uchar8 cb = (uchar8)(cbcr.s1, cbcr.s3, cbcr.s5, cbcr.s7, cbcr.s9, cbcr.sb, cbcr.sd, cbcr.sf); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore8(cb, 0, out_u.ptr); + vstore8(cr, 0, out_v.ptr); +} + +/** Convert a UYVY image to IYUV using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] uyvy_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] uyvy_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] uyvy_input_step_x uyvy_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uyvy_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] uyvy_input_step_y uyvy_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uyvy_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void UYVY422_to_IYUV_bt709( + IMAGE_DECLARATION(uyvy_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_uyvy = CONVERT_TO_IMAGE_STRUCT(uyvy_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 16 pixels every time, each line 8 pixels + uchar16 uyvy = vload16(0, in_uyvy.ptr); + uchar8 luma = (uchar8)(uyvy.s1, uyvy.s3, uyvy.s5, uyvy.s7, uyvy.s9, uyvy.sb, uyvy.sd, uyvy.sf); + ushort4 cb_0 = (ushort4)(uyvy.s0, uyvy.s4, uyvy.s8, uyvy.sc); + ushort4 cr_0 = (ushort4)(uyvy.s2, uyvy.s6, uyvy.sa, uyvy.se); + vstore8(luma, 0, out_y.ptr); + + uyvy = vload16(0, in_uyvy.ptr + uyvy_input_stride_y); + luma = (uchar8)(uyvy.s1, uyvy.s3, uyvy.s5, uyvy.s7, uyvy.s9, uyvy.sb, uyvy.sd, uyvy.sf); + ushort4 cb_1 = (ushort4)(uyvy.s0, uyvy.s4, uyvy.s8, uyvy.sc); + ushort4 cr_1 = (ushort4)(uyvy.s2, uyvy.s6, uyvy.sa, uyvy.se); + vstore8(luma, 0, out_y.ptr + luma_output_stride_y); + + uchar4 cb = convert_uchar4((cb_0 + cb_1) / (ushort4)(2)); + uchar4 cr = convert_uchar4((cr_0 + cr_1) / (ushort4)(2)); + vstore4(cb, 0, out_u.ptr); + vstore4(cr, 0, out_v.ptr); +} + +/** Convert a YUYV image to IYUV using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] yuyv_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] yuyv_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] yuyv_input_step_x yuyv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] yuyv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] yuyv_input_step_y yuyv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] yuyv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void YUYV422_to_IYUV_bt709( + IMAGE_DECLARATION(yuyv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_yuyv = CONVERT_TO_IMAGE_STRUCT(yuyv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 16 pixels every time, each line 8 pixels + uchar16 yuyv = vload16(0, in_yuyv.ptr); + uchar8 luma = (uchar8)(yuyv.s0, yuyv.s2, yuyv.s4, yuyv.s6, yuyv.s8, yuyv.sa, yuyv.sc, yuyv.se); + ushort4 cb_0 = (ushort4)(yuyv.s1, yuyv.s5, yuyv.s9, yuyv.sd); + ushort4 cr_0 = (ushort4)(yuyv.s3, yuyv.s7, yuyv.sb, yuyv.sf); + vstore8(luma, 0, out_y.ptr); + + yuyv = vload16(0, in_yuyv.ptr + yuyv_input_stride_y); + luma = (uchar8)(yuyv.s0, yuyv.s2, yuyv.s4, yuyv.s6, yuyv.s8, yuyv.sa, yuyv.sc, yuyv.se); + ushort4 cb_1 = (ushort4)(yuyv.s1, yuyv.s5, yuyv.s9, yuyv.sd); + ushort4 cr_1 = (ushort4)(yuyv.s3, yuyv.s7, yuyv.sb, yuyv.sf); + vstore8(luma, 0, out_y.ptr + luma_output_stride_y); + + uchar4 cb = convert_uchar4((cb_0 + cb_1) / (ushort4)(2)); + uchar4 cr = convert_uchar4((cr_0 + cr_1) / (ushort4)(2)); + vstore4(cb, 0, out_u.ptr); + vstore4(cr, 0, out_v.ptr); +} + +/** Convert an IYUV image to RGB888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] u_input_ptr Pointer to the source U channel. Supported Format: U8 + * @param[in] u_input_stride_x Stride of the source image U channel in X dimension (in bytes) + * @param[in] u_input_step_x u_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] u_input_step_y u_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_input_offset_first_element_in_bytes The offset of the first element in the source U channel + * @param[in] v_input_ptr Pointer to the source V channel. Supported Format: U8 + * @param[in] v_input_stride_x Stride of the source image V channel in X dimension (in bytes) + * @param[in] v_input_step_x v_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_input_stride_y Stride of the source image V channel in Y dimension (in bytes) + * @param[in] v_input_step_y v_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_input_offset_first_element_in_bytes The offset of the first element in the source image V channel + * @param[out] rgb_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgb_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgb_output_step_x rgb_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgb_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgb_output_step_y rgb_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgb_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void IYUV_to_RGB888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(u_input), + IMAGE_DECLARATION(v_input), + IMAGE_DECLARATION(rgb_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_u = CONVERT_TO_IMAGE_STRUCT(u_input); + Image in_v = CONVERT_TO_IMAGE_STRUCT(v_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgb_output); + + // handle 8 pixels every time, two lines, each line for 4 pixels + uchar4 luma_0 = vload4(0, in_y.ptr); + uchar4 luma_1 = vload4(0, in_y.ptr + luma_input_stride_y); + uchar4 cbcr = (uchar4)(vload2(0, in_u.ptr), vload2(0, in_v.ptr)); + char4 cb = (char4)(cbcr.s0, cbcr.s0, cbcr.s1, cbcr.s1) - (char4)(128); + char4 cr = (char4)(cbcr.s2, cbcr.s2, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + uchar4 rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore4(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, r_0.s1, g_0.s1, b_0.s1, r_0.s2, g_0.s2); + rgb_1 = (uchar4)(b_0.s2, r_0.s3, g_0.s3, b_0.s3); + vstore8(rgb_0, 0, out_rgb.ptr + rgb_output_stride_y); + vstore4(rgb_1, 0, out_rgb.ptr + rgb_output_stride_y + 8); +} + +/** Convert an IYUV image to RGB8888 + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] u_input_ptr Pointer to the source U channel. Supported Format: U8 + * @param[in] u_input_stride_x Stride of the source image U channel in X dimension (in bytes) + * @param[in] u_input_step_x u_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] u_input_step_y u_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_input_offset_first_element_in_bytes The offset of the first element in the source U channel + * @param[in] v_input_ptr Pointer to the source V channel. Supported Format: U8 + * @param[in] v_input_stride_x Stride of the source image V channel in X dimension (in bytes) + * @param[in] v_input_step_x v_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_input_stride_y Stride of the source image V channel in Y dimension (in bytes) + * @param[in] v_input_step_y v_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_input_offset_first_element_in_bytes The offset of the first element in the source image V channel + * @param[out] rgba_output_ptr Pointer to the destination image. Supported Format: U8 + * @param[in] rgba_output_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] rgba_output_step_x rgba_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rgba_output_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] rgba_output_step_y rgba_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] rgba_output_offset_first_element_in_bytes The offset of the first element in the destination image + */ +__kernel void IYUV_to_RGBA8888_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(u_input), + IMAGE_DECLARATION(v_input), + IMAGE_DECLARATION(rgba_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_u = CONVERT_TO_IMAGE_STRUCT(u_input); + Image in_v = CONVERT_TO_IMAGE_STRUCT(v_input); + Image out_rgb = CONVERT_TO_IMAGE_STRUCT(rgba_output); + + // handle 8 pixels every time, two lines, each line for 4 pixels + uchar4 luma_0 = vload4(0, in_y.ptr); + uchar4 luma_1 = vload4(0, in_y.ptr + luma_input_stride_y); + uchar4 cbcr = (uchar4)(vload2(0, in_u.ptr), vload2(0, in_v.ptr)); + char4 cb = (char4)(cbcr.s0, cbcr.s0, cbcr.s1, cbcr.s1) - (char4)(128); + char4 cr = (char4)(cbcr.s2, cbcr.s2, cbcr.s3, cbcr.s3) - (char4)(128); + + float4 temp0 = (float4)(0.0000f) + (float4)(0.0000f) * convert_float4(cb) + (float4)(1.5748f) * convert_float4(cr); + float4 temp1 = (float4)(0.0000f) - (float4)(0.1873f) * convert_float4(cb) - (float4)(0.4681f) * convert_float4(cr); + float4 temp2 = (float4)(0.0000f) + (float4)(1.8556f) * convert_float4(cb) + (float4)(0.0000f) * convert_float4(cr); + + float4 f_r = convert_float4(luma_0) + temp0; + float4 f_g = convert_float4(luma_0) + temp1; + float4 f_b = convert_float4(luma_0) + temp2; + + uchar4 r_0 = convert_uchar4_rtz(f_r); + uchar4 g_0 = convert_uchar4_rtz(f_g); + uchar4 b_0 = convert_uchar4_rtz(f_b); + + uchar8 rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + uchar8 rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr); + vstore8(rgb_1, 0, out_rgb.ptr + 8); + + f_r = convert_float4(luma_1) + temp0; + f_g = convert_float4(luma_1) + temp1; + f_b = convert_float4(luma_1) + temp2; + + r_0 = convert_uchar4_rtz(f_r); + g_0 = convert_uchar4_rtz(f_g); + b_0 = convert_uchar4_rtz(f_b); + + rgb_0 = (uchar8)(r_0.s0, g_0.s0, b_0.s0, 255, r_0.s1, g_0.s1, b_0.s1, 255); + rgb_1 = (uchar8)(r_0.s2, g_0.s2, b_0.s2, 255, r_0.s3, g_0.s3, b_0.s3, 255); + vstore8(rgb_0, 0, out_rgb.ptr + rgba_output_stride_y); + vstore8(rgb_1, 0, out_rgb.ptr + rgba_output_stride_y + 8); +} + +/** Convert an IYUV image to YUV444 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] u_input_ptr Pointer to the source U channel. Supported Format: U8 + * @param[in] u_input_stride_x Stride of the source image U channel in X dimension (in bytes) + * @param[in] u_input_step_x u_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] u_input_step_y u_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_input_offset_first_element_in_bytes The offset of the first element in the source U channel + * @param[in] v_input_ptr Pointer to the source V channel. Supported Format: U8 + * @param[in] v_input_stride_x Stride of the source image V channel in X dimension (in bytes) + * @param[in] v_input_step_x v_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_input_stride_y Stride of the source image V channel in Y dimension (in bytes) + * @param[in] v_input_step_y v_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_input_offset_first_element_in_bytes The offset of the first element in the source image V channel + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] u_output_ptr Pointer to the destination U channel. Supported Format: U8 + * @param[in] u_output_stride_x Stride of the destination U channel in X dimension (in bytes) + * @param[in] u_output_step_x u_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] u_output_step_y u_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_output_offset_first_element_in_bytes The offset of the first element in the destination U channel + * @param[out] v_output_ptr Pointer to the destination V channel. Supported Format: U8 + * @param[in] v_output_stride_x Stride of the destination V channel in X dimension (in bytes) + * @param[in] v_output_step_x v_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_output_stride_y Stride of the destination V channel in Y dimension (in bytes) + * @param[in] v_output_step_y v_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_output_offset_first_element_in_bytes The offset of the first element in the destination V channel + * + */ +__kernel void IYUV_to_YUV444_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(u_input), + IMAGE_DECLARATION(v_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(u_output), + IMAGE_DECLARATION(v_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_u = CONVERT_TO_IMAGE_STRUCT(u_input); + Image in_v = CONVERT_TO_IMAGE_STRUCT(v_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_u = CONVERT_TO_IMAGE_STRUCT(u_output); + Image out_v = CONVERT_TO_IMAGE_STRUCT(v_output); + + // handle 32 pixels every time, two lines, each line for 16 pixels + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar8 cb_src = vload8(0, in_u.ptr); + uchar8 cr_src = vload8(0, in_v.ptr); + uchar16 cb = (uchar16)(cb_src.s0, cb_src.s0, cb_src.s1, cb_src.s1, cb_src.s2, cb_src.s2, cb_src.s3, cb_src.s3, + cb_src.s4, cb_src.s4, cb_src.s5, cb_src.s5, cb_src.s6, cb_src.s6, cb_src.s7, cb_src.s7); + uchar16 cr = (uchar16)(cr_src.s0, cr_src.s0, cr_src.s1, cr_src.s1, cr_src.s2, cr_src.s2, cr_src.s3, cr_src.s3, + cr_src.s4, cr_src.s4, cr_src.s5, cr_src.s5, cr_src.s6, cr_src.s6, cr_src.s7, cr_src.s7); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore16(cb, 0, out_u.ptr); + vstore16(cb, 0, out_u.ptr + u_output_stride_y); + vstore16(cr, 0, out_v.ptr); + vstore16(cr, 0, out_v.ptr + v_output_stride_y); +} + +/** Convert an IYUV image to NV12 + * + * Global Workgroup Size [ DIV_CEIL(width, 16), height ] + * No offset. + * + * @param[in] luma_input_ptr Pointer to the source luma channel. Supported Format: U8 + * @param[in] luma_input_stride_x Stride of the luma image in X dimension (in bytes) + * @param[in] luma_input_step_x luma_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_input_stride_y Stride of the source luma channel in Y dimension (in bytes) + * @param[in] luma_input_step_y luma_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] u_input_ptr Pointer to the source U channel. Supported Format: U8 + * @param[in] u_input_stride_x Stride of the source image U channel in X dimension (in bytes) + * @param[in] u_input_step_x u_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] u_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] u_input_step_y u_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] u_input_offset_first_element_in_bytes The offset of the first element in the source U channel + * @param[in] v_input_ptr Pointer to the source V channel. Supported Format: U8 + * @param[in] v_input_stride_x Stride of the source image V channel in X dimension (in bytes) + * @param[in] v_input_step_x v_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] v_input_stride_y Stride of the source image V channel in Y dimension (in bytes) + * @param[in] v_input_step_y v_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] v_input_offset_first_element_in_bytes The offset of the first element in the source image V channel + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] uv_output_ptr Pointer to the destination UV channel. Supported Format: U8 + * @param[in] uv_output_stride_x Stride of the destination UV channel in X dimension (in bytes) + * @param[in] uv_output_step_x uv_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_output_stride_y Stride of the destination image U channel in Y dimension (in bytes) + * @param[in] uv_output_step_y uv_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_output_offset_first_element_in_bytes The offset of the first element in the destination UV channel + * + */ +__kernel void IYUV_to_NV12_bt709( + IMAGE_DECLARATION(luma_input), + IMAGE_DECLARATION(u_input), + IMAGE_DECLARATION(v_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(uv_output)) +{ + Image in_y = CONVERT_TO_IMAGE_STRUCT(luma_input); + Image in_u = CONVERT_TO_IMAGE_STRUCT(u_input); + Image in_v = CONVERT_TO_IMAGE_STRUCT(v_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_uv = CONVERT_TO_IMAGE_STRUCT(uv_output); + + // handle 32 pixels every time, two lines, each line for 16 pixels + uchar16 luma_0 = vload16(0, in_y.ptr); + uchar16 luma_1 = vload16(0, in_y.ptr + luma_input_stride_y); + uchar8 cb = vload8(0, in_u.ptr); + uchar8 cr = vload8(0, in_v.ptr); + uchar16 cbcr = (uchar16)(cb.s0, cr.s0, cb.s1, cr.s1, cb.s2, cr.s2, cb.s3, cr.s3, cb.s4, cr.s4, cb.s5, cr.s5, cb.s6, + cr.s6, cb.s7, cr.s7); + + vstore16(luma_0, 0, out_y.ptr); + vstore16(luma_1, 0, out_y.ptr + luma_output_stride_y); + vstore16(cbcr, 0, out_uv.ptr); +} + +/** Convert a YUYV image to NV12 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 8), height ] + * No offset. + * + * @param[in] yuyv_input_ptr Pointer to the source image. Supported Format: U8 + * @param[in] yuyv_input_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] yuyv_input_step_x yuyv_input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] yuyv_input_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] yuyv_input_step_y yuyv_input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] yuyv_input_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_output_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_output_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_output_step_x luma_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_output_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_output_step_y luma_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_output_offset_first_element_in_bytes The offset of the first element in the destination luma channel + * @param[out] uv_output_ptr Pointer to the destination UV channel. Supported Format: U8 + * @param[in] uv_output_stride_x Stride of the destination UV channel in X dimension (in bytes) + * @param[in] uv_output_step_x uv_output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_output_stride_y Stride of the destination image UV channel in Y dimension (in bytes) + * @param[in] uv_output_step_y uv_output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_output_offset_first_element_in_bytes The offset of the first element in the destination UV channel + * + */ +__kernel void YUYV422_to_NV12_bt709( + IMAGE_DECLARATION(yuyv_input), + IMAGE_DECLARATION(luma_output), + IMAGE_DECLARATION(uv_output)) +{ + Image in_yuyv = CONVERT_TO_IMAGE_STRUCT(yuyv_input); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma_output); + Image out_uv = CONVERT_TO_IMAGE_STRUCT(uv_output); + + // handle 16 pixels every time, each line 8 pixels + uchar16 yuyv = vload16(0, in_yuyv.ptr); + ushort8 cbcr_0 = (ushort8)(yuyv.s1, yuyv.s3, yuyv.s5, yuyv.s7, yuyv.s9, yuyv.sb, yuyv.sd, yuyv.sf); + uchar8 luma = (uchar8)(yuyv.s0, yuyv.s2, yuyv.s4, yuyv.s6, yuyv.s8, yuyv.sa, yuyv.sc, yuyv.se); + vstore8(luma, 0, out_y.ptr); + + yuyv = vload16(0, in_yuyv.ptr + yuyv_input_stride_y); + ushort8 cbcr_1 = (ushort8)(yuyv.s1, yuyv.s3, yuyv.s5, yuyv.s7, yuyv.s9, yuyv.sb, yuyv.sd, yuyv.sf); + luma = (uchar8)(yuyv.s0, yuyv.s2, yuyv.s4, yuyv.s6, yuyv.s8, yuyv.sa, yuyv.sc, yuyv.se); + vstore8(luma, 0, out_y.ptr + luma_output_stride_y); + + uchar8 cbcr = convert_uchar8((cbcr_0 + cbcr_1) / (ushort8)(2)); + vstore8(cbcr, 0, out_uv.ptr); +} + +/** Convert a UYVY image to NV12 using BT709 color space + * + * Global Workgroup Size [ DIV_CEIL(width, 4), height ] + * No offset. + * + * @param[in] input_uyvy_ptr Pointer to the source image. Supported Format: U8 + * @param[in] input_uyvy_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] input_uyvy_step_x input_uyvy_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_uyvy_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] input_uyvy_step_y input_uyvy_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_uyvy_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] luma_ptr Pointer to the destination luma channel. Supported Format: U8 + * @param[in] luma_stride_x Stride of the destination luma channel in X dimension (in bytes) + * @param[in] luma_step_x luma_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] luma_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] luma_step_y luma_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] luma_offset_first_element_in_bytes The offset of the first element in the destination image luma channel + * @param[out] uv_ptr Pointer to the destination uv channel. Supported Format: U8 + * @param[in] uv_stride_x Stride of the destination uv channel in X dimension (in bytes) + * @param[in] uv_step_x uv_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] uv_stride_y Stride of the destination image luma channel in Y dimension (in bytes) + * @param[in] uv_step_y uv_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] uv_offset_first_element_in_bytes The offset of the first element in the destination image uv channel + * + */ +__kernel void UYVY422_to_NV12_bt709( + IMAGE_DECLARATION(input_uyvy), + IMAGE_DECLARATION(luma), + IMAGE_DECLARATION(uv)) +{ + Image in = CONVERT_TO_IMAGE_STRUCT(input_uyvy); + Image out_y = CONVERT_TO_IMAGE_STRUCT(luma); + Image out_uv = CONVERT_TO_IMAGE_STRUCT(uv); + + // handle 16 pixels every time, each line 8 pixels + const uchar16 uyvy_t = vload16(0, in.ptr); + vstore8(uyvy_t.s13579bdf, 0, out_y.ptr); + + const uchar16 uyvy_b = vload16(0, in.ptr + input_uyvy_stride_y); + vstore8(uyvy_b.s13579bdf, 0, out_y.ptr + luma_stride_y); + + const ushort8 cbcr_t = (ushort8)(uyvy_t.s0, uyvy_t.s2, uyvy_t.s4, uyvy_t.s6, uyvy_t.s8, uyvy_t.sa, uyvy_t.sc, uyvy_t.se); + const ushort8 cbcr_b = (ushort8)(uyvy_b.s0, uyvy_b.s2, uyvy_b.s4, uyvy_b.s6, uyvy_b.s8, uyvy_b.sa, uyvy_b.sc, uyvy_b.se); + const uchar8 cbcr = convert_uchar8((cbcr_t + cbcr_b) / (ushort8)(2)); + vstore8(cbcr, 0, out_uv.ptr); +} -- cgit v1.2.1