diff options
Diffstat (limited to 'src/core/CL/cl_kernels/color_convert.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/color_convert.cl | 1911 |
1 files changed, 0 insertions, 1911 deletions
diff --git a/src/core/CL/cl_kernels/color_convert.cl b/src/core/CL/cl_kernels/color_convert.cl deleted file mode 100644 index 7a872b47b5..0000000000 --- a/src/core/CL/cl_kernels/color_convert.cl +++ /dev/null @@ -1,1911 +0,0 @@ -/* - * Copyright (c) 2016-2018 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.sf, 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 U8 - * - * Global Workgroup Size [ DIV_CEIL(width, 16), height ] - * No offset. - * - * @param[in] input_ptr Pointer to the source image. Supported Format: RGB888 - * @param[in] input_stride_x Stride of the source image in X dimension (in bytes) - * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes) - * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image - * @param[out] output_ptr Pointer to the destination image. Supported Format: U8 - * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes) - * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes) - * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image - */ -__kernel void RGB888_to_U8_bt709( - IMAGE_DECLARATION(input), - IMAGE_DECLARATION(output)) -{ - Image in = CONVERT_TO_IMAGE_STRUCT(input); - Image out = CONVERT_TO_IMAGE_STRUCT(output); - - // handle 16 pixels every time - const uchar16 rgb_0 = vload16(0, in.ptr); - const uchar16 rgb_1 = vload16(0, in.ptr + 16); - const uchar16 rgb_2 = vload16(0, in.ptr + 32); - - //Resequence values from a sequence of 16 RGB values to sequence of 16 R, 16 G, 16 B values - const uchar16 rgb_r = (uchar16)(rgb_0.s0369, rgb_0.scf, rgb_1.s258b, rgb_1.se, rgb_2.s147a, rgb_2.sd); - const uchar16 rgb_g = (uchar16)(rgb_0.s147a, rgb_0.sd, rgb_1.s0369, rgb_1.scf, rgb_2.s258b, rgb_2.se); - const uchar16 rgb_b = (uchar16)(rgb_0.s258b, rgb_0.se, rgb_1.s147a, rgb_1.sd, rgb_2.s0369, rgb_2.scf); - - const float16 rgb2u8_red_coef_bt709 = 0.2126f; - const float16 rgb2u8_green_coef_bt709 = 0.7152f; - const float16 rgb2u8_blue_coef_bt709 = 0.0722f; - - //Computation of 16 greyscale values in float - const float16 greyscale_f_0 = rgb2u8_red_coef_bt709 * convert_float16(rgb_r) + rgb2u8_green_coef_bt709 * convert_float16(rgb_g) + rgb2u8_blue_coef_bt709 * convert_float16(rgb_b); - - //Convert it to 16 grayscale uchar values - const uchar16 greyscale_u8_0 = convert_uchar16_sat_rtz(greyscale_f_0); - - vstore16(greyscale_u8_0, 0, out.ptr); -} - -/** Convert an RGB888 image to RGBX8888 - * - * Global Workgroup Size [ DIV_CEIL(width, 16), height ] - * 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 red_coef_bt709 = (float8)(1.5748f); - float8 green_coef_bt709 = (float8)(-0.1873f); - float8 green_coef2_bt709 = (float8)(-0.4681f); - float8 blue_coef_bt709 = (float8)(1.8556f); - float8 lumav = convert_float8(luma); - - float8 f_r = red_coef_bt709 * convert_float8(cr); - float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); - float8 f_b = blue_coef_bt709 * convert_float8(cb); - - f_r += lumav; - f_g += lumav; - f_b += lumav; - - uchar8 r_0 = convert_uchar8_sat_rtz(f_r); - uchar8 g_0 = convert_uchar8_sat_rtz(f_g); - uchar8 b_0 = convert_uchar8_sat_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 red_coef_bt709 = (float8)(1.5748f); - float8 green_coef_bt709 = (float8)(-0.1873f); - float8 green_coef2_bt709 = (float8)(-0.4681f); - float8 blue_coef_bt709 = (float8)(1.8556f); - float8 lumav = convert_float8(luma); - - float8 f_r = red_coef_bt709 * convert_float8(cr); - float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); - float8 f_b = blue_coef_bt709 * convert_float8(cb); - - f_r += lumav; - f_g += lumav; - f_b += lumav; - - uchar8 r_0 = convert_uchar8_sat_rtz(f_r); - uchar8 g_0 = convert_uchar8_sat_rtz(f_g); - uchar8 b_0 = convert_uchar8_sat_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 red_coef_bt709 = (float8)(1.5748f); - float8 green_coef_bt709 = (float8)(-0.1873f); - float8 green_coef2_bt709 = (float8)(-0.4681f); - float8 blue_coef_bt709 = (float8)(1.8556f); - float8 lumav = convert_float8(luma); - - float8 f_r = red_coef_bt709 * convert_float8(cr); - float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); - float8 f_b = blue_coef_bt709 * convert_float8(cb); - - f_r += lumav; - f_g += lumav; - f_b += lumav; - - uchar8 r_0 = convert_uchar8_sat_rtz(f_r); - uchar8 g_0 = convert_uchar8_sat_rtz(f_g); - uchar8 b_0 = convert_uchar8_sat_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 red_coef_bt709 = (float8)(1.5748f); - float8 green_coef_bt709 = (float8)(-0.1873f); - float8 green_coef2_bt709 = (float8)(-0.4681f); - float8 blue_coef_bt709 = (float8)(1.8556f); - float8 lumav = convert_float8(luma); - - float8 f_r = red_coef_bt709 * convert_float8(cr); - float8 f_g = green_coef_bt709 * convert_float8(cb) + green_coef2_bt709 * convert_float8(cr); - float8 f_b = blue_coef_bt709 * convert_float8(cb); - - f_r += lumav; - f_g += lumav; - f_b += lumav; - - uchar8 r_0 = convert_uchar8_sat_rtz(f_r); - uchar8 g_0 = convert_uchar8_sat_rtz(f_g); - uchar8 b_0 = convert_uchar8_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - uchar4 g_0 = convert_uchar4_sat_rtz(f_g); - uchar4 b_0 = convert_uchar4_sat_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_sat_rtz(f_r); - g_0 = convert_uchar4_sat_rtz(f_g); - b_0 = convert_uchar4_sat_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); -} |