diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.cpp
similarity index 87%
rename from subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
rename to subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.cpp
index 3d1aa11c97..c5ddac2c47 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.cpp
@@ -23,6 +23,7 @@
 
 #include "gstcudaconverter.h"
 #include <string.h>
+#include <mutex>
 
 GST_DEBUG_CATEGORY_STATIC (gst_cuda_converter_debug);
 #define GST_CAT_DEFAULT gst_cuda_converter_debug
@@ -32,13 +33,13 @@ GST_DEBUG_CATEGORY_STATIC (gst_cuda_converter_debug);
 #define DIV_UP(size,block) (((size) + ((block) - 1)) / (block))
 
 /* from GstD3D11 */
-typedef struct _GstCudaColorMatrix
+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)
@@ -598,7 +599,7 @@ gst_cuda_rgb_to_yuv_matrix_unorm (const GstVideoInfo * in_rgb_info,
   return TRUE;
 }
 
-typedef struct
+struct ColorMatrix
 {
   float coeffX[3];
   float coeffY[3];
@@ -606,9 +607,9 @@ typedef struct
   float offset[3];
   float min[3];
   float max[3];
-} ColorMatrix;
+};
 
-typedef struct
+struct ConstBuffer
 {
   ColorMatrix toRGBCoeff;
   ColorMatrix toYuvCoeff;
@@ -620,13 +621,13 @@ typedef struct
   int bottom;
   int view_width;
   int view_height;
-  float offset_x;
-  float offset_y;
   float border_x;
   float border_y;
   float border_z;
   float border_w;
-} ConstBuffer;
+  int fill_border;
+  int video_direction;
+};
 
 #define COLOR_SPACE_IDENTITY "color_space_identity"
 #define COLOR_SPACE_CONVERT "color_space_convert"
@@ -685,14 +686,6 @@ typedef struct
 #define WRITE_GBR_16 "write_gbr_16"
 #define WRITE_GBRA "write_gbra"
 #define WRITE_VUYA "write_vuya"
-#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[] =
@@ -1307,52 +1300,74 @@ WRITE_VUYA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n
 "}\n"
 "\n"
 "__device__ inline float2\n"
-ROTATE_IDENTITY "(float x, float y)\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"
+"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"
+"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"
+"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"
+"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"
+"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"
+"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"
+"rotate_ur_ll (float x, float y)\n"
 "{\n"
 "  return make_float2(1.0 - y, 1.0 - x);\n"
 "}\n"
+"__device__ inline float2\n"
+"do_rotate (float x, float y, int direction)"
+"{\n"
+"  switch (direction) {\n"
+"    case 1:\n"
+"      return rotate_90r (x, y);\n"
+"    case 2:\n"
+"      return rotate_180 (x, y);\n"
+"    case 3:\n"
+"      return rotate_90l (x, y);\n"
+"    case 4:\n"
+"      return rotate_horiz (x, y);\n"
+"    case 5:\n"
+"      return rotate_vert (x, y);\n"
+"    case 6:\n"
+"      return rotate_ul_lr (x, y);\n"
+"    case 7:\n"
+"      return rotate_ur_ll (x, y);\n"
+"    default:\n"
+"      return rotate_identity (x, y);\n"
+"  }\n"
+"}\n"
 "\n";
 
 #define GST_CUDA_KERNEL_UNPACK_FUNC "gst_cuda_kernel_unpack_func"
@@ -1455,12 +1470,12 @@ static const gchar TEMPLATE_KERNEL[] =
 "  int bottom;\n"
 "  int view_width;\n"
 "  int view_height;\n"
-"  float offset_x;\n"
-"  float offset_y;\n"
 "  float border_x;\n"
 "  float border_y;\n"
 "  float border_z;\n"
 "  float border_w;\n"
+"  int fill_border;\n"
+"  int video_direction;\n"
 "};\n"
 "\n"
 "extern \"C\" {\n"
@@ -1468,32 +1483,28 @@ static const gchar TEMPLATE_KERNEL[] =
 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, ConstBuffer * const_buf)\n"
+"    int stride0, int stride1, ConstBuffer * const_buf, int off_x, int off_y)\n"
 "{\n"
