mirrors/gstreamer
mirrors/gstreamer
1
1
Fork 0
mirror of https://gitlab.freedesktop.org/gstreamer/gstreamer.git synced 2024-05-17 07:52:47 +00:00
Code Issues Projects Releases Wiki Activity

cuda: Add support for VUYA format

Browse source 
Part-of: <https://gitlab.freedesktop.org/gstreamer/gstreamer/-/merge_requests/6417>
This commit is contained in:
Seungha Yang 2024-03-22 01:19:53 +09:00 committed by GStreamer Marge Bot
parent 707ac69f5f
commit e6f496a240
4 changed files with 30 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp
View file

@ -264,6 +264,7 @@ gst_cuda_allocator_update_info (const GstVideoInfo * reference,
case GST_VIDEO_FORMAT_RGB10A2_LE:
case GST_VIDEO_FORMAT_YUY2:
case GST_VIDEO_FORMAT_UYVY:
case GST_VIDEO_FORMAT_VUYA:
ret.stride[0] = pitch;
ret.offset[0] = 0;
break;

26
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c
View file

@ -638,6 +638,7 @@ typedef struct
#define SAMPLE_GBR_10 "sample_gbr_10"
#define SAMPLE_GBR_12 "sample_gbr_12"
#define SAMPLE_GBRA "sample_gbra"
#define SAMPLE_VUYA "sample_vuya"
#define WRITE_I420 "write_i420"
#define WRITE_YV12 "write_yv12"
@ -670,6 +671,7 @@ typedef struct
#define WRITE_GBR_12 "write_gbr_12"
#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"
@ -931,6 +933,14 @@ SAMPLE_GBRA "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n"
" return make_float4 (r, g, b, a);\n"
"}\n"
"\n"
"__device__ inline float4\n"
SAMPLE_VUYA "(cudaTextureObject_t tex0, cudaTextureObject_t tex1,\n"
" cudaTextureObject_t tex2, cudaTextureObject_t tex3, float x, float y)\n"
"{\n"
" float4 vuya = tex2D<float4>(tex0, x, y);\n"
" return make_float4 (vuya.z, vuya.y, vuya.x, vuya.w);\n"
"}\n"
"\n"
"__device__ inline void\n"
WRITE_I420 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n"
" unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n"
@ -1271,6 +1281,18 @@ WRITE_GBRA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n
" dst2[pos] = scale_to_uchar (sample.x);\n"
" dst3[pos] = scale_to_uchar (sample.w);\n"
"}\n"
"\n"
"__device__ inline void\n"
WRITE_VUYA "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n"
" unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n"
"{\n"
" int pos = x * 4 + y * stride0;\n"
" dst0[pos] = scale_to_uchar (sample.z);\n"
" dst0[pos + 1] = scale_to_uchar (sample.y);\n"
" dst0[pos + 2] = scale_to_uchar (sample.x);\n"
" dst0[pos + 3] = scale_to_uchar (sample.w);\n"
"}\n"
"\n"
"__device__ inline float2\n"
ROTATE_IDENTITY "(float x, float y)\n"
"{\n"
@ -1520,6 +1542,7 @@ static const TextureFormat format_map[] = {
MAKE_FORMAT_RGBP (GBR_12LE, UNSIGNED_INT16, SAMPLE_GBR_12),
MAKE_FORMAT_RGBP (GBR_16LE, UNSIGNED_INT16, SAMPLE_GBR),
MAKE_FORMAT_RGBAP (GBRA, UNSIGNED_INT8, SAMPLE_GBRA),
MAKE_FORMAT_RGB (VUYA, UNSIGNED_INT8, SAMPLE_VUYA),
};
typedef struct _TextureBuffer
@ -1819,6 +1842,9 @@ gst_cuda_converter_setup (GstCudaConverter * self)
case GST_VIDEO_FORMAT_GBRA:
write_func = WRITE_GBRA;
break;
case GST_VIDEO_FORMAT_VUYA:
write_func = WRITE_VUYA;
break;
default:
break;
}

2
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c
View file

@ -33,7 +33,7 @@ GST_DEBUG_CATEGORY_STATIC (gst_cuda_base_convert_debug);
"{ I420, YV12, NV12, NV21, P010_10LE, P012_LE, P016_LE, I420_10LE, I420_12LE, Y444, " \
"Y444_10LE, Y444_12LE, Y444_16LE, BGRA, RGBA, RGBx, BGRx, ARGB, ABGR, RGB, " \
"BGR, BGR10A2_LE, RGB10A2_LE, Y42B, I422_10LE, I422_12LE, RGBP, BGRP, GBR, " \
"GBRA, GBR_10LE, GBR_12LE, GBR_16LE }"
"GBRA, GBR_10LE, GBR_12LE, GBR_16LE, VUYA }"
static GstStaticPadTemplate sink_template = GST_STATIC_PAD_TEMPLATE ("sink",
GST_PAD_SINK,

4
subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h
View file

@ -27,12 +27,12 @@ G_BEGIN_DECLS
"{ I420, YV12, NV12, NV21, P010_10LE, P012_LE, P016_LE, I420_10LE, I420_12LE, Y444, " \
"Y444_10LE, Y444_12LE, Y444_16LE, BGRA, RGBA, RGBx, BGRx, ARGB, ABGR, RGB, " \
"BGR, BGR10A2_LE, RGB10A2_LE, Y42B, I422_10LE, I422_12LE, YUY2, UYVY, RGBP, " \
"BGRP, GBR, GBR_10LE, GBR_12LE, GBR_16LE, GBRA }"
"BGRP, GBR, GBR_10LE, GBR_12LE, GBR_16LE, GBRA, VUYA }"
#define GST_CUDA_GL_FORMATS \
"{ I420, YV12, NV12, NV21, P010_10LE, P012_LE, P016_LE, Y444, " \
"BGRA, RGBA, RGBx, BGRx, ARGB, ABGR, RGB, BGR, BGR10A2_LE, RGB10A2_LE, " \
"YUY2, UYVY, RGBP, BGRP, GBR, GBRA }"
"YUY2, UYVY, RGBP, BGRP, GBR, GBRA, VUYA }"
#define GST_CUDA_D3D11_FORMATS \
"{ I420, YV12, I420_10LE, I420_12LE, Y444, Y444_10LE, Y444_12LE, Y444_16LE, " \