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>

3 files changed, 368 insertions, 20 deletions

diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
index dc334ca513..d4120e1ab5 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
@@ -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;
diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
index 1be9d8a9b3..ddf54519ed 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
@@ -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;
diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
index 829eb648c3..cdb7868d72 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
@@ -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;
 }
@@ -1242,19 +1299,71 @@ gst_cuda_base_convert_decide_allocation (GstBaseTransform * trans,
 }
 
 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,
     GstVideoInfo * out_info)
 {
   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