-"  int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
-"  int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+"  int x_pos = blockIdx.x * blockDim.x + threadIdx.x + off_x;\n"
+"  int y_pos = blockIdx.y * blockDim.y + threadIdx.y + off_y;\n"
 "  float4 sample;\n"
-"  if (x_pos >= const_buf->width || y_pos >= const_buf->height)\n"
+"  if (x_pos >= const_buf->width || y_pos >= const_buf->height ||\n"
+"      const_buf->view_width <= 0 || const_buf->view_height <= 0)\n"
 "    return;\n"
 "  if (x_pos < const_buf->left || x_pos >= const_buf->right ||\n"
 "      y_pos < const_buf->top || y_pos >= const_buf->bottom) {\n"
+"    if (!const_buf->fill_border)\n"
+"      return;\n"
 "    sample = make_float4 (const_buf->border_x, const_buf->border_y,\n"
 "       const_buf->border_z, const_buf->border_w);\n"
 "  } else {\n"
-"    float x, y;"
-"    if (const_buf->view_width > 0) {\n"
-"      x = const_buf->offset_x +\n"
-"        __int2float_rz (x_pos - const_buf->left) / const_buf->view_width;\n"
-"    } else {\n"
-"      x = const_buf->offset_x;\n"
-"    }\n"
-"    if (const_buf->view_height > 0) {\n"
-"      y = const_buf->offset_y +\n"
-"        __int2float_rz (y_pos - const_buf->top) / const_buf->view_height;\n"
-"    } else {\n"
-"      y = const_buf->offset_y;\n"
-"    }\n"
-"    float2 rotated = %s (x, y);\n"
+"    float x = (__int2float_rz (x_pos - const_buf->left) + 0.5) / const_buf->view_width;\n"
+"    if (x < 0.0 || x > 1.0)\n"
+"      return;\n"
+"    float y = (__int2float_rz (y_pos - const_buf->top) + 0.5) / const_buf->view_height;\n"
+"    if (y < 0.0 || y > 1.0)\n"
+"      return;\n"
+"    float2 rotated = do_rotate (x, y, const_buf->video_direction);\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, &const_buf->toRGBCoeff);\n"
@@ -1513,7 +1524,7 @@ typedef struct _TextureFormat
   const gchar *sample_func;
 } TextureFormat;
 
-#define CU_AD_FORMAT_NONE 0
+#define CU_AD_FORMAT_NONE ((CUarray_format)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 }
@@ -1564,60 +1575,111 @@ static const TextureFormat format_map[] = {
   MAKE_FORMAT_RGB (VUYA, UNSIGNED_INT8, SAMPLE_VUYA),
 };
 
-typedef struct _TextureBuffer
+struct TextureBuffer
 {
-  CUdeviceptr ptr;
-  gsize stride;
-  CUtexObject texture;
-} TextureBuffer;
+  CUdeviceptr ptr = 0;
+  gsize stride = 0;
+  CUtexObject texture = 0;
+};
 
-typedef struct
+enum
 {
-  gint x;
-  gint y;
-  gint width;
-  gint height;
-} ConverterRect;
+  PROP_0,
+  PROP_DEST_X,
+  PROP_DEST_Y,
+  PROP_DEST_WIDTH,
+  PROP_DEST_HEIGHT,
+  PROP_FILL_BORDER,
+  PROP_VIDEO_DIRECTION,
+};
 
 struct _GstCudaConverterPrivate
 {
+  _GstCudaConverterPrivate ()
+  {
+    config = gst_structure_new_empty ("converter-config");
+  }
+
+   ~_GstCudaConverterPrivate ()
+  {
+    if (config)
+      gst_structure_free (config);
+  }
+
+  std::mutex lock;
+
   GstVideoInfo in_info;
   GstVideoInfo out_info;
 
-  GstVideoOrientationMethod method;
-
-  GstStructure *config;
+  GstStructure *config = nullptr;
 
   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;
-  ConstBuffer *const_buf_staging;
-  CUdeviceptr const_buf;
+  ConstBuffer *const_buf_staging = nullptr;
+  CUdeviceptr const_buf = 0;
 
-  CUmodule module;
-  CUfunction main_func;
-  CUfunction unpack_func;
+  CUmodule module = nullptr;
+  CUfunction main_func = nullptr;
+  CUfunction unpack_func = nullptr;
+
+  gboolean update_const_buf = FALSE;
+
+  /* properties */
+  gint dest_x = 0;
+  gint dest_y = 0;
+  gint dest_width = 0;
+  gint dest_height = 0;
+  GstVideoOrientationMethod video_direction = GST_VIDEO_ORIENTATION_IDENTITY;
+  gboolean fill_border = FALSE;
+  CUfilter_mode filter_mode = CU_TR_FILTER_MODE_LINEAR;
 };
 
 static void gst_cuda_converter_dispose (GObject * object);
 static void gst_cuda_converter_finalize (GObject * object);
