mirror of
https://gitlab.freedesktop.org/gstreamer/gstreamer.git
synced 2024-12-23 00:36:51 +00:00
cudaconvertscale: Add support for flip/rotation
Similar to the d3d11convert element, colorspace conversion, resizing and flip/rotation operations can be done in a single kernel function call Part-of: <https://gitlab.freedesktop.org/gstreamer/gstreamer/-/merge_requests/4640>
This commit is contained in:
Seungha Yang
GStreamer Marge Bot
parent
d335eb8c7c
commit
7a3be74b63
3 changed files with 368 additions and 20 deletions
|
|
@ -664,6 +664,14 @@ typedef struct
|
|||
#define WRITE_BGRP "write_bgrp"
|
||||
#define WRITE_GBR "write_gbr"
|
||||
#define WRITE_GBRA "write_gbra"
|
||||
#define ROTATE_IDENTITY "rotate_identity"
|
||||
#define ROTATE_90R "rotate_90r"
|
||||
#define ROTATE_180 "rotate_180"
|
||||
#define ROTATE_90L "rotate_90l"
|
||||
#define ROTATE_HORIZ "rotate_horiz"
|
||||
#define ROTATE_VERT "rotate_vert"
|
||||
#define ROTATE_UL_LR "rotate_ul_lr"
|
||||
#define ROTATE_UR_LL "rotate_ur_ll"
|
||||
|
||||
/* *INDENT-OFF* */
|
||||
const static gchar KERNEL_COMMON[] =
|
||||
|
|
@ -1173,7 +1181,55 @@ WRITE_GBRA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\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";
|
||||
"}\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_IDENTITY "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(x, y);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_90R "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(y, 1.0 - x);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_180 "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(1.0 - x, 1.0 - y);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_90L "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(1.0 - y, x);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_HORIZ "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(1.0 - x, y);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_VERT "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(x, 1.0 - y);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_UL_LR "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(y, x);\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__device__ inline float2\n"
|
||||
ROTATE_UR_LL "(float x, float y)\n"
|
||||
"{\n"
|
||||
" return make_float2(1.0 - y, 1.0 - x);\n"
|
||||
"}\n"
|
||||
"\n";
|
||||
|
||||
#define GST_CUDA_KERNEL_UNPACK_FUNC "gst_cuda_kernel_unpack_func"
|
||||
static const gchar RGB_TO_RGBx[] =
|
||||
|
|
@ -1307,7 +1363,8 @@ GST_CUDA_KERNEL_MAIN_FUNC "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\
|
|||
" } 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"
|
||||
" float2 rotated = %s (x, y);\n"
|
||||
" float4 s = %s (tex0, tex1, tex2, tex3, rotated.x, rotated.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"
|
||||
|
|
@ -1389,6 +1446,8 @@ struct _GstCudaConverterPrivate
|
|||
GstVideoInfo in_info;
|
||||
GstVideoInfo out_info;
|
||||
|
||||
GstVideoOrientationMethod method;
|
||||
|
||||
GstStructure *config;
|
||||
|
||||
GstVideoInfo texture_info;
|
||||
|
|
@ -1549,6 +1608,7 @@ gst_cuda_converter_setup (GstCudaConverter * self)
|
|||
const gchar *write_func = NULL;
|
||||
const gchar *to_rgb_func = COLOR_SPACE_IDENTITY;
|
||||
const gchar *to_yuv_func = COLOR_SPACE_IDENTITY;
|
||||
const gchar *rotate_func = ROTATE_IDENTITY;
|
||||
const GstVideoColorimetry *in_color;
|
||||
const GstVideoColorimetry *out_color;
|
||||
gchar *str;
|
||||
|
|
@ -1810,6 +1870,32 @@ gst_cuda_converter_setup (GstCudaConverter * self)
|
|||
g_ascii_formatd (offset_y, G_ASCII_DTOSTR_BUF_SIZE, "%f",
|
||||
(gdouble) 0.5 / priv->dest_rect.height);
|
||||
|
||||
switch (priv->method) {
|
||||
case GST_VIDEO_ORIENTATION_90R:
|
||||
rotate_func = ROTATE_90R;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_180:
|
||||
rotate_func = ROTATE_180;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_90L:
|
||||
rotate_func = ROTATE_90L;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_HORIZ:
|
||||
rotate_func = ROTATE_HORIZ;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_VERT:
|
||||
rotate_func = ROTATE_VERT;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_UL_LR:
|
||||
rotate_func = ROTATE_UL_LR;
|
||||
break;
|
||||
case GST_VIDEO_ORIENTATION_UR_LL:
|
||||
rotate_func = ROTATE_UR_LL;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
str = g_strdup_printf (TEMPLETA_KERNEL, KERNEL_COMMON,
|
||||
unpack_function ? unpack_function : "",
|
||||
/* TO RGB matrix */
|
||||
|
|
@ -1860,6 +1946,8 @@ gst_cuda_converter_setup (GstCudaConverter * self)
|
|||
/* border colors */
|
||||
border_color_str[0], border_color_str[1],
|
||||
border_color_str[2], border_color_str[3],
|
||||
/* adjust coord before sampling */
|
||||
rotate_func,
|
||||
/* sampler function name */
|
||||
priv->texture_fmt->sample_func,
|
||||
/* TO RGB conversion function name */
|
||||
|
|
@ -2003,6 +2091,7 @@ gst_cuda_converter_new (const GstVideoInfo * in_info,
|
|||
{
|
||||
GstCudaConverter *self;
|
||||
GstCudaConverterPrivate *priv;
|
||||
gint method;
|
||||
|
||||
g_return_val_if_fail (in_info != NULL, NULL);
|
||||
g_return_val_if_fail (out_info != NULL, NULL);
|
||||
|
|
@ -2029,6 +2118,12 @@ gst_cuda_converter_new (const GstVideoInfo * in_info,
|
|||
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_structure_get_enum (priv->config,
|
||||
GST_CUDA_CONVERTER_OPT_ORIENTATION_METHOD,
|
||||
GST_TYPE_VIDEO_ORIENTATION_METHOD, &method)) {
|
||||
priv->method = method;
|
||||
GST_DEBUG_OBJECT (self, "Selected orientation method %d", method);
|
||||
}
|
||||
|
||||
if (!gst_cuda_converter_setup (self))
|
||||
goto error;
|
||||
|
|
|
|||
|
|
@ -64,6 +64,13 @@ typedef struct _GstCudaConverterPrivate GstCudaConverterPrivate;
|
|||
*/
|
||||
#define GST_CUDA_CONVERTER_OPT_DEST_HEIGHT "GstCudaConverter.dest-height"
|
||||
|
||||
/**
|
||||
* GST_CUDA_CONVERTER_OPT_ORIENTATION_METHOD:
|
||||
*
|
||||
* #GstVideoOrientationMethod, default #GST_VIDEO_ORIENTATION_IDENTITY
|
||||
*/
|
||||
#define GST_CUDA_CONVERTER_OPT_ORIENTATION_METHOD "GstCudaConverter.orientation-method"
|
||||
|
||||
struct _GstCudaConverter
|
||||
{
|
||||
GstObject parent;
|
||||
|
|
|
|||
|
|
@ -60,9 +60,22 @@ struct _GstCudaBaseConvert
|
|||
gint borders_h;
|
||||
gint borders_w;
|
||||
gboolean add_borders;
|
||||
|
||||
/* orientation */
|
||||
/* method configured via property */
|
||||
GstVideoOrientationMethod method;
|
||||
/* method parsed from tag */
|
||||
GstVideoOrientationMethod tag_method;
|
||||
/* method currently selected based on "method" and "tag_method" */
|
||||
GstVideoOrientationMethod selected_method;
|
||||
/* method previously selected and used for negotiation */
|
||||
GstVideoOrientationMethod active_method;
|
||||
|
||||
GMutex lock;
|
||||
};
|
||||
|
||||
static void gst_cuda_base_convert_dispose (GObject * object);
|
||||
static void gst_cuda_base_convert_finalize (GObject * object);
|
||||
static GstCaps *gst_cuda_base_convert_transform_caps (GstBaseTransform * trans,
|
||||
GstPadDirection direction, GstCaps * caps, GstCaps * filter);
|
||||
static GstCaps *gst_cuda_base_convert_fixate_caps (GstBaseTransform * trans,
|
||||
|
|
@ -103,6 +116,7 @@ gst_cuda_base_convert_class_init (GstCudaBaseConvertClass * klass)
|
|||
GST_CUDA_BASE_TRANSFORM_CLASS (klass);
|
||||
|
||||
gobject_class->dispose = gst_cuda_base_convert_dispose;
|
||||
gobject_class->finalize = gst_cuda_base_convert_finalize;
|
||||
|
||||
gst_element_class_add_static_pad_template (element_class, &sink_template);
|
||||
gst_element_class_add_static_pad_template (element_class, &src_template);
|
||||
|
|
@ -130,6 +144,7 @@ static void
|
|||
gst_cuda_base_convert_init (GstCudaBaseConvert * self)
|
||||
{
|
||||
self->add_borders = DEFAULT_ADD_BORDERS;
|
||||
g_mutex_init (&self->lock);
|
||||
}
|
||||
|
||||
static void
|
||||
|
|
@ -143,6 +158,16 @@ gst_cuda_base_convert_dispose (GObject * object)
|
|||
G_OBJECT_CLASS (parent_class)->dispose (object);
|
||||
}
|
||||
|
||||
static void
|
||||
gst_cuda_base_convert_finalize (GObject * object)
|
||||
{
|
||||
GstCudaBaseConvert *self = GST_CUDA_BASE_CONVERT (object);
|
||||
|
||||
g_mutex_clear (&self->lock);
|
||||
|
||||
G_OBJECT_CLASS (parent_class)->finalize (object);
|
||||
}
|
||||
|
||||
static GstCaps *
|
||||
gst_cuda_base_convert_caps_remove_format_info (GstCaps * caps)
|
||||
{
|
||||
|
|
@ -610,9 +635,11 @@ static GstCaps *
|
|||
gst_cuda_base_convert_fixate_size (GstBaseTransform * base,
|
||||
GstPadDirection direction, GstCaps * caps, GstCaps * othercaps)
|
||||
{
|
||||
GstCudaBaseConvert *self = GST_CUDA_BASE_CONVERT (base);
|
||||
GstStructure *ins, *outs;
|
||||
const GValue *from_par, *to_par;
|
||||
GValue fpar = G_VALUE_INIT, tpar = G_VALUE_INIT;
|
||||
gboolean rotate = FALSE;
|
||||
|
||||
othercaps = gst_caps_truncate (othercaps);
|
||||
othercaps = gst_caps_make_writable (othercaps);
|
||||
|
|
@ -622,6 +649,19 @@ gst_cuda_base_convert_fixate_size (GstBaseTransform * base,
|
|||
from_par = gst_structure_get_value (ins, "pixel-aspect-ratio");
|
||||
to_par = gst_structure_get_value (outs, "pixel-aspect-ratio");
|
||||
|
||||
g_mutex_lock (&self->lock);
|
||||
switch (self->selected_method) {
|
||||
case GST_VIDEO_ORIENTATION_90R:
|
||||
case GST_VIDEO_ORIENTATION_90L:
|
||||
case GST_VIDEO_ORIENTATION_UL_LR:
|
||||
case GST_VIDEO_ORIENTATION_UR_LL:
|
||||
rotate = TRUE;
|
||||
break;
|
||||
default:
|
||||
rotate = FALSE;
|
||||
break;
|
||||
}
|
||||
|
||||
if (direction == GST_PAD_SINK) {
|
||||
if (!from_par) {
|
||||
g_value_init (&fpar, GST_TYPE_FRACTION);
|
||||
|
|
@ -650,8 +690,13 @@ gst_cuda_base_convert_fixate_size (GstBaseTransform * base,
|
|||
if (!to_par) {
|
||||
gint to_par_n, to_par_d;
|
||||
|
||||
to_par_n = from_par_n;
|
||||
to_par_d = from_par_d;
|
||||
if (rotate) {
|
||||
to_par_n = from_par_n;
|
||||
to_par_d = from_par_d;
|
||||
} else {
|
||||
to_par_n = from_par_n;
|
||||
to_par_d = from_par_d;
|
||||
}
|
||||
|
||||
g_value_init (&tpar, GST_TYPE_FRACTION);
|
||||
gst_value_set_fraction (&tpar, to_par_n, to_par_d);
|
||||
|
|
@ -681,6 +726,17 @@ gst_cuda_base_convert_fixate_size (GstBaseTransform * base,
|
|||
gst_structure_get_int (outs, "width", &w);
|
||||
gst_structure_get_int (outs, "height", &h);
|
||||
|
||||
/* swap dimensions when it's rotated */
|
||||
if (rotate) {
|
||||
gint _tmp = from_w;
|
||||
from_w = from_h;
|
||||
from_h = _tmp;
|
||||
|
||||
_tmp = from_par_n;
|
||||
from_par_n = from_par_d;
|
||||
from_par_d = _tmp;
|
||||
}
|
||||
|
||||
/* if both width and height are already fixed, we can't do anything
|
||||
* about it anymore */
|
||||
if (w && h) {
|
||||
|
|
@ -1046,6 +1102,7 @@ done:
|
|||
g_value_unset (&fpar);
|
||||
if (to_par == &tpar)
|
||||
g_value_unset (&tpar);
|
||||
g_mutex_unlock (&self->lock);
|
||||
|
||||
return othercaps;
|
||||
}
|
||||
|
|
@ -1241,6 +1298,32 @@ gst_cuda_base_convert_decide_allocation (GstBaseTransform * trans,
|
|||
query);
|
||||
}
|
||||
|
||||
static gboolean
|
||||
needs_color_convert (const GstVideoInfo * in_info,
|
||||
const GstVideoInfo * out_info)
|
||||
{
|
||||
const GstVideoColorimetry *in_cinfo = &in_info->colorimetry;
|
||||
const GstVideoColorimetry *out_cinfo = &out_info->colorimetry;
|
||||
|
||||
if (in_cinfo->range != out_cinfo->range ||
|
||||
in_cinfo->matrix != out_cinfo->matrix) {
|
||||
return TRUE;
|
||||
}
|
||||
|
||||
if (!gst_video_color_primaries_is_equivalent (in_cinfo->primaries,
|
||||
out_cinfo->primaries)) {
|
||||
return TRUE;
|
||||
}
|
||||
|
||||
if (!gst_video_transfer_function_is_equivalent (in_cinfo->transfer,
|
||||
GST_VIDEO_INFO_COMP_DEPTH (in_info, 0), out_cinfo->transfer,
|
||||
GST_VIDEO_INFO_COMP_DEPTH (out_info, 0))) {
|
||||
return TRUE;
|
||||
}
|
||||
|
||||
return FALSE;
|
||||
}
|
||||
|
||||
static gboolean
|
||||
gst_cuda_base_convert_set_info (GstCudaBaseTransform * btrans,
|
||||
GstCaps * incaps, GstVideoInfo * in_info, GstCaps * outcaps,
|
||||
|
|
@ -1248,13 +1331,39 @@ gst_cuda_base_convert_set_info (GstCudaBaseTransform * btrans,
|
|||
{
|
||||
GstCudaBaseConvert *self = GST_CUDA_BASE_CONVERT (btrans);
|
||||
gint from_dar_n, from_dar_d, to_dar_n, to_dar_d;
|
||||
GstVideoInfo tmp_info;
|
||||
gboolean need_flip = FALSE;
|
||||
gint in_width, in_height, in_par_n, in_par_d;
|
||||
GstVideoOrientationMethod active_method;
|
||||
|
||||
gst_clear_object (&self->converter);
|
||||
|
||||
if (!gst_util_fraction_multiply (in_info->width,
|
||||
in_info->height, in_info->par_n, in_info->par_d, &from_dar_n,
|
||||
&from_dar_d)) {
|
||||
g_mutex_lock (&self->lock);
|
||||
active_method = self->active_method = self->selected_method;
|
||||
g_mutex_unlock (&self->lock);
|
||||
|
||||
if (active_method != GST_VIDEO_ORIENTATION_IDENTITY)
|
||||
need_flip = TRUE;
|
||||
|
||||
switch (active_method) {
|
||||
case GST_VIDEO_ORIENTATION_90R:
|
||||
case GST_VIDEO_ORIENTATION_90L:
|
||||
case GST_VIDEO_ORIENTATION_UL_LR:
|
||||
case GST_VIDEO_ORIENTATION_UR_LL:
|
||||
in_width = in_info->height;
|
||||
in_height = in_info->width;
|
||||
in_par_n = in_info->par_d;
|
||||
in_par_d = in_info->par_n;
|
||||
break;
|
||||
default:
|
||||
in_width = in_info->width;
|
||||
in_height = in_info->height;
|
||||
in_par_n = in_info->par_n;
|
||||
in_par_d = in_info->par_d;
|
||||
break;
|
||||
}
|
||||
|
||||
if (!gst_util_fraction_multiply (in_width,
|
||||
in_height, in_par_n, in_par_d, &from_dar_n, &from_dar_d)) {
|
||||
from_dar_n = from_dar_d = -1;
|
||||
}
|
||||
|
||||
|
|
@ -1296,14 +1405,10 @@ gst_cuda_base_convert_set_info (GstCudaBaseTransform * btrans,
|
|||
return FALSE;
|
||||
}
|
||||
|
||||
/* if the only thing different in the caps is the transfer function, and
|
||||
* we're converting between equivalent transfer functions, do passthrough */
|
||||
tmp_info = *in_info;
|
||||
tmp_info.colorimetry.transfer = out_info->colorimetry.transfer;
|
||||
if (gst_video_info_is_equal (&tmp_info, out_info) &&
|
||||
gst_video_transfer_function_is_equivalent (in_info->colorimetry.transfer,
|
||||
in_info->finfo->bits, out_info->colorimetry.transfer,
|
||||
out_info->finfo->bits)) {
|
||||
if (in_width == out_info->width && in_height == out_info->height
|
||||
&& in_info->finfo == out_info->finfo && self->borders_w == 0 &&
|
||||
self->borders_h == 0 && !need_flip &&
|
||||
!needs_color_convert (in_info, out_info)) {
|
||||
gst_base_transform_set_passthrough (GST_BASE_TRANSFORM (self), TRUE);
|
||||
} else {
|
||||
GstStructure *config;
|
||||
|
|
@ -1317,7 +1422,9 @@ gst_cuda_base_convert_set_info (GstCudaBaseTransform * btrans,
|
|||
GST_CUDA_CONVERTER_OPT_DEST_WIDTH,
|
||||
G_TYPE_INT, out_info->width - self->borders_w,
|
||||
GST_CUDA_CONVERTER_OPT_DEST_HEIGHT,
|
||||
G_TYPE_INT, out_info->height - self->borders_h, NULL);
|
||||
G_TYPE_INT, out_info->height - self->borders_h,
|
||||
GST_CUDA_CONVERTER_OPT_ORIENTATION_METHOD,
|
||||
GST_TYPE_VIDEO_ORIENTATION_METHOD, active_method, NULL);
|
||||
|
||||
self->converter = gst_cuda_converter_new (in_info,
|
||||
out_info, btrans->context, config);
|
||||
|
|
@ -1463,6 +1570,37 @@ gst_cuda_base_convert_set_add_border (GstCudaBaseConvert * self,
|
|||
gst_base_transform_reconfigure_src (GST_BASE_TRANSFORM_CAST (self));
|
||||
}
|
||||
|
||||
static void
|
||||
gst_cuda_base_convert_set_orientation (GstCudaBaseConvert * self,
|
||||
GstVideoOrientationMethod method, gboolean from_tag)
|
||||
{
|
||||
if (method == GST_VIDEO_ORIENTATION_CUSTOM) {
|
||||
GST_WARNING_OBJECT (self, "Unsupported custom orientation");
|
||||
return;
|
||||
}
|
||||
|
||||
g_mutex_lock (&self->lock);
|
||||
if (from_tag)
|
||||
self->tag_method = method;
|
||||
else
|
||||
self->method = method;
|
||||
|
||||
if (self->method == GST_VIDEO_ORIENTATION_AUTO) {
|
||||
self->selected_method = self->tag_method;
|
||||
} else {
|
||||
self->selected_method = self->method;
|
||||
}
|
||||
|
||||
if (self->selected_method != self->active_method) {
|
||||
GST_DEBUG_OBJECT (self, "Rotation orientation %d -> %d",
|
||||
self->active_method, self->selected_method);
|
||||
|
||||
gst_base_transform_reconfigure_src (GST_BASE_TRANSFORM (self));
|
||||
}
|
||||
|
||||
g_mutex_unlock (&self->lock);
|
||||
}
|
||||
|
||||
/**
|
||||
* SECTION:element-cudaconvertscale
|
||||
* @title: cudaconvertscale
|
||||
|
|
@ -1486,6 +1624,7 @@ enum
|
|||
{
|
||||
PROP_CONVERT_SCALE_0,
|
||||
PROP_CONVERT_SCALE_ADD_BORDERS,
|
||||
PROP_CONVERT_SCALE_VIDEO_DIRECTION,
|
||||
};
|
||||
|
||||
struct _GstCudaConvertScale
|
||||
|
|
@ -1493,19 +1632,34 @@ struct _GstCudaConvertScale
|
|||
GstCudaBaseConvert parent;
|
||||
};
|
||||
|
||||
static void
|
||||
gst_cuda_convert_scale_video_direction_interface_init
|
||||
(GstVideoDirectionInterface * iface)
|
||||
{
|
||||
}
|
||||
|
||||
static void gst_cuda_convert_scale_set_property (GObject * object,
|
||||
guint prop_id, const GValue * value, GParamSpec * pspec);
|
||||
static void gst_cuda_convert_scale_get_property (GObject * object,
|
||||
guint prop_id, GValue * value, GParamSpec * pspec);
|
||||
static gboolean gst_cuda_convert_scale_sink_event (GstBaseTransform * trans,
|
||||
GstEvent * event);
|
||||
|
||||
G_DEFINE_TYPE (GstCudaConvertScale, gst_cuda_convert_scale,
|
||||
GST_TYPE_CUDA_BASE_CONVERT);
|
||||
#define gst_cuda_convert_scale_parent_class convert_scale_parent_class
|
||||
G_DEFINE_TYPE_WITH_CODE (GstCudaConvertScale, gst_cuda_convert_scale,
|
||||
GST_TYPE_CUDA_BASE_CONVERT,
|
||||
G_IMPLEMENT_INTERFACE (GST_TYPE_VIDEO_DIRECTION,
|
||||
gst_cuda_convert_scale_video_direction_interface_init));
|
||||
|
||||
static void gst_cuda_convert_scale_before_transform (GstBaseTransform * trans,
|
||||
GstBuffer * buffer);
|
||||
|
||||
static void
|
||||
gst_cuda_convert_scale_class_init (GstCudaConvertScaleClass * klass)
|
||||
{
|
||||
GObjectClass *gobject_class = G_OBJECT_CLASS (klass);
|
||||
GstElementClass *element_class = GST_ELEMENT_CLASS (klass);
|
||||
GstBaseTransformClass *transform_class = GST_BASE_TRANSFORM_CLASS (klass);
|
||||
|
||||
gobject_class->set_property = gst_cuda_convert_scale_set_property;
|
||||
gobject_class->get_property = gst_cuda_convert_scale_get_property;
|
||||
|
|
@ -1517,11 +1671,27 @@ gst_cuda_convert_scale_class_init (GstCudaConvertScaleClass * klass)
|
|||
DEFAULT_ADD_BORDERS, (GParamFlags) (GST_PARAM_MUTABLE_PLAYING |
|
||||
G_PARAM_READWRITE | G_PARAM_STATIC_STRINGS)));
|
||||
|
||||
/**
|
||||
* GstCudaConvertScale:video-direction:
|
||||
*
|
||||
* Video rotation/flip method to use
|
||||
*
|
||||
* Since: 1.24
|
||||
*/
|
||||
g_object_class_override_property (gobject_class,
|
||||
PROP_CONVERT_SCALE_VIDEO_DIRECTION, "video-direction");
|
||||
|
||||
gst_element_class_set_static_metadata (element_class,
|
||||
"CUDA colorspace converter and scaler",
|
||||
"Filter/Converter/Video/Scaler/Colorspace/Hardware",
|
||||
"Filter/Converter/Video/Scaler/Colorspace/Effect/Hardware",
|
||||
"Resizes video and allow color conversion using CUDA",
|
||||
"Seungha Yang <seungha@centricular.com>");
|
||||
|
||||
transform_class->passthrough_on_same_caps = FALSE;
|
||||
transform_class->before_transform =
|
||||
GST_DEBUG_FUNCPTR (gst_cuda_convert_scale_before_transform);
|
||||
transform_class->sink_event =
|
||||
GST_DEBUG_FUNCPTR (gst_cuda_convert_scale_sink_event);
|
||||
}
|
||||
|
||||
static void
|
||||
|
|
@ -1539,6 +1709,10 @@ gst_cuda_convert_scale_set_property (GObject * object, guint prop_id,
|
|||
case PROP_CONVERT_SCALE_ADD_BORDERS:
|
||||
gst_cuda_base_convert_set_add_border (base, g_value_get_boolean (value));
|
||||
break;
|
||||
case PROP_CONVERT_SCALE_VIDEO_DIRECTION:
|
||||
gst_cuda_base_convert_set_orientation (base, g_value_get_enum (value),
|
||||
FALSE);
|
||||
break;
|
||||
default:
|
||||
G_OBJECT_WARN_INVALID_PROPERTY_ID (object, prop_id, pspec);
|
||||
break;
|
||||
|
|
@ -1555,6 +1729,8 @@ gst_cuda_convert_scale_get_property (GObject * object, guint prop_id,
|
|||
case PROP_CONVERT_SCALE_ADD_BORDERS:
|
||||
g_value_set_boolean (value, base->add_borders);
|
||||
break;
|
||||
case PROP_CONVERT_SCALE_VIDEO_DIRECTION:
|
||||
g_value_set_enum (value, base->method);
|
||||
break;
|
||||
default:
|
||||
G_OBJECT_WARN_INVALID_PROPERTY_ID (object, prop_id, pspec);
|
||||
|
|
@ -1562,6 +1738,76 @@ gst_cuda_convert_scale_get_property (GObject * object, guint prop_id,
|
|||
}
|
||||
}
|
||||
|
||||
static void
|
||||
gst_cuda_convert_scale_before_transform (GstBaseTransform * trans,
|
||||
GstBuffer * buffer)
|
||||
{
|
||||
GstCudaBaseConvert *base = GST_CUDA_BASE_CONVERT (trans);
|
||||
gboolean update = FALSE;
|
||||
GstCaps *in_caps;
|
||||
GstCaps *out_caps;
|
||||
GstBaseTransformClass *klass;
|
||||
|
||||
GST_BASE_TRANSFORM_CLASS (convert_scale_parent_class)->before_transform
|
||||
(trans, buffer);
|
||||
|
||||
g_mutex_lock (&base->lock);
|
||||
if (base->selected_method != base->active_method)
|
||||
update = TRUE;
|
||||
g_mutex_unlock (&base->lock);
|
||||
|
||||
if (!update)
|
||||
return;
|
||||
|
||||
/* basetransform wouldn't call set_caps if in/out caps were not changed.
|
||||
* Update it manually here */
|
||||
GST_DEBUG_OBJECT (base, "Updating caps for direction change");
|
||||
|
||||
in_caps = gst_pad_get_current_caps (GST_BASE_TRANSFORM_SINK_PAD (trans));
|
||||
if (!in_caps) {
|
||||
GST_WARNING_OBJECT (trans, "sinkpad has no current caps");
|
||||
return;
|
||||
}
|
||||
|
||||
out_caps = gst_pad_get_current_caps (GST_BASE_TRANSFORM_SRC_PAD (trans));
|
||||
if (!out_caps) {
|
||||
GST_WARNING_OBJECT (trans, "srcpad has no current caps");
|
||||
gst_caps_unref (in_caps);
|
||||
return;
|
||||
}
|
||||
|
||||
klass = GST_BASE_TRANSFORM_GET_CLASS (trans);
|
||||
klass->set_caps (trans, in_caps, out_caps);
|
||||
gst_caps_unref (in_caps);
|
||||
gst_caps_unref (out_caps);
|
||||
|
||||
gst_base_transform_reconfigure_src (trans);
|
||||
}
|
||||
|
||||
static gboolean
|
||||
gst_cuda_convert_scale_sink_event (GstBaseTransform * trans, GstEvent * event)
|
||||
{
|
||||
GstCudaBaseConvert *base = GST_CUDA_BASE_CONVERT (trans);
|
||||
|
||||
switch (GST_EVENT_TYPE (event)) {
|
||||
case GST_EVENT_TAG:{
|
||||
GstTagList *taglist;
|
||||
GstVideoOrientationMethod method = GST_VIDEO_ORIENTATION_IDENTITY;
|
||||
|
||||
gst_event_parse_tag (event, &taglist);
|
||||
if (gst_video_orientation_from_tag (taglist, &method))
|
||||
gst_cuda_base_convert_set_orientation (base, method, TRUE);
|
||||
break;
|
||||
}
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return
|
||||
GST_BASE_TRANSFORM_CLASS (convert_scale_parent_class)->sink_event
|
||||
(trans, event);
|
||||
}
|
||||
|
||||
/**
|
||||
* SECTION:element-cudaconvert
|
||||
* @title: cudaconvert
|
||||
|
|
|
|||
Loading…
Reference in a new issue