/* GStreamer * Copyright (C) 2022 Seungha Yang * * This library is free software; you can redistribute it and/or * modify it under the terms of the GNU Library General Public * License as published by the Free Software Foundation; either * version 2 of the License, or (at your option) any later version. * * This library is distributed in the hope that it will be useful, * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU * Library General Public License for more details. * * You should have received a copy of the GNU Library General Public * License along with this library; if not, write to the * Free Software Foundation, Inc., 51 Franklin St, Fifth Floor, * Boston, MA 02110-1301, USA. */ #ifdef HAVE_CONFIG_H #include "config.h" #endif #include "gstcudaconverter.h" #include GST_DEBUG_CATEGORY_STATIC (gst_cuda_converter_debug); #define GST_CAT_DEFAULT gst_cuda_converter_debug #define CUDA_BLOCK_X 16 #define CUDA_BLOCK_Y 16 #define DIV_UP(size,block) (((size) + ((block) - 1)) / (block)) /* from GstD3D11 */ typedef struct _GstCudaColorMatrix { gdouble matrix[3][3]; gdouble offset[3]; gdouble min[3]; gdouble max[3]; } GstCudaColorMatrix; static gchar * gst_cuda_dump_color_matrix (GstCudaColorMatrix * matrix) { /* *INDENT-OFF* */ static const gchar format[] = "[MATRIX]\n" "|% .6f, % .6f, % .6f|\n" "|% .6f, % .6f, % .6f|\n" "|% .6f, % .6f, % .6f|\n" "[OFFSET]\n" "|% .6f, % .6f, % .6f|\n" "[MIN]\n" "|% .6f, % .6f, % .6f|\n" "[MAX]\n" "|% .6f, % .6f, % .6f|"; /* *INDENT-ON* */ return g_strdup_printf (format, matrix->matrix[0][0], matrix->matrix[0][1], matrix->matrix[0][2], matrix->matrix[1][0], matrix->matrix[1][1], matrix->matrix[1][2], matrix->matrix[2][0], matrix->matrix[2][1], matrix->matrix[2][2], matrix->offset[0], matrix->offset[1], matrix->offset[2], matrix->min[0], matrix->min[1], matrix->min[2], matrix->max[0], matrix->max[1], matrix->max[2]); } static void color_matrix_copy (GstCudaColorMatrix * dst, const GstCudaColorMatrix * src) { for (guint i = 0; i < 3; i++) { for (guint j = 0; j < 3; j++) { dst->matrix[i][j] = src->matrix[i][j]; } } } static void color_matrix_multiply (GstCudaColorMatrix * dst, GstCudaColorMatrix * a, GstCudaColorMatrix * b) { GstCudaColorMatrix tmp; for (guint i = 0; i < 3; i++) { for (guint j = 0; j < 3; j++) { gdouble val = 0; for (guint k = 0; k < 3; k++) { val += a->matrix[i][k] * b->matrix[k][j]; } tmp.matrix[i][j] = val; } } color_matrix_copy (dst, &tmp); } static void color_matrix_identity (GstCudaColorMatrix * m) { for (guint i = 0; i < 3; i++) { for (guint j = 0; j < 3; j++) { if (i == j) m->matrix[i][j] = 1.0; else m->matrix[i][j] = 0; } } } /** * gst_cuda_color_range_adjust_matrix_unorm: * @in_info: a #GstVideoInfo * @out_info: a #GstVideoInfo * @matrix: a #GstCudaColorMatrix * * Calculates matrix for color range adjustment. Both input and output * signals are in normalized [0.0..1.0] space. * * Resulting values can be calculated by * | Yout | | Yin | | matrix.offset[0] | * | Uout | = clamp ( matrix.matrix * | Uin | + | matrix.offset[1] |, matrix.min, matrix.max ) * | Vout | | Vin | | matrix.offset[2] | * * Returns: %TRUE if successful */ static gboolean gst_cuda_color_range_adjust_matrix_unorm (const GstVideoInfo * in_info, const GstVideoInfo * out_info, GstCudaColorMatrix * matrix) { gboolean in_rgb, out_rgb; gint in_offset[GST_VIDEO_MAX_COMPONENTS]; gint in_scale[GST_VIDEO_MAX_COMPONENTS]; gint out_offset[GST_VIDEO_MAX_COMPONENTS]; gint out_scale[GST_VIDEO_MAX_COMPONENTS]; GstVideoColorRange in_range; GstVideoColorRange out_range; gdouble src_fullscale, dst_fullscale; memset (matrix, 0, sizeof (GstCudaColorMatrix)); for (guint i = 0; i < 3; i++) { matrix->matrix[i][i] = 1.0; matrix->matrix[i][i] = 1.0; matrix->matrix[i][i] = 1.0; matrix->max[i] = 1.0; } in_rgb = GST_VIDEO_INFO_IS_RGB (in_info); out_rgb = GST_VIDEO_INFO_IS_RGB (out_info); if (in_rgb != out_rgb) { GST_WARNING ("Invalid format conversion"); return FALSE; } in_range = in_info->colorimetry.range; out_range = out_info->colorimetry.range; if (in_range == GST_VIDEO_COLOR_RANGE_UNKNOWN) { GST_WARNING ("Unknown input color range"); if (in_rgb || GST_VIDEO_INFO_IS_GRAY (in_info)) in_range = GST_VIDEO_COLOR_RANGE_0_255; else in_range = GST_VIDEO_COLOR_RANGE_16_235; } if (out_range == GST_VIDEO_COLOR_RANGE_UNKNOWN) { GST_WARNING ("Unknown output color range"); if (out_rgb || GST_VIDEO_INFO_IS_GRAY (out_info)) out_range = GST_VIDEO_COLOR_RANGE_0_255; else out_range = GST_VIDEO_COLOR_RANGE_16_235; } src_fullscale = (gdouble) ((1 << in_info->finfo->depth[0]) - 1); dst_fullscale = (gdouble) ((1 << out_info->finfo->depth[0]) - 1); gst_video_color_range_offsets (in_range, in_info->finfo, in_offset, in_scale); gst_video_color_range_offsets (out_range, out_info->finfo, out_offset, out_scale); matrix->min[0] = matrix->min[1] = matrix->min[2] = (gdouble) out_offset[0] / dst_fullscale; matrix->max[0] = (out_scale[0] + out_offset[0]) / dst_fullscale; matrix->max[1] = matrix->max[2] = (out_scale[1] + out_offset[0]) / dst_fullscale; if (in_info->colorimetry.range == out_info->colorimetry.range) { GST_DEBUG ("Same color range"); return TRUE; } /* Formula * * 1) Scales and offset compensates input to [0..1] range * SRC_NORM[i] = (src[i] * src_fullscale - in_offset[i]) / in_scale[i] * = (src[i] * src_fullscale / in_scale[i]) - in_offset[i] / in_scale[i] * * 2) Reverse to output UNIT scale * DST_UINT[i] = SRC_NORM[i] * out_scale[i] + out_offset[i] * = src[i] * src_fullscale * out_scale[i] / in_scale[i] * - in_offset[i] * out_scale[i] / in_scale[i] * + out_offset[i] * * 3) Back to [0..1] scale * dst[i] = DST_UINT[i] / dst_fullscale * = COEFF[i] * src[i] + OFF[i] * where * src_fullscale * out_scale[i] * COEFF[i] = ------------------------------ * dst_fullscale * in_scale[i] * * out_offset[i] in_offset[i] * out_scale[i] * OFF[i] = -------------- - ------------------------------ * dst_fullscale dst_fullscale * in_scale[i] */ for (guint i = 0; i < 3; i++) { matrix->matrix[i][i] = (src_fullscale * out_scale[i]) / (dst_fullscale * in_scale[i]); matrix->offset[i] = (out_offset[i] / dst_fullscale) - ((gdouble) in_offset[i] * out_scale[i] / (dst_fullscale * in_scale[i])); } return TRUE; } /** * gst_cuda_yuv_to_rgb_matrix_unorm: * @in_yuv_info: a #GstVideoInfo of input YUV signal * @out_rgb_info: a #GstVideoInfo of output RGB signal * @matrix: a #GstCudaColorMatrix * * Calculates transform matrix from YUV to RGB conversion. Both input and output * signals are in normalized [0.0..1.0] space and additional gamma decoding * or primary/transfer function transform is not performed by this matrix. * * Resulting non-linear RGB values can be calculated by * | R' | | Y' | | matrix.offset[0] | * | G' | = clamp ( matrix.matrix * | Cb | + | matrix.offset[1] | matrix.min, matrix.max ) * | B' | | Cr | | matrix.offset[2] | * * Returns: %TRUE if successful */ static gboolean gst_cuda_yuv_to_rgb_matrix_unorm (const GstVideoInfo * in_yuv_info, const GstVideoInfo * out_rgb_info, GstCudaColorMatrix * matrix) { gint offset[4], scale[4]; gdouble Kr, Kb, Kg; /* * * * Input: Unsigned normalized Y'CbCr(unorm), [0.0..1.0] range * Output: Unsigned normalized non-linear R'G'B'(unorm), [0.0..1.0] range * * 1) Y'CbCr(unorm) to scaled Y'CbCr * | Y' | | Y'(unorm) | * | Cb | = S | Cb(unorm) | * | Cb | | Cr(unorm) | * where S = (2 ^ bitdepth) - 1 * * 2) Y'CbCr to YPbPr * Y = (Y' - offsetY ) / scaleY * Pb = [(Cb - offsetCbCr) / scaleCbCr] * Pr = [(Cr - offsetCrCr) / scaleCrCr] * => * Y = Y'(unorm) * Sy + Oy * Pb = Cb(unorm) * Suv + Ouv * Pb = Cr(unorm) * Suv + Ouv * where * Sy = S / scaleY * Suv = S / scaleCbCr * Oy = -(offsetY / scaleY) * Ouv = -(offsetCbCr / scaleCbCr) * * 3) YPbPr to R'G'B' * | R' | | Y | * | G' | = M *| Pb | * | B' | | Pr | * where * | vecR | * M = | vecG | * | vecB | * vecR = | 1, 0 , 2(1 - Kr) | * vecG = | 1, -(Kb/Kg) * 2(1 - Kb), -(Kr/Kg) * 2(1 - Kr) | * vecB = | 1, 2(1 - Kb) , 0 | * => * R' = dot(vecR, (Syuv * Y'CbCr(unorm))) + dot(vecR, Offset) * G' = dot(vecG, (Svuy * Y'CbCr(unorm))) + dot(vecG, Offset) * B' = dot(vecB, (Syuv * Y'CbCr(unorm)) + dot(vecB, Offset) * where * | Sy, 0, 0 | * Syuv = | 0, Suv, 0 | * | 0 0, Suv | * * | Oy | * Offset = | Ouv | * | Ouv | * * 4) YUV -> RGB matrix * | R' | | Y'(unorm) | | offsetA | * | G' | = Matrix * | Cb(unorm) | + | offsetB | * | B' | | Cr(unorm) | | offsetC | * * where * | vecR | * Matrix = | vecG | * Syuv * | vecB | * * offsetA = dot(vecR, Offset) * offsetB = dot(vecG, Offset) * offsetC = dot(vecB, Offset) * * 4) Consider 16-235 scale RGB * RGBfull(0..255) -> RGBfull(16..235) matrix is represented by * | Rs | | Rf | | Or | * | Gs | = Ms | Gf | + | Og | * | Bs | | Bf | | Ob | * * Combining all matrix into * | Rs | | Y'(unorm) | | offsetA | | Or | * | Gs | = Ms * ( Matrix * | Cb(unorm) | + | offsetB | ) + | Og | * | Bs | | Cr(unorm) | | offsetC | | Ob | * * | Y'(unorm) | | offsetA | | Or | * = Ms * Matrix * | Cb(unorm) | + Ms | offsetB | + | Og | * | Cr(unorm) | | offsetC | | Ob | */ memset (matrix, 0, sizeof (GstCudaColorMatrix)); for (guint i = 0; i < 3; i++) matrix->max[i] = 1.0; gst_video_color_range_offsets (in_yuv_info->colorimetry.range, in_yuv_info->finfo, offset, scale); if (gst_video_color_matrix_get_Kr_Kb (in_yuv_info->colorimetry.matrix, &Kr, &Kb)) { guint S; gdouble Sy, Suv; gdouble Oy, Ouv; gdouble vecR[3], vecG[3], vecB[3]; Kg = 1.0 - Kr - Kb; vecR[0] = 1.0; vecR[1] = 0; vecR[2] = 2 * (1 - Kr); vecG[0] = 1.0; vecG[1] = -(Kb / Kg) * 2 * (1 - Kb); vecG[2] = -(Kr / Kg) * 2 * (1 - Kr); vecB[0] = 1.0; vecB[1] = 2 * (1 - Kb); vecB[2] = 0; /* Assume all components has the same bitdepth */ S = (1 << in_yuv_info->finfo->depth[0]) - 1; Sy = (gdouble) S / scale[0]; Suv = (gdouble) S / scale[1]; Oy = -((gdouble) offset[0] / scale[0]); Ouv = -((gdouble) offset[1] / scale[1]); matrix->matrix[0][0] = Sy * vecR[0]; matrix->matrix[1][0] = Sy * vecG[0]; matrix->matrix[2][0] = Sy * vecB[0]; matrix->matrix[0][1] = Suv * vecR[1]; matrix->matrix[1][1] = Suv * vecG[1]; matrix->matrix[2][1] = Suv * vecB[1]; matrix->matrix[0][2] = Suv * vecR[2]; matrix->matrix[1][2] = Suv * vecG[2]; matrix->matrix[2][2] = Suv * vecB[2]; matrix->offset[0] = vecR[0] * Oy + vecR[1] * Ouv + vecR[2] * Ouv; matrix->offset[1] = vecG[0] * Oy + vecG[1] * Ouv + vecG[2] * Ouv; matrix->offset[2] = vecB[0] * Oy + vecB[1] * Ouv + vecB[2] * Ouv; /* Apply RGB range scale matrix */ if (out_rgb_info->colorimetry.range == GST_VIDEO_COLOR_RANGE_16_235) { GstCudaColorMatrix scale_matrix, rst; GstVideoInfo full_rgb = *out_rgb_info; full_rgb.colorimetry.range = GST_VIDEO_COLOR_RANGE_0_255; if (gst_cuda_color_range_adjust_matrix_unorm (&full_rgb, out_rgb_info, &scale_matrix)) { /* Ms * Matrix */ color_matrix_multiply (&rst, &scale_matrix, matrix); /* Ms * transform offsets */ for (guint i = 0; i < 3; i++) { gdouble val = 0; for (guint j = 0; j < 3; j++) { val += scale_matrix.matrix[i][j] * matrix->offset[j]; } rst.offset[i] = val + scale_matrix.offset[i]; } /* copy back to output matrix */ for (guint i = 0; i < 3; i++) { for (guint j = 0; j < 3; j++) { matrix->matrix[i][j] = rst.matrix[i][j]; } matrix->offset[i] = rst.offset[i]; matrix->min[i] = scale_matrix.min[i]; matrix->max[i] = scale_matrix.max[i]; } } } } else { /* Unknown matrix */ matrix->matrix[0][0] = 1.0; matrix->matrix[1][1] = 1.0; matrix->matrix[2][2] = 1.0; } return TRUE; } /** * gst_cuda_rgb_to_yuv_matrix_unorm: * @in_rgb_info: a #GstVideoInfo of input RGB signal * @out_yuv_info: a #GstVideoInfo of output YUV signal * @matrix: a #GstCudaColorMatrix * * Calculates transform matrix from RGB to YUV conversion. Both input and output * signals are in normalized [0.0..1.0] space and additional gamma decoding * or primary/transfer function transform is not performed by this matrix. * * Resulting RGB values can be calculated by * | Y' | | R' | | matrix.offset[0] | * | Cb | = clamp ( matrix.matrix * | G' | + | matrix.offset[1] |, matrix.min, matrix.max ) * | Cr | | B' | | matrix.offset[2] | * * Returns: %TRUE if successful */ static gboolean gst_cuda_rgb_to_yuv_matrix_unorm (const GstVideoInfo * in_rgb_info, const GstVideoInfo * out_yuv_info, GstCudaColorMatrix * matrix) { gint offset[4], scale[4]; gdouble Kr, Kb, Kg; /* * * * Input: Unsigned normalized non-linear R'G'B'(unorm), [0.0..1.0] range * Output: Unsigned normalized Y'CbCr(unorm), [0.0..1.0] range * * 1) R'G'B' to YPbPr * | Y | | R' | * | Pb | = M *| G' | * | Pr | | B' | * where * | vecY | * M = | vecU | * | vecV | * vecY = | Kr , Kg , Kb | * vecU = | -0.5*Kr/(1-Kb), -0.5*Kg/(1-Kb), 0.5 | * vecV = | 0.5 , -0.5*Kg/(1-Kr), -0.5*Kb(1-Kr) | * * 2) YPbPr to Y'CbCr(unorm) * Y'(unorm) = (Y * scaleY + offsetY) / S * Cb(unorm) = (Pb * scaleCbCr + offsetCbCr) / S * Cr(unorm) = (Pr * scaleCbCr + offsetCbCr) / S * => * Y'(unorm) = (Y * scaleY / S) + (offsetY / S) * Cb(unorm) = (Pb * scaleCbCr / S) + (offsetCbCr / S) * Cr(unorm) = (Pb * scaleCbCr / S) + (offsetCbCr / S) * where S = (2 ^ bitdepth) - 1 * * 3) RGB -> YUV matrix * | Y'(unorm) | | R' | | offsetA | * | Cb(unorm) | = Matrix * | G' | + | offsetB | * | Cr(unorm) | | B' | | offsetC | * * where * | (scaleY/S) * vecY | * Matrix = | (scaleCbCr/S) * vecU | * | (scaleCbCr/S) * vecV | * * offsetA = offsetY / S * offsetB = offsetCbCr / S * offsetC = offsetCbCr / S * * 4) Consider 16-235 scale RGB * RGBstudio(16..235) -> RGBfull(0..255) matrix is represented by * | Rf | | Rs | | Or | * | Gf | = Ms | Gs | + | Og | * | Bf | | Bs | | Ob | * * Combining all matrix into * | Y'(unorm) | | Rs | | Or | | offsetA | * | Cb(unorm) | = Matrix * ( Ms | Gs | + | Og | ) + | offsetB | * | Cr(unorm) | | Bs | | Ob | | offsetC | * * | Rs | | Or | | offsetA | * = Matrix * Ms | Gs | + Matrix | Og | + | offsetB | * | Bs | | Ob | | offsetB | */ memset (matrix, 0, sizeof (GstCudaColorMatrix)); for (guint i = 0; i < 3; i++) matrix->max[i] = 1.0; gst_video_color_range_offsets (out_yuv_info->colorimetry.range, out_yuv_info->finfo, offset, scale); if (gst_video_color_matrix_get_Kr_Kb (out_yuv_info->colorimetry.matrix, &Kr, &Kb)) { guint S; gdouble Sy, Suv; gdouble Oy, Ouv; gdouble vecY[3], vecU[3], vecV[3]; Kg = 1.0 - Kr - Kb; vecY[0] = Kr; vecY[1] = Kg; vecY[2] = Kb; vecU[0] = -0.5 * Kr / (1 - Kb); vecU[1] = -0.5 * Kg / (1 - Kb); vecU[2] = 0.5; vecV[0] = 0.5; vecV[1] = -0.5 * Kg / (1 - Kr); vecV[2] = -0.5 * Kb / (1 - Kr); /* Assume all components has the same bitdepth */ S = (1 << out_yuv_info->finfo->depth[0]) - 1; Sy = (gdouble) scale[0] / S; Suv = (gdouble) scale[1] / S; Oy = (gdouble) offset[0] / S; Ouv = (gdouble) offset[1] / S; for (guint i = 0; i < 3; i++) { matrix->matrix[0][i] = Sy * vecY[i]; matrix->matrix[1][i] = Suv * vecU[i]; matrix->matrix[2][i] = Suv * vecV[i]; } matrix->offset[0] = Oy; matrix->offset[1] = Ouv; matrix->offset[2] = Ouv; matrix->min[0] = Oy; matrix->min[1] = Oy; matrix->min[2] = Oy; matrix->max[0] = ((gdouble) scale[0] + offset[0]) / S; matrix->max[1] = ((gdouble) scale[1] + offset[0]) / S; matrix->max[2] = ((gdouble) scale[1] + offset[0]) / S; /* Apply RGB range scale matrix */ if (in_rgb_info->colorimetry.range == GST_VIDEO_COLOR_RANGE_16_235) { GstCudaColorMatrix scale_matrix, rst; GstVideoInfo full_rgb = *in_rgb_info; full_rgb.colorimetry.range = GST_VIDEO_COLOR_RANGE_0_255; if (gst_cuda_color_range_adjust_matrix_unorm (in_rgb_info, &full_rgb, &scale_matrix)) { /* Matrix * Ms */ color_matrix_multiply (&rst, matrix, &scale_matrix); /* Matrix * scale offsets */ for (guint i = 0; i < 3; i++) { gdouble val = 0; for (guint j = 0; j < 3; j++) { val += matrix->matrix[i][j] * scale_matrix.offset[j]; } rst.offset[i] = val + matrix->offset[i]; } /* copy back to output matrix */ for (guint i = 0; i < 3; i++) { for (guint j = 0; j < 3; j++) { matrix->matrix[i][j] = rst.matrix[i][j]; } matrix->offset[i] = rst.offset[i]; } } } } else { /* Unknown matrix */ matrix->matrix[0][0] = 1.0; matrix->matrix[1][1] = 1.0; matrix->matrix[2][2] = 1.0; } return TRUE; } typedef struct { float coeffX[3]; float coeffY[3]; float coeffZ[3]; float offset[3]; float min[3]; float max[3]; } ColorMatrix; typedef struct { ColorMatrix toRGBCoeff; ColorMatrix toYuvCoeff; ColorMatrix primariesCoeff; } ConstBuffer; #define COLOR_SPACE_IDENTITY "color_space_identity" #define COLOR_SPACE_CONVERT "color_space_convert" #define SAMPLE_YUV_PLANAR "sample_yuv_planar" #define SAMPLE_YV12 "sample_yv12" #define SAMPLE_YUV_PLANAR_10BIS "sample_yuv_planar_10bits" #define SAMPLE_YUV_PLANAR_12BIS "sample_yuv_planar_12bits" #define SAMPLE_SEMI_PLANAR "sample_semi_planar" #define SAMPLE_SEMI_PLANAR_SWAP "sample_semi_planar_swap" #define SAMPLE_RGBA "sample_rgba" #define SAMPLE_BGRA "sample_bgra" #define SAMPLE_RGBx "sample_rgbx" #define SAMPLE_BGRx "sample_bgrx" #define SAMPLE_ARGB "sample_argb" /* same as ARGB */ #define SAMPLE_ARGB64 "sample_argb" #define SAMPLE_AGBR "sample_abgr" #define SAMPLE_RGBP "sample_rgbp" #define SAMPLE_BGRP "sample_bgrp" #define SAMPLE_GBR "sample_gbr" #define SAMPLE_GBRA "sample_gbra" #define WRITE_I420 "write_i420" #define WRITE_YV12 "write_yv12" #define WRITE_NV12 "write_nv12" #define WRITE_NV21 "write_nv21" #define WRITE_P010 "write_p010" /* same as P010 */ #define WRITE_P016 "write_p010" #define WRITE_I420_10 "write_i420_10" #define WRITE_Y444 "write_y444" #define WRITE_Y444_16 "write_y444_16" #define WRITE_RGBA "write_rgba" #define WRITE_RGBx "write_rgbx" #define WRITE_BGRA "write_bgra" #define WRITE_BGRx "write_bgrx" #define WRITE_ARGB "write_argb" #define WRITE_ABGR "write_abgr" #define WRITE_RGB "write_rgb" #define WRITE_BGR "write_bgr" #define WRITE_RGB10A2 "write_rgb10a2" #define WRITE_BGR10A2 "write_bgr10a2" #define WRITE_Y42B "write_y42b" #define WRITE_I422_10 "write_i422_10" #define WRITE_I422_12 "write_i422_12" #define WRITE_RGBP "write_rgbp" #define WRITE_BGRP "write_bgrp" #define WRITE_GBR "write_gbr" #define WRITE_GBRA "write_gbra" /* *INDENT-OFF* */ const static gchar KERNEL_COMMON[] = "struct ColorMatrix\n" "{\n" " float CoeffX[3];\n" " float CoeffY[3];\n" " float CoeffZ[3];\n" " float Offset[3];\n" " float Min[3];\n" " float Max[3];\n" "};\n" "\n" "__device__ inline float\n" "dot (const float coeff[3], float3 val)\n" "{\n" " return coeff[0] * val.x + coeff[1] * val.y + coeff[2] * val.z;\n" "}\n" "\n" "__device__ inline float\n" "clamp (float val, float min_val, float max_val)\n" "{\n" " return max (min_val, min (val, max_val));\n" "}\n" "\n" "__device__ inline float3\n" "clamp3 (float3 val, const float min_val[3], const float max_val[3])\n" "{\n" " return make_float3 (clamp (val.x, min_val[0], max_val[0]),\n" " clamp (val.y, min_val[1], max_val[2]),\n" " clamp (val.z, min_val[1], max_val[2]));\n" "}\n" "\n" "__device__ inline unsigned char\n" "scale_to_2bits (float val)\n" "{\n" " return (unsigned short) __float2int_rz (val * 3.0);\n" "}\n" "\n" "__device__ inline unsigned char\n" "scale_to_uchar (float val)\n" "{\n" " return (unsigned char) __float2int_rz (val * 255.0);\n" "}\n" "\n" "__device__ inline unsigned short\n" "scale_to_ushort (float val)\n" "{\n" " return (unsigned short) __float2int_rz (val * 65535.0);\n" "}\n" "\n" "__device__ inline unsigned short\n" "scale_to_10bits (float val)\n" "{\n" " return (unsigned short) __float2int_rz (val * 1023.0);\n" "}\n" "\n" "__device__ inline unsigned short\n" "scale_to_12bits (float val)\n" "{\n" " return (unsigned short) __float2int_rz (val * 4095.0);\n" "}\n" "\n" "__device__ inline float3\n" COLOR_SPACE_IDENTITY "(float3 sample, const ColorMatrix * matrix)\n" "{\n" " return sample;\n" "}\n" "\n" "__device__ inline float3\n" COLOR_SPACE_CONVERT "(float3 sample, const ColorMatrix * matrix)\n" "{\n" " float3 out;\n" " out.x = dot (matrix->CoeffX, sample);\n" " out.y = dot (matrix->CoeffY, sample);\n" " out.z = dot (matrix->CoeffZ, sample);\n" " out.x += matrix->Offset[0];\n" " out.y += matrix->Offset[1];\n" " out.z += matrix->Offset[2];\n" " return clamp3 (out, matrix->Min, matrix->Max);\n" "}\n" "/* All 8bits yuv planar except for yv12 */\n" "__device__ inline float4\n" SAMPLE_YUV_PLANAR "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float u = tex2D(tex1, x, y);\n" " float v = tex2D(tex2, x, y);\n" " return make_float4 (luma, u, v, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_YV12 "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float u = tex2D(tex2, x, y);\n" " float v = tex2D(tex1, x, y);\n" " return make_float4 (luma, u, v, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_YUV_PLANAR_10BIS "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float u = tex2D(tex1, x, y);\n" " float v = tex2D(tex2, x, y);\n" " /* (1 << 6) to scale [0, 1.0) range */\n" " return make_float4 (luma * 64, u * 64, v * 64, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_YUV_PLANAR_12BIS "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float u = tex2D(tex1, x, y);\n" " float v = tex2D(tex2, x, y);\n" " /* (1 << 6) to scale [0, 1.0) range */\n" " return make_float4 (luma * 16, u * 16, v * 16, 1);\n" "}\n" "\n" "/* NV12, P010, and P016 */\n" "__device__ inline float4\n" SAMPLE_SEMI_PLANAR "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float2 uv = tex2D(tex1, x, y);\n" " return make_float4 (luma, uv.x, uv.y, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_SEMI_PLANAR_SWAP "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float luma = tex2D(tex0, x, y);\n" " float2 vu = tex2D(tex1, x, y);\n" " return make_float4 (luma, vu.y, vu.x, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_RGBA "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " return tex2D(tex0, x, y);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_BGRA "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float4 bgra = tex2D(tex0, x, y);\n" " return make_float4 (bgra.z, bgra.y, bgra.x, bgra.w);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_RGBx "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float4 rgbx = tex2D(tex0, x, y);\n" " rgbx.w = 1;\n" " return rgbx;\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_BGRx "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float4 bgrx = tex2D(tex0, x, y);\n" " return make_float4 (bgrx.z, bgrx.y, bgrx.x, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_ARGB "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float4 argb = tex2D(tex0, x, y);\n" " return make_float4 (argb.y, argb.z, argb.w, argb.x);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_AGBR "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float4 abgr = tex2D(tex0, x, y);\n" " return make_float4 (abgr.w, abgr.z, abgr.y, abgr.x);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_RGBP "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float r = tex2D(tex0, x, y);\n" " float g = tex2D(tex1, x, y);\n" " float b = tex2D(tex2, x, y);\n" " return make_float4 (r, g, b, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_BGRP "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float b = tex2D(tex0, x, y);\n" " float g = tex2D(tex1, x, y);\n" " float r = tex2D(tex2, x, y);\n" " return make_float4 (r, g, b, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_GBR "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float g = tex2D(tex0, x, y);\n" " float b = tex2D(tex1, x, y);\n" " float r = tex2D(tex2, x, y);\n" " return make_float4 (r, g, b, 1);\n" "}\n" "\n" "__device__ inline float4\n" SAMPLE_GBRA "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n" "{\n" " float g = tex2D(tex0, x, y);\n" " float b = tex2D(tex1, x, y);\n" " float r = tex2D(tex2, x, y);\n" " float a = tex2D(tex3, x, y);\n" " return make_float4 (r, g, b, a);\n" "}\n" "\n" "__device__ inline void\n" WRITE_I420 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " dst0[x + y * stride0] = scale_to_uchar (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x / 2 + (y / 2) * stride1;\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst2[pos] = scale_to_uchar (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_YV12 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " dst0[x + y * stride0] = scale_to_uchar (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x / 2 + (y / 2) * stride1;\n" " dst1[pos] = scale_to_uchar (sample.z);\n" " dst2[pos] = scale_to_uchar (sample.y);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_NV12 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " dst0[x + y * stride0] = scale_to_uchar (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x + (y / 2) * stride1;\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst1[pos + 1] = scale_to_uchar (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_NV21 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " dst0[x + y * stride0] = scale_to_uchar (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x + (y / 2) * stride1;\n" " dst1[pos] = scale_to_uchar (sample.z);\n" " dst1[pos + 1] = scale_to_uchar (sample.y);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_P010 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " *(unsigned short *) &dst0[x * 2 + y * stride0] = scale_to_ushort (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x * 2 + (y / 2) * stride1;\n" " *(unsigned short *) &dst1[pos] = scale_to_ushort (sample.y);\n" " *(unsigned short *) &dst1[pos + 2] = scale_to_ushort (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_I420_10 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " *(unsigned short *) &dst0[x * 2 + y * stride0] = scale_to_10bits (sample.x);\n" " if (x % 2 == 0 && y % 2 == 0) {\n" " unsigned int pos = x + (y / 2) * stride1;\n" " *(unsigned short *) &dst1[pos] = scale_to_10bits (sample.y);\n" " *(unsigned short *) &dst2[pos] = scale_to_10bits (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_Y444 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.x);\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst2[pos] = scale_to_uchar (sample.z);\n" "}\n" "\n" "__device__ inline void\n" WRITE_Y444_16 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 2 + y * stride0;\n" " *(unsigned short *) &dst0[pos] = scale_to_ushort (sample.x);\n" " *(unsigned short *) &dst1[pos] = scale_to_ushort (sample.y);\n" " *(unsigned short *) &dst2[pos] = scale_to_ushort (sample.z);\n" "}\n" "\n" "__device__ inline void\n" WRITE_RGBA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.x);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.z);\n" " dst0[pos + 3] = scale_to_uchar (sample.w);\n" "}\n" "\n" "__device__ inline void\n" WRITE_RGBx "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.x);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.z);\n" " dst0[pos + 3] = 255;\n" "}\n" "\n" "__device__ inline void\n" WRITE_BGRA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.z);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.x);\n" " dst0[pos + 3] = scale_to_uchar (sample.w);\n" "}\n" "\n" "__device__ inline void\n" WRITE_BGRx "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.z);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.x);\n" " dst0[pos + 3] = 255;\n" "}\n" "\n" "__device__ inline void\n" WRITE_ARGB "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.w);\n" " dst0[pos + 1] = scale_to_uchar (sample.x);\n" " dst0[pos + 2] = scale_to_uchar (sample.y);\n" " dst0[pos + 3] = scale_to_uchar (sample.z);\n" "}\n" "\n" "__device__ inline void\n" WRITE_ABGR "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 4 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.w);\n" " dst0[pos + 1] = scale_to_uchar (sample.z);\n" " dst0[pos + 2] = scale_to_uchar (sample.y);\n" " dst0[pos + 3] = scale_to_uchar (sample.x);\n" "}\n" "\n" "__device__ inline void\n" WRITE_RGB "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 3 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.x);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.z);\n" "}\n" "\n" "__device__ inline void\n" WRITE_BGR "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x * 3 + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.z);\n" " dst0[pos + 1] = scale_to_uchar (sample.y);\n" " dst0[pos + 2] = scale_to_uchar (sample.x);\n" "}\n" "\n" "__device__ inline void\n" WRITE_RGB10A2 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " unsigned int alpha = (unsigned int) scale_to_2bits (sample.x);\n" " unsigned int packed_rgb = alpha << 30;\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.x));\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.y)) << 10;\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.z)) << 20;\n" " *(unsigned int *) &dst0[x * 4 + y * stride0] = packed_rgb;\n" "}\n" "\n" "__device__ inline void\n" WRITE_BGR10A2 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " unsigned int alpha = (unsigned int) scale_to_2bits (sample.x);\n" " unsigned int packed_rgb = alpha << 30;\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.x)) << 20;\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.y)) << 10;\n" " packed_rgb |= ((unsigned int) scale_to_10bits (sample.z));\n" " *(unsigned int *) &dst0[x * 4 + y * stride0] = packed_rgb;\n" "}\n" "\n" "__device__ inline void\n" WRITE_Y42B "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " dst0[x + y * stride0] = scale_to_uchar (sample.x);\n" " if (x % 2 == 0) {\n" " unsigned int pos = x / 2 + y * stride1;\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst2[pos] = scale_to_uchar (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_I422_10 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " *(unsigned short *) &dst0[x * 2 + y * stride0] = scale_to_10bits (sample.x);\n" " if (x % 2 == 0) {\n" " unsigned int pos = x + y * stride1;\n" " *(unsigned short *) &dst1[pos] = scale_to_10bits (sample.y);\n" " *(unsigned short *) &dst2[pos] = scale_to_10bits (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_I422_12 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " *(unsigned short *) &dst0[x * 2 + y * stride0] = scale_to_12bits (sample.x);\n" " if (x % 2 == 0) {\n" " unsigned int pos = x + y * stride1;\n" " *(unsigned short *) &dst1[pos] = scale_to_12bits (sample.y);\n" " *(unsigned short *) &dst2[pos] = scale_to_12bits (sample.z);\n" " }\n" "}\n" "\n" "__device__ inline void\n" WRITE_RGBP "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.x);\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst2[pos] = scale_to_uchar (sample.z);\n" "}\n" "\n" "__device__ inline void\n" WRITE_BGRP "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.z);\n" " dst1[pos] = scale_to_uchar (sample.y);\n" " dst2[pos] = scale_to_uchar (sample.x);\n" "}\n" "\n" "__device__ inline void\n" WRITE_GBR "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.y);\n" " dst1[pos] = scale_to_uchar (sample.z);\n" " dst2[pos] = scale_to_uchar (sample.x);\n" "}\n" "\n" "__device__ inline void\n" WRITE_GBRA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" "{\n" " int pos = x + y * stride0;\n" " dst0[pos] = scale_to_uchar (sample.y);\n" " dst1[pos] = scale_to_uchar (sample.z);\n" " dst2[pos] = scale_to_uchar (sample.x);\n" " dst3[pos] = scale_to_uchar (sample.w);\n" "}\n"; #define GST_CUDA_KERNEL_UNPACK_FUNC "gst_cuda_kernel_unpack_func" static const gchar RGB_TO_RGBx[] = "extern \"C\" {\n" "__global__ void\n" GST_CUDA_KERNEL_UNPACK_FUNC "(unsigned char *src, unsigned char *dst, int width, int height,\n" " int src_stride, int dst_stride)\n" "{\n" " int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n" " int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n" " if (x_pos < width && y_pos < height) {\n" " int dst_pos = x_pos * 4 + y_pos * dst_stride;\n" " int src_pos = x_pos * 3 + y_pos * src_stride;\n" " dst[dst_pos] = src[src_pos];\n" " dst[dst_pos + 1] = src[src_pos + 1];\n" " dst[dst_pos + 2] = src[src_pos + 2];\n" " dst[dst_pos + 3] = 0xff;\n" " }\n" "}\n" "}\n"; static const gchar RGB10A2_TO_ARGB64[] = "extern \"C\" {\n" "__global__ void\n" GST_CUDA_KERNEL_UNPACK_FUNC "(unsigned char *src, unsigned char *dst, int width, int height,\n" " int src_stride, int dst_stride)\n" "{\n" " int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n" " int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n" " if (x_pos < width && y_pos < height) {\n" " unsigned short a, r, g, b;\n" " unsigned int val;\n" " int dst_pos = x_pos * 8 + y_pos * dst_stride;\n" " val = *(unsigned int *)&src[x_pos * 4 + y_pos * src_stride];\n" " a = (val >> 30) & 0x03;\n" " a = (a << 14) | (a << 12) | (a << 10) | (a << 8) | (a << 6) | (a << 4) | (a << 2) | (a << 0);\n" " r = (val & 0x3ff);\n" " r = (r << 6) | (r >> 4);\n" " g = ((val >> 10) & 0x3ff);\n" " g = (g << 6) | (g >> 4);\n" " b = ((val >> 20) & 0x3ff);\n" " b = (b << 6) | (b >> 4);\n" " *(unsigned short *) &dst[dst_pos] = a;\n" " *(unsigned short *) &dst[dst_pos + 2] = r;\n" " *(unsigned short *) &dst[dst_pos + 4] = g;\n" " *(unsigned short *) &dst[dst_pos + 6] = b;\n" " }\n" "}\n" "}\n"; static const gchar BGR10A2_TO_ARGB64[] = "extern \"C\" {\n" "__global__ void\n" GST_CUDA_KERNEL_UNPACK_FUNC "(unsigned char *src, unsigned char *dst, int width, int height,\n" " int src_stride, int dst_stride)\n" "{\n" " int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n" " int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n" " if (x_pos < width && y_pos < height) {\n" " unsigned short a, r, g, b;\n" " unsigned int val;\n" " int dst_pos = x_pos * 8 + y_pos * dst_stride;\n" " val = *(unsigned int *)&src[x_pos * 4 + y_pos * src_stride];\n" " a = (val >> 30) & 0x03;\n" " a = (a << 14) | (a << 12) | (a << 10) | (a << 8) | (a << 6) | (a << 4) | (a << 2) | (a << 0);\n" " b = (val & 0x3ff);\n" " b = (b << 6) | (b >> 4);\n" " g = ((val >> 10) & 0x3ff);\n" " g = (g << 6) | (g >> 4);\n" " r = ((val >> 20) & 0x3ff);\n" " r = (r << 6) | (r >> 4);\n" " *(unsigned short *) &dst[dst_pos] = a;\n" " *(unsigned short *) &dst[dst_pos + 2] = r;\n" " *(unsigned short *) &dst[dst_pos + 4] = g;\n" " *(unsigned short *) &dst[dst_pos + 6] = b;\n" " }\n" "}\n" "}\n"; #define GST_CUDA_KERNEL_MAIN_FUNC "gst_cuda_converter_main" static const gchar TEMPLETA_KERNEL[] = /* KERNEL_COMMON */ "%s\n" /* UNPACK FUNCTION */ "%s\n" "__constant__ ColorMatrix TO_RGB_MATRIX = { { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s } };\n" "__constant__ ColorMatrix TO_YUV_MATRIX = { { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s },\n" " { %s, %s, %s } };\n" "__constant__ int WIDTH = %d;\n" "__constant__ int HEIGHT = %d;\n" "__constant__ int LEFT = %d;\n" "__constant__ int TOP = %d;\n" "__constant__ int RIGHT = %d;\n" "__constant__ int BOTTOM = %d;\n" "__constant__ int VIEW_WIDTH = %d;\n" "__constant__ int VIEW_HEIGHT = %d;\n" "__constant__ float OFFSET_X = %s;\n" "__constant__ float OFFSET_Y = %s;\n" "__constant__ float BORDER_X = %s;\n" "__constant__ float BORDER_Y = %s;\n" "__constant__ float BORDER_Z = %s;\n" "__constant__ float BORDER_W = %s;\n" "\n" "extern \"C\" {\n" "__global__ void\n" GST_CUDA_KERNEL_MAIN_FUNC "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n" " cudaTextureObject_t tex2, cudaTextureObject_t tex3, unsigned char * dst0,\n" " unsigned char * dst1, unsigned char * dst2, unsigned char * dst3,\n" " int stride0, int stride1)\n" "{\n" " int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n" " int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n" " float4 sample;\n" " if (x_pos >= WIDTH || y_pos >= HEIGHT)\n" " return;\n" " if (x_pos < LEFT || x_pos >= RIGHT || y_pos < TOP || y_pos >= BOTTOM) {\n" " sample = make_float4 (BORDER_X, BORDER_Y, BORDER_Z, BORDER_W);\n" " } else {\n" " float x = OFFSET_X + (float) (x_pos - LEFT) / VIEW_WIDTH;\n" " float y = OFFSET_Y + (float) (y_pos - TOP) / VIEW_HEIGHT;\n" " float4 s = %s (tex0, tex1, tex2, tex3, x, y);\n" " float3 xyz = make_float3 (s.x, s.y, s.z);\n" " float3 rgb = %s (xyz, &TO_RGB_MATRIX);\n" " float3 yuv = %s (rgb, &TO_YUV_MATRIX);\n" " sample = make_float4 (yuv.x, yuv.y, yuv.z, s.w);\n" " }\n" " %s (dst0, dst1, dst2, dst3, sample, x_pos, y_pos, stride0, stride1);\n" "}\n" "}\n"; /* *INDENT-ON* */ typedef struct _TextureFormat { GstVideoFormat format; CUarray_format array_format[GST_VIDEO_MAX_COMPONENTS]; guint channels[GST_VIDEO_MAX_COMPONENTS]; const gchar *sample_func; } TextureFormat; #define CU_AD_FORMAT_NONE 0 #define MAKE_FORMAT_YUV_PLANAR(f,cf,sample_func) \ { GST_VIDEO_FORMAT_ ##f, { CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_ ##cf, \ CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_NONE }, {1, 1, 1, 0}, sample_func } #define MAKE_FORMAT_YUV_SEMI_PLANAR(f,cf,sample_func) \ { GST_VIDEO_FORMAT_ ##f, { CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_ ##cf, \ CU_AD_FORMAT_NONE, CU_AD_FORMAT_NONE }, {1, 2, 0, 0}, sample_func } #define MAKE_FORMAT_RGB(f,cf,sample_func) \ { GST_VIDEO_FORMAT_ ##f, { CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_NONE, \ CU_AD_FORMAT_NONE, CU_AD_FORMAT_NONE }, {4, 0, 0, 0}, sample_func } #define MAKE_FORMAT_RGBP(f,cf,sample_func) \ { GST_VIDEO_FORMAT_ ##f, { CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_ ##cf, \ CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_NONE }, {1, 1, 1, 0}, sample_func } #define MAKE_FORMAT_RGBAP(f,cf,sample_func) \ { GST_VIDEO_FORMAT_ ##f, { CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_ ##cf, \ CU_AD_FORMAT_ ##cf, CU_AD_FORMAT_ ##cf }, {1, 1, 1, 1}, sample_func } static const TextureFormat format_map[] = { MAKE_FORMAT_YUV_PLANAR (I420, UNSIGNED_INT8, SAMPLE_YUV_PLANAR), MAKE_FORMAT_YUV_PLANAR (YV12, UNSIGNED_INT8, SAMPLE_YV12), MAKE_FORMAT_YUV_SEMI_PLANAR (NV12, UNSIGNED_INT8, SAMPLE_SEMI_PLANAR), MAKE_FORMAT_YUV_SEMI_PLANAR (NV21, UNSIGNED_INT8, SAMPLE_SEMI_PLANAR_SWAP), MAKE_FORMAT_YUV_SEMI_PLANAR (P010_10LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR), MAKE_FORMAT_YUV_SEMI_PLANAR (P016_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR), MAKE_FORMAT_YUV_PLANAR (I420_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS), MAKE_FORMAT_YUV_PLANAR (Y444, UNSIGNED_INT8, SAMPLE_YUV_PLANAR), MAKE_FORMAT_YUV_PLANAR (Y444_16LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR), MAKE_FORMAT_RGB (RGBA, UNSIGNED_INT8, SAMPLE_RGBA), MAKE_FORMAT_RGB (BGRA, UNSIGNED_INT8, SAMPLE_BGRA), MAKE_FORMAT_RGB (RGBx, UNSIGNED_INT8, SAMPLE_RGBx), MAKE_FORMAT_RGB (BGRx, UNSIGNED_INT8, SAMPLE_BGRx), MAKE_FORMAT_RGB (ARGB, UNSIGNED_INT8, SAMPLE_ARGB), MAKE_FORMAT_RGB (ARGB64, UNSIGNED_INT16, SAMPLE_ARGB64), MAKE_FORMAT_RGB (ABGR, UNSIGNED_INT8, SAMPLE_AGBR), MAKE_FORMAT_YUV_PLANAR (Y42B, UNSIGNED_INT8, SAMPLE_YUV_PLANAR), MAKE_FORMAT_YUV_PLANAR (I422_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS), MAKE_FORMAT_YUV_PLANAR (I422_12LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_12BIS), MAKE_FORMAT_RGBP (RGBP, UNSIGNED_INT8, SAMPLE_RGBP), MAKE_FORMAT_RGBP (BGRP, UNSIGNED_INT8, SAMPLE_BGRP), MAKE_FORMAT_RGBP (GBR, UNSIGNED_INT8, SAMPLE_GBR), MAKE_FORMAT_RGBAP (GBRA, UNSIGNED_INT8, SAMPLE_GBRA), }; typedef struct _TextureBuffer { CUdeviceptr ptr; gsize stride; CUtexObject texture; } TextureBuffer; typedef struct { gint x; gint y; gint width; gint height; } ConverterRect; struct _GstCudaConverterPrivate { GstVideoInfo in_info; GstVideoInfo out_info; GstStructure *config; GstVideoInfo texture_info; const TextureFormat *texture_fmt; gint texture_align; ConverterRect dest_rect; TextureBuffer fallback_buffer[GST_VIDEO_MAX_COMPONENTS]; CUfilter_mode filter_mode[GST_VIDEO_MAX_COMPONENTS]; TextureBuffer unpack_buffer; CUmodule module; CUfunction main_func; CUfunction unpack_func; }; static void gst_cuda_converter_dispose (GObject * object); static void gst_cuda_converter_finalize (GObject * object); #define gst_cuda_converter_parent_class parent_class G_DEFINE_TYPE_WITH_PRIVATE (GstCudaConverter, gst_cuda_converter, GST_TYPE_OBJECT); static void gst_cuda_converter_class_init (GstCudaConverterClass * klass) { GObjectClass *object_class = G_OBJECT_CLASS (klass); object_class->dispose = gst_cuda_converter_dispose; object_class->finalize = gst_cuda_converter_finalize; GST_DEBUG_CATEGORY_INIT (gst_cuda_converter_debug, "cudaconverter", 0, "cudaconverter"); } static void gst_cuda_converter_init (GstCudaConverter * self) { GstCudaConverterPrivate *priv; self->priv = priv = gst_cuda_converter_get_instance_private (self); priv->config = gst_structure_new_empty ("GstCudaConverter"); } static void gst_cuda_converter_dispose (GObject * object) { GstCudaConverter *self = GST_CUDA_CONVERTER (object); GstCudaConverterPrivate *priv = self->priv; guint i; if (self->context && gst_cuda_context_push (self->context)) { if (priv->module) { CuModuleUnload (priv->module); priv->module = NULL; } for (i = 0; i < G_N_ELEMENTS (priv->fallback_buffer); i++) { if (priv->fallback_buffer[i].ptr) { if (priv->fallback_buffer[i].texture) { CuTexObjectDestroy (priv->fallback_buffer[i].texture); priv->fallback_buffer[i].texture = 0; } CuMemFree (priv->fallback_buffer[i].ptr); priv->fallback_buffer[i].ptr = 0; } } if (priv->unpack_buffer.ptr) { if (priv->unpack_buffer.texture) { CuTexObjectDestroy (priv->unpack_buffer.texture); priv->unpack_buffer.texture = 0; } CuMemFree (priv->unpack_buffer.ptr); priv->unpack_buffer.ptr = 0; } gst_cuda_context_pop (NULL); } gst_clear_object (&self->context); G_OBJECT_CLASS (parent_class)->dispose (object); } static void gst_cuda_converter_finalize (GObject * object) { GstCudaConverter *self = GST_CUDA_CONVERTER (object); GstCudaConverterPrivate *priv = self->priv; gst_structure_free (priv->config); G_OBJECT_CLASS (parent_class)->finalize (object); } static const gchar * get_color_range_name (GstVideoColorRange range) { switch (range) { case GST_VIDEO_COLOR_RANGE_0_255: return "FULL"; case GST_VIDEO_COLOR_RANGE_16_235: return "STUDIO"; default: break; } return "UNKNOWN"; } typedef struct _GstCudaColorMatrixString { gchar matrix[3][3][G_ASCII_DTOSTR_BUF_SIZE]; gchar offset[3][G_ASCII_DTOSTR_BUF_SIZE]; gchar min[3][G_ASCII_DTOSTR_BUF_SIZE]; gchar max[3][G_ASCII_DTOSTR_BUF_SIZE]; } GstCudaColorMatrixString; static void color_matrix_to_string (const GstCudaColorMatrix * m, GstCudaColorMatrixString * str) { guint i, j; for (i = 0; i < 3; i++) { for (j = 0; j < 3; j++) { g_ascii_formatd (str->matrix[i][j], G_ASCII_DTOSTR_BUF_SIZE, "%f", m->matrix[i][j]); } g_ascii_formatd (str->offset[i], G_ASCII_DTOSTR_BUF_SIZE, "%f", m->offset[i]); g_ascii_formatd (str->min[i], G_ASCII_DTOSTR_BUF_SIZE, "%f", m->min[i]); g_ascii_formatd (str->max[i], G_ASCII_DTOSTR_BUF_SIZE, "%f", m->max[i]); } } static gboolean gst_cuda_converter_setup (GstCudaConverter * self) { GstCudaConverterPrivate *priv = self->priv; const GstVideoInfo *in_info; const GstVideoInfo *out_info; const GstVideoInfo *texture_info; GstCudaColorMatrix to_rgb_matrix; GstCudaColorMatrix to_yuv_matrix; GstCudaColorMatrix border_color_matrix; GstCudaColorMatrixString to_rgb_matrix_str; GstCudaColorMatrixString to_yuv_matrix_str; gchar border_color_str[4][G_ASCII_DTOSTR_BUF_SIZE]; gdouble border_color[4]; gchar offset_x[G_ASCII_DTOSTR_BUF_SIZE]; gchar offset_y[G_ASCII_DTOSTR_BUF_SIZE]; gint i, j; const gchar *unpack_function = NULL; const gchar *write_func = NULL; const gchar *to_rgb_func = COLOR_SPACE_IDENTITY; const gchar *to_yuv_func = COLOR_SPACE_IDENTITY; const GstVideoColorimetry *in_color; const GstVideoColorimetry *out_color; gchar *str; gchar *ptx; CUresult ret; in_info = &priv->in_info; out_info = &priv->out_info; texture_info = &priv->texture_info; in_color = &in_info->colorimetry; out_color = &out_info->colorimetry; memset (&to_rgb_matrix, 0, sizeof (GstCudaColorMatrix)); color_matrix_identity (&to_rgb_matrix); memset (&to_yuv_matrix, 0, sizeof (GstCudaColorMatrix)); color_matrix_identity (&to_yuv_matrix); switch (GST_VIDEO_INFO_FORMAT (out_info)) { case GST_VIDEO_FORMAT_I420: write_func = WRITE_I420; break; case GST_VIDEO_FORMAT_YV12: write_func = WRITE_YV12; break; case GST_VIDEO_FORMAT_NV12: write_func = WRITE_NV12; break; case GST_VIDEO_FORMAT_NV21: write_func = WRITE_NV21; break; case GST_VIDEO_FORMAT_P010_10LE: write_func = WRITE_P010; break; case GST_VIDEO_FORMAT_P016_LE: write_func = WRITE_P016; break; case GST_VIDEO_FORMAT_I420_10LE: write_func = WRITE_I420_10; break; case GST_VIDEO_FORMAT_Y444: write_func = WRITE_Y444; break; case GST_VIDEO_FORMAT_Y444_16LE: write_func = WRITE_Y444_16; break; case GST_VIDEO_FORMAT_RGBA: write_func = WRITE_RGBA; break; case GST_VIDEO_FORMAT_RGBx: write_func = WRITE_RGBx; break; case GST_VIDEO_FORMAT_BGRA: write_func = WRITE_BGRA; break; case GST_VIDEO_FORMAT_BGRx: write_func = WRITE_BGRx; break; case GST_VIDEO_FORMAT_ARGB: write_func = WRITE_ARGB; break; case GST_VIDEO_FORMAT_ABGR: write_func = WRITE_ABGR; break; case GST_VIDEO_FORMAT_RGB: write_func = WRITE_RGB; break; case GST_VIDEO_FORMAT_BGR: write_func = WRITE_BGR; break; case GST_VIDEO_FORMAT_RGB10A2_LE: write_func = WRITE_RGB10A2; break; case GST_VIDEO_FORMAT_BGR10A2_LE: write_func = WRITE_BGR10A2; break; case GST_VIDEO_FORMAT_Y42B: write_func = WRITE_Y42B; break; case GST_VIDEO_FORMAT_I422_10LE: write_func = WRITE_I422_10; break; case GST_VIDEO_FORMAT_I422_12LE: write_func = WRITE_I422_12; break; case GST_VIDEO_FORMAT_RGBP: write_func = WRITE_RGBP; break; case GST_VIDEO_FORMAT_BGRP: write_func = WRITE_BGRP; break; case GST_VIDEO_FORMAT_GBR: write_func = WRITE_GBR; break; case GST_VIDEO_FORMAT_GBRA: write_func = WRITE_GBRA; break; default: break; } if (!write_func) { GST_ERROR_OBJECT (self, "Unknown write function for format %s", gst_video_format_to_string (GST_VIDEO_INFO_FORMAT (out_info))); return FALSE; } /* Decide texture info to use, 3 channel RGB or 10bits packed RGB * need be converted to other format */ priv->texture_info = priv->in_info; switch (GST_VIDEO_INFO_FORMAT (in_info)) { case GST_VIDEO_FORMAT_RGB: gst_video_info_set_format (&priv->texture_info, GST_VIDEO_FORMAT_RGBx, GST_VIDEO_INFO_WIDTH (in_info), GST_VIDEO_INFO_HEIGHT (in_info)); unpack_function = RGB_TO_RGBx; break; case GST_VIDEO_FORMAT_BGR: gst_video_info_set_format (&priv->texture_info, GST_VIDEO_FORMAT_BGRx, GST_VIDEO_INFO_WIDTH (in_info), GST_VIDEO_INFO_HEIGHT (in_info)); unpack_function = RGB_TO_RGBx; break; case GST_VIDEO_FORMAT_RGB10A2_LE: gst_video_info_set_format (&priv->texture_info, GST_VIDEO_FORMAT_ARGB64, GST_VIDEO_INFO_WIDTH (in_info), GST_VIDEO_INFO_HEIGHT (in_info)); unpack_function = RGB10A2_TO_ARGB64; break; case GST_VIDEO_FORMAT_BGR10A2_LE: gst_video_info_set_format (&priv->texture_info, GST_VIDEO_FORMAT_ARGB64, GST_VIDEO_INFO_WIDTH (in_info), GST_VIDEO_INFO_HEIGHT (in_info)); unpack_function = BGR10A2_TO_ARGB64; break; default: break; } for (i = 0; i < G_N_ELEMENTS (format_map); i++) { if (format_map[i].format == GST_VIDEO_INFO_FORMAT (texture_info)) { priv->texture_fmt = &format_map[i]; break; } } if (!priv->texture_fmt) { GST_ERROR_OBJECT (self, "Couldn't find texture format for %s (%s)", gst_video_format_to_string (GST_VIDEO_INFO_FORMAT (in_info)), gst_video_format_to_string (GST_VIDEO_INFO_FORMAT (texture_info))); return FALSE; } /* calculate black color * TODO: add support border color */ if (GST_VIDEO_INFO_IS_RGB (out_info)) { GstVideoInfo rgb_info = *out_info; rgb_info.colorimetry.range = GST_VIDEO_COLOR_RANGE_0_255; gst_cuda_color_range_adjust_matrix_unorm (&rgb_info, out_info, &border_color_matrix); } else { GstVideoInfo rgb_info; gst_video_info_set_format (&rgb_info, GST_VIDEO_FORMAT_RGBA64_LE, out_info->width, out_info->height); gst_cuda_rgb_to_yuv_matrix_unorm (&rgb_info, out_info, &border_color_matrix); } for (i = 0; i < 3; i++) { /* TODO: property */ gdouble border_rgba[4] = { 0, 0, 0 }; border_color[i] = 0; for (j = 0; j < 3; j++) border_color[i] += border_color_matrix.matrix[i][j] * border_rgba[i]; border_color[i] = border_color_matrix.offset[i]; border_color[i] = CLAMP (border_color[i], border_color_matrix.min[i], border_color_matrix.max[i]); g_ascii_formatd (border_color_str[i], G_ASCII_DTOSTR_BUF_SIZE, "%f", border_color[i]); } g_ascii_formatd (border_color_str[3], G_ASCII_DTOSTR_BUF_SIZE, "%f", 1); /* FIXME: handle primaries and transfer functions */ if (GST_VIDEO_INFO_IS_RGB (texture_info)) { if (GST_VIDEO_INFO_IS_RGB (out_info)) { /* RGB -> RGB */ if (in_color->range == out_color->range) { GST_DEBUG_OBJECT (self, "RGB -> RGB conversion without matrix"); } else { if (!gst_cuda_color_range_adjust_matrix_unorm (in_info, out_info, &to_rgb_matrix)) { GST_ERROR_OBJECT (self, "Failed to get RGB range adjust matrix"); return FALSE; } str = gst_cuda_dump_color_matrix (&to_rgb_matrix); GST_DEBUG_OBJECT (self, "RGB range adjust %s -> %s\n%s", get_color_range_name (in_color->range), get_color_range_name (out_color->range), str); g_free (str); to_rgb_func = COLOR_SPACE_CONVERT; } } else { /* RGB -> YUV */ if (!gst_cuda_rgb_to_yuv_matrix_unorm (in_info, out_info, &to_yuv_matrix)) { GST_ERROR_OBJECT (self, "Failed to get RGB -> YUV transform matrix"); return FALSE; } str = gst_cuda_dump_color_matrix (&to_yuv_matrix); GST_DEBUG_OBJECT (self, "RGB -> YUV matrix:\n%s", str); g_free (str); to_yuv_func = COLOR_SPACE_CONVERT; } } else { if (GST_VIDEO_INFO_IS_RGB (out_info)) { /* YUV -> RGB */ if (!gst_cuda_yuv_to_rgb_matrix_unorm (in_info, out_info, &to_rgb_matrix)) { GST_ERROR_OBJECT (self, "Failed to get YUV -> RGB transform matrix"); return FALSE; } str = gst_cuda_dump_color_matrix (&to_rgb_matrix); GST_DEBUG_OBJECT (self, "YUV -> RGB matrix:\n%s", str); g_free (str); to_rgb_func = COLOR_SPACE_CONVERT; } else { /* YUV -> YUV */ if (in_color->range == out_color->range) { GST_DEBUG_OBJECT (self, "YUV -> YU conversion without matrix"); } else { if (!gst_cuda_color_range_adjust_matrix_unorm (in_info, out_info, &to_yuv_matrix)) { GST_ERROR_OBJECT (self, "Failed to get GRAY range adjust matrix"); return FALSE; } str = gst_cuda_dump_color_matrix (&to_yuv_matrix); GST_DEBUG_OBJECT (self, "YUV range adjust matrix:\n%s", str); g_free (str); to_yuv_func = COLOR_SPACE_CONVERT; } } } color_matrix_to_string (&to_rgb_matrix, &to_rgb_matrix_str); color_matrix_to_string (&to_yuv_matrix, &to_yuv_matrix_str); /* half pixel offset, to sample texture at center of the pixel position */ g_ascii_formatd (offset_x, G_ASCII_DTOSTR_BUF_SIZE, "%f", (gdouble) 0.5 / priv->dest_rect.width); g_ascii_formatd (offset_y, G_ASCII_DTOSTR_BUF_SIZE, "%f", (gdouble) 0.5 / priv->dest_rect.height); str = g_strdup_printf (TEMPLETA_KERNEL, KERNEL_COMMON, unpack_function ? unpack_function : "", /* TO RGB matrix */ to_rgb_matrix_str.matrix[0][0], to_rgb_matrix_str.matrix[0][1], to_rgb_matrix_str.matrix[0][2], to_rgb_matrix_str.matrix[1][0], to_rgb_matrix_str.matrix[1][1], to_rgb_matrix_str.matrix[1][2], to_rgb_matrix_str.matrix[2][0], to_rgb_matrix_str.matrix[2][1], to_rgb_matrix_str.matrix[2][2], to_rgb_matrix_str.offset[0], to_rgb_matrix_str.offset[1], to_rgb_matrix_str.offset[2], to_rgb_matrix_str.min[0], to_rgb_matrix_str.min[1], to_rgb_matrix_str.min[2], to_rgb_matrix_str.max[0], to_rgb_matrix_str.max[1], to_rgb_matrix_str.max[2], /* TO YUV matrix */ to_yuv_matrix_str.matrix[0][0], to_yuv_matrix_str.matrix[0][1], to_yuv_matrix_str.matrix[0][2], to_yuv_matrix_str.matrix[1][0], to_yuv_matrix_str.matrix[1][1], to_yuv_matrix_str.matrix[1][2], to_yuv_matrix_str.matrix[2][0], to_yuv_matrix_str.matrix[2][1], to_yuv_matrix_str.matrix[2][2], to_yuv_matrix_str.offset[0], to_yuv_matrix_str.offset[1], to_yuv_matrix_str.offset[2], to_yuv_matrix_str.min[0], to_yuv_matrix_str.min[1], to_yuv_matrix_str.min[2], to_yuv_matrix_str.max[0], to_yuv_matrix_str.max[1], to_yuv_matrix_str.max[2], /* width/height */ GST_VIDEO_INFO_WIDTH (out_info), GST_VIDEO_INFO_HEIGHT (out_info), /* viewport */ priv->dest_rect.x, priv->dest_rect.y, priv->dest_rect.x + priv->dest_rect.width, priv->dest_rect.y + priv->dest_rect.height, priv->dest_rect.width, priv->dest_rect.height, /* half pixel offsets */ offset_x, offset_y, /* border colors */ border_color_str[0], border_color_str[1], border_color_str[2], border_color_str[3], /* sampler function name */ priv->texture_fmt->sample_func, /* TO RGB conversion function name */ to_rgb_func, /* TO YUV conversion function name */ to_yuv_func, /* write function name */ write_func); GST_LOG_OBJECT (self, "kernel code:\n%s\n", str); ptx = gst_cuda_nvrtc_compile (str); g_free (str); if (!ptx) { GST_ERROR_OBJECT (self, "Could not compile code"); return FALSE; } if (priv->dest_rect.x != 0 || priv->dest_rect.y != 0 || priv->dest_rect.width != out_info->width || priv->dest_rect.height != out_info->height || in_info->width != out_info->width || in_info->height != out_info->height) { for (i = 0; i < G_N_ELEMENTS (priv->filter_mode); i++) priv->filter_mode[i] = CU_TR_FILTER_MODE_LINEAR; } else { for (i = 0; i < G_N_ELEMENTS (priv->filter_mode); i++) priv->filter_mode[i] = CU_TR_FILTER_MODE_POINT; } if (!gst_cuda_context_push (self->context)) { GST_ERROR_OBJECT (self, "Couldn't push context"); return FALSE; } /* Allocates intermediate memory for texture */ if (unpack_function) { CUDA_TEXTURE_DESC texture_desc; CUDA_RESOURCE_DESC resource_desc; CUtexObject texture = 0; memset (&texture_desc, 0, sizeof (CUDA_TEXTURE_DESC)); memset (&resource_desc, 0, sizeof (CUDA_RESOURCE_DESC)); ret = CuMemAllocPitch (&priv->unpack_buffer.ptr, &priv->unpack_buffer.stride, GST_VIDEO_INFO_COMP_WIDTH (texture_info, 0) * GST_VIDEO_INFO_COMP_PSTRIDE (texture_info, 0), GST_VIDEO_INFO_HEIGHT (texture_info), 16); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Couldn't allocate unpack buffer"); goto error; } resource_desc.resType = CU_RESOURCE_TYPE_PITCH2D; resource_desc.res.pitch2D.format = priv->texture_fmt->array_format[0]; resource_desc.res.pitch2D.numChannels = 4; resource_desc.res.pitch2D.width = in_info->width; resource_desc.res.pitch2D.height = in_info->height; resource_desc.res.pitch2D.pitchInBytes = priv->unpack_buffer.stride; resource_desc.res.pitch2D.devPtr = priv->unpack_buffer.ptr; texture_desc.filterMode = priv->filter_mode[0]; texture_desc.flags = 0x2; texture_desc.addressMode[0] = 1; texture_desc.addressMode[1] = 1; texture_desc.addressMode[2] = 1; ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, NULL); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Couldn't create unpack texture"); goto error; } priv->unpack_buffer.texture = texture; } ret = CuModuleLoadData (&priv->module, ptx); g_free (ptx); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Could not load module"); priv->module = NULL; goto error; } ret = CuModuleGetFunction (&priv->main_func, priv->module, GST_CUDA_KERNEL_MAIN_FUNC); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Could not get main function"); goto error; } if (unpack_function) { ret = CuModuleGetFunction (&priv->unpack_func, priv->module, GST_CUDA_KERNEL_UNPACK_FUNC); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Could not get unpack function"); goto error; } } gst_cuda_context_pop (NULL); return TRUE; error: gst_cuda_context_pop (NULL); return FALSE; } static gboolean copy_config (GQuark field_id, const GValue * value, gpointer user_data) { GstCudaConverter *self = (GstCudaConverter *) user_data; gst_structure_id_set_value (self->priv->config, field_id, value); return TRUE; } static void gst_cuda_converter_set_config (GstCudaConverter * self, GstStructure * config) { gst_structure_foreach (config, copy_config, self); gst_structure_free (config); } static gint get_opt_int (GstCudaConverter * self, const gchar * opt, gint def) { gint res; if (!gst_structure_get_int (self->priv->config, opt, &res)) res = def; return res; } GstCudaConverter * gst_cuda_converter_new (const GstVideoInfo * in_info, const GstVideoInfo * out_info, GstCudaContext * context, GstStructure * config) { GstCudaConverter *self; GstCudaConverterPrivate *priv; g_return_val_if_fail (in_info != NULL, NULL); g_return_val_if_fail (out_info != NULL, NULL); g_return_val_if_fail (GST_IS_CUDA_CONTEXT (context), NULL); self = g_object_new (GST_TYPE_CUDA_CONVERTER, NULL); if (!GST_IS_CUDA_CONTEXT (context)) { GST_WARNING_OBJECT (self, "Not a valid cuda context object"); goto error; } self->context = gst_object_ref (context); priv = self->priv; priv->in_info = *in_info; priv->out_info = *out_info; if (config) gst_cuda_converter_set_config (self, config); priv->dest_rect.x = get_opt_int (self, GST_CUDA_CONVERTER_OPT_DEST_X, 0); priv->dest_rect.y = get_opt_int (self, GST_CUDA_CONVERTER_OPT_DEST_Y, 0); priv->dest_rect.width = get_opt_int (self, GST_CUDA_CONVERTER_OPT_DEST_WIDTH, out_info->width); priv->dest_rect.height = get_opt_int (self, GST_CUDA_CONVERTER_OPT_DEST_HEIGHT, out_info->height); if (!gst_cuda_converter_setup (self)) goto error; priv->texture_align = gst_cuda_context_get_texture_alignment (context); gst_object_ref_sink (self); return self; error: gst_object_unref (self); return NULL; } static CUtexObject gst_cuda_converter_create_texture_unchecked (GstCudaConverter * self, CUdeviceptr src, gint width, gint height, CUarray_format format, guint channels, gint stride, gint plane, CUfilter_mode mode) { CUDA_TEXTURE_DESC texture_desc; CUDA_RESOURCE_DESC resource_desc; CUtexObject texture = 0; CUresult cuda_ret; memset (&texture_desc, 0, sizeof (CUDA_TEXTURE_DESC)); memset (&resource_desc, 0, sizeof (CUDA_RESOURCE_DESC)); resource_desc.resType = CU_RESOURCE_TYPE_PITCH2D; resource_desc.res.pitch2D.format = format; resource_desc.res.pitch2D.numChannels = channels; resource_desc.res.pitch2D.width = width; resource_desc.res.pitch2D.height = height; resource_desc.res.pitch2D.pitchInBytes = stride; resource_desc.res.pitch2D.devPtr = src; texture_desc.filterMode = mode; /* Will read texture value as a normalized [0, 1] float value * with [0, 1) coordinates */ /* CU_TRSF_NORMALIZED_COORDINATES */ texture_desc.flags = 0x2; /* CU_TR_ADDRESS_MODE_CLAMP */ texture_desc.addressMode[0] = 1; texture_desc.addressMode[1] = 1; texture_desc.addressMode[2] = 1; cuda_ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, NULL); if (!gst_cuda_result (cuda_ret)) { GST_ERROR_OBJECT (self, "Could not create texture"); return 0; } return texture; } static gboolean ensure_fallback_buffer (GstCudaConverter * self, gint width_in_bytes, gint height, guint plane) { GstCudaConverterPrivate *priv = self->priv; CUresult ret; if (priv->fallback_buffer[plane].ptr) return TRUE; ret = CuMemAllocPitch (&priv->fallback_buffer[plane].ptr, &priv->fallback_buffer[plane].stride, width_in_bytes, height, 16); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Couldn't allocate fallback buffer"); return FALSE; } return TRUE; } static CUtexObject gst_cuda_converter_create_texture (GstCudaConverter * self, CUdeviceptr src, gint width, gint height, gint stride, CUfilter_mode mode, CUarray_format format, guint channles, gint plane, CUstream stream) { GstCudaConverterPrivate *priv = self->priv; CUresult ret; CUdeviceptr src_ptr; CUDA_MEMCPY2D params = { 0, }; if (!ensure_fallback_buffer (self, stride, height, plane)) return 0; params.srcMemoryType = CU_MEMORYTYPE_DEVICE; params.srcPitch = stride; params.srcDevice = (CUdeviceptr) src; params.dstMemoryType = CU_MEMORYTYPE_DEVICE; params.dstPitch = priv->fallback_buffer[plane].stride; params.dstDevice = priv->fallback_buffer[plane].ptr; params.WidthInBytes = GST_VIDEO_INFO_COMP_WIDTH (&priv->in_info, plane) * GST_VIDEO_INFO_COMP_PSTRIDE (&priv->in_info, plane), params.Height = GST_VIDEO_INFO_COMP_HEIGHT (&priv->in_info, plane); ret = CuMemcpy2DAsync (¶ms, stream); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Couldn't copy to fallback buffer"); return 0; } if (!priv->fallback_buffer[plane].texture) { src_ptr = priv->fallback_buffer[plane].ptr; stride = priv->fallback_buffer[plane].stride; priv->fallback_buffer[plane].texture = gst_cuda_converter_create_texture_unchecked (self, src_ptr, width, height, format, channles, stride, plane, mode); } return priv->fallback_buffer[plane].texture; } static gboolean gst_cuda_converter_unpack_rgb (GstCudaConverter * self, GstVideoFrame * src_frame, CUstream stream) { GstCudaConverterPrivate *priv = self->priv; CUdeviceptr src; gint width, height, src_stride, dst_stride; CUresult ret; gpointer args[] = { &src, &priv->unpack_buffer.ptr, &width, &height, &src_stride, &dst_stride }; g_assert (priv->unpack_buffer.ptr); g_assert (priv->unpack_buffer.stride > 0); src = (CUdeviceptr) GST_VIDEO_FRAME_PLANE_DATA (src_frame, 0); width = GST_VIDEO_FRAME_WIDTH (src_frame); height = GST_VIDEO_FRAME_HEIGHT (src_frame); src_stride = GST_VIDEO_FRAME_PLANE_STRIDE (src_frame, 0); dst_stride = (gint) priv->unpack_buffer.stride; ret = CuLaunchKernel (priv->unpack_func, DIV_UP (width, CUDA_BLOCK_X), DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0, stream, args, NULL); if (!gst_cuda_result (ret)) { GST_ERROR_OBJECT (self, "Couldn't unpack source RGB"); return FALSE; } return TRUE; } gboolean gst_cuda_converter_convert_frame (GstCudaConverter * converter, GstVideoFrame * src_frame, GstVideoFrame * dst_frame, CUstream stream, gboolean * synchronized) { GstCudaConverterPrivate *priv; const TextureFormat *format; CUtexObject texture[GST_VIDEO_MAX_COMPONENTS] = { 0, }; guint8 *dst[GST_VIDEO_MAX_COMPONENTS] = { NULL, }; gint stride[2] = { 0, }; gint i; gboolean ret = FALSE; CUresult cuda_ret; gint width, height; gpointer args[] = { &texture[0], &texture[1], &texture[2], &texture[3], &dst[0], &dst[1], &dst[2], &dst[3], &stride[0], &stride[1] }; gboolean need_sync = FALSE; GstCudaMemory *cmem; g_return_val_if_fail (GST_IS_CUDA_CONVERTER (converter), FALSE); g_return_val_if_fail (src_frame != NULL, FALSE); g_return_val_if_fail (dst_frame != NULL, FALSE); priv = converter->priv; format = priv->texture_fmt; g_assert (format); cmem = (GstCudaMemory *) gst_buffer_peek_memory (src_frame->buffer, 0); g_return_val_if_fail (gst_is_cuda_memory (GST_MEMORY_CAST (cmem)), FALSE); if (!gst_cuda_context_push (converter->context)) { GST_ERROR_OBJECT (converter, "Couldn't push context"); return FALSE; } if (priv->unpack_func) { if (!gst_cuda_converter_unpack_rgb (converter, src_frame, stream)) goto out; texture[0] = priv->unpack_buffer.texture; if (!texture[0]) { GST_ERROR_OBJECT (converter, "Unpack texture is unavailable"); goto out; } } else { for (i = 0; i < GST_VIDEO_FRAME_N_PLANES (src_frame); i++) { if (!gst_cuda_memory_get_texture (cmem, i, priv->filter_mode[i], &texture[i])) { CUdeviceptr src; src = (CUdeviceptr) GST_VIDEO_FRAME_PLANE_DATA (src_frame, i); texture[i] = gst_cuda_converter_create_texture (converter, src, GST_VIDEO_FRAME_COMP_WIDTH (src_frame, i), GST_VIDEO_FRAME_COMP_HEIGHT (src_frame, i), GST_VIDEO_FRAME_PLANE_STRIDE (src_frame, i), priv->filter_mode[i], format->array_format[i], format->channels[i], i, stream); need_sync = TRUE; } if (!texture[i]) { GST_ERROR_OBJECT (converter, "Couldn't create texture %d", i); goto out; } } } width = GST_VIDEO_FRAME_WIDTH (dst_frame); height = GST_VIDEO_FRAME_HEIGHT (dst_frame); for (i = 0; i < GST_VIDEO_FRAME_N_PLANES (dst_frame); i++) dst[i] = GST_VIDEO_FRAME_PLANE_DATA (dst_frame, i); stride[0] = stride[1] = GST_VIDEO_FRAME_PLANE_STRIDE (dst_frame, 0); if (GST_VIDEO_FRAME_N_PLANES (dst_frame) > 1) stride[1] = GST_VIDEO_FRAME_PLANE_STRIDE (dst_frame, 1); cuda_ret = CuLaunchKernel (priv->main_func, DIV_UP (width, CUDA_BLOCK_X), DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0, stream, args, NULL); if (!gst_cuda_result (cuda_ret)) { GST_ERROR_OBJECT (converter, "Couldn't convert frame"); goto out; } if (need_sync) CuStreamSynchronize (stream); if (synchronized) *synchronized = need_sync; ret = TRUE; out: gst_cuda_context_pop (NULL); return ret; }