+static void gst_cuda_converter_set_property (GObject * object, guint prop_id,
+    const GValue * value, GParamSpec * pspec);
+static void gst_cuda_converter_get_property (GObject * object, guint prop_id,
+    GValue * value, GParamSpec * pspec);
 
 #define gst_cuda_converter_parent_class parent_class
-G_DEFINE_TYPE_WITH_PRIVATE (GstCudaConverter, gst_cuda_converter,
-    GST_TYPE_OBJECT);
+G_DEFINE_TYPE (GstCudaConverter, gst_cuda_converter, GST_TYPE_OBJECT);
 
 static void
 gst_cuda_converter_class_init (GstCudaConverterClass * klass)
 {
-  GObjectClass *object_class = G_OBJECT_CLASS (klass);
+  auto object_class = G_OBJECT_CLASS (klass);
+  auto param_flags = (GParamFlags) (G_PARAM_READWRITE | G_PARAM_STATIC_STRINGS);
 
   object_class->dispose = gst_cuda_converter_dispose;
   object_class->finalize = gst_cuda_converter_finalize;
+  object_class->set_property = gst_cuda_converter_set_property;
+  object_class->get_property = gst_cuda_converter_get_property;
+
+  g_object_class_install_property (object_class, PROP_DEST_X,
+      g_param_spec_int ("dest-x", "Dest-X",
+          "x poisition in the destination frame", G_MININT, G_MAXINT, 0,
+          param_flags));
+  g_object_class_install_property (object_class, PROP_DEST_Y,
+      g_param_spec_int ("dest-y", "Dest-Y",
+          "y poisition in the destination frame", G_MININT, G_MAXINT, 0,
+          param_flags));
+  g_object_class_install_property (object_class, PROP_DEST_WIDTH,
+      g_param_spec_int ("dest-width", "Dest-Width",
+          "Width in the destination frame", 0, G_MAXINT, 0, param_flags));
+  g_object_class_install_property (object_class, PROP_DEST_HEIGHT,
+      g_param_spec_int ("dest-height", "Dest-Height",
+          "Height in the destination frame", 0, G_MAXINT, 0, param_flags));
+  g_object_class_install_property (object_class, PROP_FILL_BORDER,
+      g_param_spec_boolean ("fill-border", "Fill border",
+          "Fill border", FALSE, param_flags));
+  g_object_class_install_property (object_class, PROP_VIDEO_DIRECTION,
+      g_param_spec_enum ("video-direction", "Video Direction",
+          "Video direction", GST_TYPE_VIDEO_ORIENTATION_METHOD,
+          GST_VIDEO_ORIENTATION_IDENTITY, param_flags));
 
   GST_DEBUG_CATEGORY_INIT (gst_cuda_converter_debug,
       "cudaconverter", 0, "cudaconverter");
@@ -1626,26 +1688,22 @@ gst_cuda_converter_class_init (GstCudaConverterClass * klass)
 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");
+  self->priv = new GstCudaConverterPrivate ();
 }
 
 static void
 gst_cuda_converter_dispose (GObject * object)
 {
-  GstCudaConverter *self = GST_CUDA_CONVERTER (object);
-  GstCudaConverterPrivate *priv = self->priv;
-  guint i;
+  auto self = GST_CUDA_CONVERTER (object);
+  auto priv = self->priv;
 
   if (self->context && gst_cuda_context_push (self->context)) {
     if (priv->module) {
       CuModuleUnload (priv->module);
-      priv->module = NULL;
+      priv->module = nullptr;
     }
 
-    for (i = 0; i < G_N_ELEMENTS (priv->fallback_buffer); i++) {
+    for (guint 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);
@@ -1669,7 +1727,7 @@ gst_cuda_converter_dispose (GObject * object)
 
     if (priv->const_buf_staging) {
       CuMemFreeHost (priv->const_buf_staging);
-      priv->const_buf_staging = NULL;
+      priv->const_buf_staging = nullptr;
     }
 
     if (priv->const_buf) {
@@ -1677,7 +1735,7 @@ gst_cuda_converter_dispose (GObject * object)
       priv->const_buf = 0;
     }
 
-    gst_cuda_context_pop (NULL);
+    gst_cuda_context_pop (nullptr);
   }
 
   gst_clear_object (&self->context);
@@ -1688,14 +1746,126 @@ gst_cuda_converter_dispose (GObject * object)
 static void
 gst_cuda_converter_finalize (GObject * object)
 {
-  GstCudaConverter *self = GST_CUDA_CONVERTER (object);
-  GstCudaConverterPrivate *priv = self->priv;
+  auto self = GST_CUDA_CONVERTER (object);
 
-  gst_structure_free (priv->config);
+  delete self->priv;
 
   G_OBJECT_CLASS (parent_class)->finalize (object);
 }
 
+static void
+gst_cuda_converter_set_property (GObject * object, guint prop_id,
+    const GValue * value, GParamSpec * pspec)
+{
+  auto self = GST_CUDA_CONVERTER (object);
+  auto priv = self->priv;
+
+  std::lock_guard < std::mutex > lk (priv->lock);
+  switch (prop_id) {
+    case PROP_DEST_X:
+    {
+      auto dest_x = g_value_get_int (value);
+      if (priv->dest_x != dest_x) {
+        priv->update_const_buf = TRUE;
+        priv->dest_x = dest_x;
+        priv->const_buf_staging->left = dest_x;
+        priv->const_buf_staging->right = priv->dest_x + priv->dest_width;
+      }
+      break;
+    }
+    case PROP_DEST_Y:
+    {
+      auto dest_y = g_value_get_int (value);
+      if (priv->dest_y != dest_y) {
+        priv->update_const_buf = TRUE;
+        priv->dest_y = dest_y;
+        priv->const_buf_staging->top = dest_y;
+        priv->const_buf_staging->bottom = priv->dest_y + priv->dest_height;
+      }
+      break;
+    }
+    case PROP_DEST_WIDTH:
+    {
+      auto dest_width = g_value_get_int (value);
+      if (priv->dest_width != dest_width) {
+        priv->update_const_buf = TRUE;
+        priv->dest_width = dest_width;
+        priv->const_buf_staging->right = priv->dest_x + dest_width;
+        priv->const_buf_staging->view_width = dest_width;
+      }
+      break;
+    }
+    case PROP_DEST_HEIGHT:
+    {
+      auto dest_height = g_value_get_int (value);
+      if (priv->dest_height != dest_height) {
+        priv->update_const_buf = TRUE;
+        priv->dest_height = dest_height;
+        priv->const_buf_staging->bottom = priv->dest_y + dest_height;
+        priv->const_buf_staging->view_height = dest_height;
+      }
+      break;
+    }
+    case PROP_FILL_BORDER:
+    {
+      auto fill_border = g_value_get_boolean (value);
+      if (priv->fill_border != fill_border) {
+        priv->update_const_buf = TRUE;
+        priv->fill_border = fill_border;
+        priv->const_buf_staging->fill_border = fill_border;
+      }
+      break;
+    }
+    case PROP_VIDEO_DIRECTION:
+    {
+      auto video_direction =
+          (GstVideoOrientationMethod) g_value_get_enum (value);
+      if (priv->video_direction != video_direction) {
+        priv->update_const_buf = TRUE;
+        priv->video_direction = video_direction;
+        priv->const_buf_staging->video_direction = video_direction;
+      }
+      break;
+    }
+    default:
+      G_OBJECT_WARN_INVALID_PROPERTY_ID (object, prop_id, pspec);
+      break;
+  }
+}
+
+static void
+gst_cuda_converter_get_property (GObject * object, guint prop_id,
+    GValue * value, GParamSpec * pspec)
+{
+  auto self = GST_CUDA_CONVERTER (object);
+  auto priv = self->priv;
+
+  std::lock_guard < std::mutex > lk (priv->lock);
+  switch (prop_id) {
+    case PROP_DEST_X:
+      g_value_set_int (value, priv->dest_x);
+      break;
+    case PROP_DEST_Y:
+      g_value_set_int (value, priv->dest_y);
+      break;
+    case PROP_DEST_WIDTH:
+      g_value_set_int (value, priv->dest_width);
+      break;
+    case PROP_DEST_HEIGHT:
+      g_value_set_int (value, priv->dest_height);
+      break;
+    case PROP_FILL_BORDER:
+      g_value_set_boolean (value, priv->fill_border);
+      break;
+    case PROP_VIDEO_DIRECTION:
+      g_value_set_enum (value, priv->video_direction);
+      break;
+    default:
+      G_OBJECT_WARN_INVALID_PROPERTY_ID (object, prop_id, pspec);
+      break;
+  }
+}
+
 static const gchar *
 get_color_range_name (GstVideoColorRange range)
 {
@@ -1722,16 +1892,15 @@ gst_cuda_converter_setup (GstCudaConverter * self)
   GstCudaColorMatrix to_yuv_matrix;
   GstCudaColorMatrix border_color_matrix;
   gdouble border_color[4];
-  gint i, j;
-  const gchar *unpack_function = NULL;
-  const gchar *write_func = NULL;
+  guint i, j;
+  const gchar *unpack_function = nullptr;
+  const gchar *write_func = nullptr;
   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;
-  gchar *program = NULL;
+  gchar *program = nullptr;
   CUresult ret;
 
   in_info = &priv->in_info;
@@ -2014,51 +2183,21 @@ gst_cuda_converter_setup (GstCudaConverter * self)
 
   priv->const_buf_staging->width = out_info->width;
   priv->const_buf_staging->height = out_info->height;
-  priv->const_buf_staging->left = priv->dest_rect.x;
-  priv->const_buf_staging->top = priv->dest_rect.y;
-  priv->const_buf_staging->right = priv->dest_rect.x + priv->dest_rect.width;
-  priv->const_buf_staging->bottom = priv->dest_rect.y + priv->dest_rect.height;
-  priv->const_buf_staging->view_width = priv->dest_rect.width;
-  priv->const_buf_staging->view_height = priv->dest_rect.height;
-  priv->const_buf_staging->offset_x = priv->dest_rect.width != 0 ?
-      0.5 / priv->dest_rect.width : 0;
-  priv->const_buf_staging->offset_y = priv->dest_rect.height != 0 ?
-      0.5 / priv->dest_rect.height : 0;
+  priv->const_buf_staging->left = 0;
+  priv->const_buf_staging->top = 0;
+  priv->const_buf_staging->right = out_info->width;
+  priv->const_buf_staging->bottom = out_info->height;
+  priv->const_buf_staging->view_width = out_info->width;
+  priv->const_buf_staging->view_height = out_info->height;
   priv->const_buf_staging->border_x = border_color[0];
   priv->const_buf_staging->border_y = border_color[1];
   priv->const_buf_staging->border_z = border_color[2];
   priv->const_buf_staging->border_w = border_color[3];
-
-  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;
-  }
+  priv->const_buf_staging->fill_border = 0;
+  priv->const_buf_staging->video_direction = 0;
 
   str = g_strdup_printf (TEMPLATE_KERNEL, KERNEL_COMMON,
       unpack_function ? unpack_function : "",
-      /* adjust coord before sampling */
-      rotate_func,
       /* sampler function name */
       priv->texture_fmt->sample_func,
       /* TO RGB conversion function name */
@@ -2070,7 +2209,7 @@ gst_cuda_converter_setup (GstCudaConverter * self)
 
   GST_LOG_OBJECT (self, "kernel code:\n%s\n", str);
   gint cuda_device;
-  g_object_get (self->context, "cuda-device-id", &cuda_device, NULL);
+  g_object_get (self->context, "cuda-device-id", &cuda_device, nullptr);
   program = gst_cuda_nvrtc_compile_cubin (str, cuda_device);
   if (!program) {
     GST_WARNING_OBJECT (self, "Couldn't compile to cubin, trying ptx");
@@ -2083,18 +2222,6 @@ gst_cuda_converter_setup (GstCudaConverter * self)
     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");
     g_free (program);
@@ -2128,13 +2255,13 @@ gst_cuda_converter_setup (GstCudaConverter * self)
     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.filterMode = priv->filter_mode;
     texture_desc.flags = 0x2;
-    texture_desc.addressMode[0] = 1;
-    texture_desc.addressMode[1] = 1;
-    texture_desc.addressMode[2] = 1;
+    texture_desc.addressMode[0] = (CUaddress_mode) 1;
+    texture_desc.addressMode[1] = (CUaddress_mode) 1;
+    texture_desc.addressMode[2] = (CUaddress_mode) 1;
 
-    ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, NULL);
+    ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, nullptr);
     if (!gst_cuda_result (ret)) {
       GST_ERROR_OBJECT (self, "Couldn't create unpack texture");
       goto error;
@@ -2147,7 +2274,7 @@ gst_cuda_converter_setup (GstCudaConverter * self)
   g_clear_pointer (&program, g_free);
   if (!gst_cuda_result (ret)) {
     GST_ERROR_OBJECT (self, "Could not load module");
-    priv->module = NULL;
+    priv->module = nullptr;
     goto error;
   }
 
@@ -2174,12 +2301,12 @@ gst_cuda_converter_setup (GstCudaConverter * self)
     goto error;
   }
 
-  gst_cuda_context_pop (NULL);
+  gst_cuda_context_pop (nullptr);
 
   return TRUE;
 
 error:
-  gst_cuda_context_pop (NULL);
+  gst_cuda_context_pop (nullptr);
   g_free (program);
 
   return FALSE;
@@ -2203,15 +2330,6 @@ gst_cuda_converter_set_config (GstCudaConverter * self, GstStructure * config)
   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,
@@ -2219,41 +2337,29 @@ gst_cuda_converter_new (const GstVideoInfo * in_info,
 {
   GstCudaConverter *self;
   GstCudaConverterPrivate *priv;
-  gint method;
   CUresult cuda_ret;
 
-  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);
+  g_return_val_if_fail (in_info != nullptr, nullptr);
+  g_return_val_if_fail (out_info != nullptr, nullptr);
+  g_return_val_if_fail (GST_IS_CUDA_CONTEXT (context), nullptr);
 
-  self = g_object_new (GST_TYPE_CUDA_CONVERTER, NULL);
+  self = (GstCudaConverter *) g_object_new (GST_TYPE_CUDA_CONVERTER, nullptr);
 
   if (!GST_IS_CUDA_CONTEXT (context)) {
     GST_WARNING_OBJECT (self, "Not a valid cuda context object");
     goto error;
   }
 
-  self->context = gst_object_ref (context);
+  self->context = (GstCudaContext *) gst_object_ref (context);
   priv = self->priv;
   priv->in_info = *in_info;
   priv->out_info = *out_info;
+  priv->dest_width = out_info->width;
+  priv->dest_height = out_info->height;
 
   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_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_context_push (context)) {
     GST_ERROR_OBJECT (self, "Couldn't push context");
     goto error;
@@ -2263,12 +2369,12 @@ gst_cuda_converter_new (const GstVideoInfo * in_info,
       sizeof (ConstBuffer));
   if (!gst_cuda_result (cuda_ret)) {
     GST_ERROR_OBJECT (self, "Couldn't allocate staging const buf");
-    gst_cuda_context_pop (NULL);
+    gst_cuda_context_pop (nullptr);
     goto error;
   }
 
   cuda_ret = CuMemAlloc (&priv->const_buf, sizeof (ConstBuffer));
-  gst_cuda_context_pop (NULL);
+  gst_cuda_context_pop (nullptr);
   if (!gst_cuda_result (cuda_ret)) {
     GST_ERROR_OBJECT (self, "Couldn't allocate const buf");
     goto error;
@@ -2284,10 +2390,9 @@ gst_cuda_converter_new (const GstVideoInfo * in_info,
 
 error:
   gst_object_unref (self);
-  return NULL;
+  return nullptr;
 }
 
-
 static CUtexObject
 gst_cuda_converter_create_texture_unchecked (GstCudaConverter * self,
     CUdeviceptr src, gint width, gint height, CUarray_format format,
@@ -2315,11 +2420,12 @@ gst_cuda_converter_create_texture_unchecked (GstCudaConverter * self,
   /* 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;
+  texture_desc.addressMode[0] = (CUaddress_mode) 1;
+  texture_desc.addressMode[1] = (CUaddress_mode) 1;
+  texture_desc.addressMode[2] = (CUaddress_mode) 1;
 
-  cuda_ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, NULL);
+  cuda_ret =
+      CuTexObjectCreate (&texture, &resource_desc, &texture_desc, nullptr);
 
   if (!gst_cuda_result (cuda_ret)) {
     GST_ERROR_OBJECT (self, "Could not create texture");
@@ -2415,7 +2521,7 @@ gst_cuda_converter_unpack_rgb (GstCudaConverter * self,
 
   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);
+      stream, args, nullptr);
 
   if (!gst_cuda_result (ret)) {
     GST_ERROR_OBJECT (self, "Couldn't unpack source RGB");
@@ -2433,27 +2539,44 @@ gst_cuda_converter_convert_frame (GstCudaConverter * converter,
   GstCudaConverterPrivate *priv;
   const TextureFormat *format;
   CUtexObject texture[GST_VIDEO_MAX_COMPONENTS] = { 0, };
-  guint8 *dst[GST_VIDEO_MAX_COMPONENTS] = { NULL, };
+  guint8 *dst[GST_VIDEO_MAX_COMPONENTS] = { nullptr, };
   gint stride[2] = { 0, };
-  gint i;
+  guint i;
   gboolean ret = FALSE;
   CUresult cuda_ret;
   gint width, height;
   gboolean need_sync = FALSE;
   GstCudaMemory *cmem;
+  gint off_x = 0;
+  gint off_y = 0;
 
   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);
+  g_return_val_if_fail (src_frame != nullptr, FALSE);
+  g_return_val_if_fail (dst_frame != nullptr, FALSE);
 
   priv = converter->priv;
   format = priv->texture_fmt;
 
   g_assert (format);
 
+  std::lock_guard < std::mutex > lk (priv->lock);
+  if (!priv->fill_border && (priv->dest_width <= 0 || priv->dest_height <= 0))
+    return TRUE;
+
+  if (priv->update_const_buf) {
+    priv->update_const_buf = FALSE;
+    cuda_ret = CuMemcpyHtoDAsync (priv->const_buf, priv->const_buf_staging,
+        sizeof (ConstBuffer), stream);
+
+    if (!gst_cuda_result (cuda_ret)) {
+      GST_ERROR_OBJECT (converter, "Couldn't upload const buffer");
+      return FALSE;
+    }
+  }
+
   gpointer args[] = { &texture[0], &texture[1], &texture[2], &texture[3],
     &dst[0], &dst[1], &dst[2], &dst[3], &stride[0], &stride[1],
-    &priv->const_buf
+    &priv->const_buf, &off_x, &off_y
   };
 
   cmem = (GstCudaMemory *) gst_buffer_peek_memory (src_frame->buffer, 0);
@@ -2476,14 +2599,14 @@ gst_cuda_converter_convert_frame (GstCudaConverter * converter,
   } 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])) {
+              i, priv->filter_mode, &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],
+            priv->filter_mode, format->array_format[i], format->channels[i],
             i, stream);
         need_sync = TRUE;
       }
@@ -2498,8 +2621,20 @@ gst_cuda_converter_convert_frame (GstCudaConverter * converter,
   width = GST_VIDEO_FRAME_WIDTH (dst_frame);
   height = GST_VIDEO_FRAME_HEIGHT (dst_frame);
 
+  if (!priv->fill_border) {
+    if (priv->dest_width < width) {
+      off_x = priv->dest_x;
+      width = priv->dest_width;
+    }
+
+    if (priv->dest_height < height) {
+      off_y = priv->dest_y;
+      height = priv->dest_height;
+    }
+  }
+
   for (i = 0; i < GST_VIDEO_FRAME_N_PLANES (dst_frame); i++)
-    dst[i] = GST_VIDEO_FRAME_PLANE_DATA (dst_frame, i);
+    dst[i] = (guint8 *) 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)
@@ -2507,7 +2642,7 @@ gst_cuda_converter_convert_frame (GstCudaConverter * converter,
 
   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);
+      stream, args, nullptr);
 
   if (!gst_cuda_result (cuda_ret)) {
     GST_ERROR_OBJECT (converter, "Couldn't convert frame");
@@ -2523,6 +2658,6 @@ gst_cuda_converter_convert_frame (GstCudaConverter * converter,
   ret = TRUE;
 
 out:
-  gst_cuda_context_pop (NULL);
+  gst_cuda_context_pop (nullptr);
   return ret;
 }
diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
index ddf54519ed..fa446a5b1b 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
@@ -36,41 +36,6 @@ typedef struct _GstCudaConverter GstCudaConverter;
 typedef struct _GstCudaConverterClass GstCudaConverterClass;
 typedef struct _GstCudaConverterPrivate GstCudaConverterPrivate;
 
-/**
- * GST_CUDA_CONVERTER_OPT_DEST_X:
- *
- * #G_TYPE_INT, x position in the destination frame, default 0
- */
-#define GST_CUDA_CONVERTER_OPT_DEST_X   "GstCudaConverter.dest-x"
-
-/**
- * GST_CUDA_CONVERTER_OPT_DEST_Y:
- *
- * #G_TYPE_INT, y position in the destination frame, default 0
- */
-#define GST_CUDA_CONVERTER_OPT_DEST_Y   "GstCudaConverter.dest-y"
-
-/**
- * GST_CUDA_CONVERTER_OPT_DEST_WIDTH:
- *
- * #G_TYPE_INT, width in the destination frame, default destination width
- */
-#define GST_CUDA_CONVERTER_OPT_DEST_WIDTH   "GstCudaConverter.dest-width"
-
-/**
- * GST_CUDA_CONVERTER_OPT_DEST_HEIGHT:
- *
- * #G_TYPE_INT, height in the destination frame, default destination height
- */
-#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 6767cb8e0c..3bb4b95d8a 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
@@ -1415,27 +1415,20 @@ gst_cuda_base_convert_set_info (GstCudaBaseTransform * btrans,
       !needs_color_convert (in_info, out_info)) {
     gst_base_transform_set_passthrough (GST_BASE_TRANSFORM (self), TRUE);
   } else {
-    GstStructure *config;
-
     gst_base_transform_set_passthrough (GST_BASE_TRANSFORM (self), FALSE);
 
-    config = gst_structure_new_empty ("GstCudaConverter");
-    gst_structure_set (config,
-        GST_CUDA_CONVERTER_OPT_DEST_X, G_TYPE_INT, self->borders_w / 2,
-        GST_CUDA_CONVERTER_OPT_DEST_Y, G_TYPE_INT, self->borders_h / 2,
-        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,
-        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);
+        out_info, btrans->context, NULL);
     if (!self->converter) {
       GST_ERROR_OBJECT (self, "Couldn't create converter");
       return FALSE;
     }
+
+    g_object_set (self->converter, "dest-x", self->borders_w / 2,
+        "dest-y", self->borders_h / 2,
+        "dest-width", out_info->width - self->borders_w,
+        "dest-height", out_info->height - self->borders_h,
+        "fill-border", TRUE, "video-direction", active_method, NULL);
   }
 
   GST_DEBUG_OBJECT (self, "%s from=%dx%d (par=%d/%d dar=%d/%d), size %"
diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/meson.build b/subprojects/gst-plugins-bad/sys/nvcodec/meson.build
index fd8c59678e..b7d69ab6cf 100644
--- a/subprojects/gst-plugins-bad/sys/nvcodec/meson.build
+++ b/subprojects/gst-plugins-bad/sys/nvcodec/meson.build
@@ -1,6 +1,6 @@
 nvcodec_sources = [
   'gstcudabasetransform.c',
-  'gstcudaconverter.c',
+  'gstcudaconverter.cpp',
   'gstcudaconvertscale.c',
   'gstcudaipc.cpp',
   'gstcudaipcclient.cpp',