FaSTTUBe/GST-Tensordecoder-ov_ep
37
0
Fork 0
Code Issues Pull Requests 1 Actions Packages Projects Releases Wiki Activity

cudaconverter: Add support for configuration update

Browse Source 
Allow updating various configuration values via property

Part-of: <https://gitlab.freedesktop.org/gstreamer/gstreamer/-/merge_requests/8170>
This commit is contained in:
Seungha Yang 2024-12-16 01:32:36 +09:00
parent 1fa51046d1
commit d761196bb7
4 changed files with 337 additions and 244 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

523
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c → subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.cpp
View File

@ -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;
}

35
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.h
View File

@ -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;

21
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
View File

@ -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 %"

2
subprojects/gst-plugins-bad/sys/nvcodec/meson.build
View File

@ -1,6 +1,6 @@
nvcodec_sources = [
'gstcudabasetransform.c',
'gstcudaconverter.c',
'gstcudaconverter.cpp',
'gstcudaconvertscale.c',
'gstcudaipc.cpp',
'gstcudaipcclient.cpp',