mirrors/gstreamer
mirrors/gstreamer
1
1
Fork 0
mirror of https://gitlab.freedesktop.org/gstreamer/gstreamer.git synced 2024-11-26 03:31:05 +00:00
Code Issues 1 Projects Releases Wiki Activity

cuda: Add support for I420_12LE format

Browse source 
Part-of: <https://gitlab.freedesktop.org/gstreamer/gstreamer/-/merge_requests/5409>
This commit is contained in:
Seungha Yang 2023-09-28 00:11:53 +09:00 committed by GStreamer Marge Bot
parent e7fc0400bc
commit c818906236
4 changed files with 24 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -195,6 +195,7 @@ gst_cuda_allocator_update_info (const GstVideoInfo * reference,
case GST_VIDEO_FORMAT_I420: case GST_VIDEO_FORMAT_I420:
case GST_VIDEO_FORMAT_YV12: case GST_VIDEO_FORMAT_YV12:
case GST_VIDEO_FORMAT_I420_10LE: case GST_VIDEO_FORMAT_I420_10LE:
case GST_VIDEO_FORMAT_I420_12LE:
/* we are wasting space yes, but required so that this memory /* we are wasting space yes, but required so that this memory
* can be used in kernel function */ * can be used in kernel function */
ret.stride[0] = pitch; ret.stride[0] = pitch;
@ -738,6 +739,7 @@ static const TextureFormat format_map[] = {
MAKE_FORMAT_YUV_SEMI_PLANAR (P012_LE, UNSIGNED_INT16), MAKE_FORMAT_YUV_SEMI_PLANAR (P012_LE, UNSIGNED_INT16),
MAKE_FORMAT_YUV_SEMI_PLANAR (P016_LE, UNSIGNED_INT16), MAKE_FORMAT_YUV_SEMI_PLANAR (P016_LE, UNSIGNED_INT16),
MAKE_FORMAT_YUV_PLANAR (I420_10LE, UNSIGNED_INT16), MAKE_FORMAT_YUV_PLANAR (I420_10LE, UNSIGNED_INT16),
MAKE_FORMAT_YUV_PLANAR (I420_12LE, UNSIGNED_INT16),
MAKE_FORMAT_YUV_PLANAR (Y444, UNSIGNED_INT8), MAKE_FORMAT_YUV_PLANAR (Y444, UNSIGNED_INT8),
MAKE_FORMAT_YUV_PLANAR (Y444_10LE, UNSIGNED_INT16), MAKE_FORMAT_YUV_PLANAR (Y444_10LE, UNSIGNED_INT16),
MAKE_FORMAT_YUV_PLANAR (Y444_12LE, UNSIGNED_INT16), MAKE_FORMAT_YUV_PLANAR (Y444_12LE, UNSIGNED_INT16),
@ -1028,6 +1030,7 @@ gst_cuda_allocator_calculate_alloc_height (const GstVideoInfo * info)
case GST_VIDEO_FORMAT_P012_LE: case GST_VIDEO_FORMAT_P012_LE:
case GST_VIDEO_FORMAT_P016_LE: case GST_VIDEO_FORMAT_P016_LE:
case GST_VIDEO_FORMAT_I420_10LE: case GST_VIDEO_FORMAT_I420_10LE:
case GST_VIDEO_FORMAT_I420_12LE:
alloc_height = GST_ROUND_UP_2 (alloc_height); alloc_height = GST_ROUND_UP_2 (alloc_height);
break; break;
default: default:
@ -1038,6 +1041,7 @@ gst_cuda_allocator_calculate_alloc_height (const GstVideoInfo * info)
case GST_VIDEO_FORMAT_I420: case GST_VIDEO_FORMAT_I420:
case GST_VIDEO_FORMAT_YV12: case GST_VIDEO_FORMAT_YV12:
case GST_VIDEO_FORMAT_I420_10LE: case GST_VIDEO_FORMAT_I420_10LE:
case GST_VIDEO_FORMAT_I420_12LE:
alloc_height *= 2; alloc_height *= 2;
break; break;
case GST_VIDEO_FORMAT_NV12: case GST_VIDEO_FORMAT_NV12:

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

@ -645,6 +645,7 @@ typedef struct
#define WRITE_NV21 "write_nv21" #define WRITE_NV21 "write_nv21"
#define WRITE_P010 "write_p010" #define WRITE_P010 "write_p010"
#define WRITE_I420_10 "write_i420_10" #define WRITE_I420_10 "write_i420_10"
#define WRITE_I420_12 "write_i420_12"
#define WRITE_Y444 "write_y444" #define WRITE_Y444 "write_y444"
#define WRITE_Y444_10 "write_y444_10" #define WRITE_Y444_10 "write_y444_10"
#define WRITE_Y444_12 "write_y444_12" #define WRITE_Y444_12 "write_y444_12"
@ -1003,6 +1004,18 @@ WRITE_I420_10 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2
"}\n" "}\n"
"\n" "\n"
"__device__ inline void\n" "__device__ inline void\n"
WRITE_I420_12 "(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"
" *(unsigned short *) &dst0[x * 2 + y * stride0] = scale_to_12bits (sample.x);\n"
" if (x % 2 == 0 && y % 2 == 0) {\n"
" unsigned int pos = x + (y / 2) * stride1;\n"
" *(unsigned short *) &dst1[pos] = scale_to_12bits (sample.y);\n"
" *(unsigned short *) &dst2[pos] = scale_to_12bits (sample.z);\n"
" }\n"
"}\n"
"\n"
"__device__ inline void\n"
WRITE_Y444 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n" WRITE_Y444 "(unsigned char * dst0, unsigned char * dst1, unsigned char * dst2,\n"
" unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\n"
"{\n" "{\n"
@ -1485,6 +1498,7 @@ static const TextureFormat format_map[] = {
MAKE_FORMAT_YUV_SEMI_PLANAR (P012_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR), MAKE_FORMAT_YUV_SEMI_PLANAR (P012_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR),
MAKE_FORMAT_YUV_SEMI_PLANAR (P016_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR), MAKE_FORMAT_YUV_SEMI_PLANAR (P016_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR),
MAKE_FORMAT_YUV_PLANAR (I420_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS), MAKE_FORMAT_YUV_PLANAR (I420_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS),
MAKE_FORMAT_YUV_PLANAR (I420_12LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_12BIS),
MAKE_FORMAT_YUV_PLANAR (Y444, UNSIGNED_INT8, SAMPLE_YUV_PLANAR), MAKE_FORMAT_YUV_PLANAR (Y444, UNSIGNED_INT8, SAMPLE_YUV_PLANAR),
MAKE_FORMAT_YUV_PLANAR (Y444_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS), MAKE_FORMAT_YUV_PLANAR (Y444_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS),
MAKE_FORMAT_YUV_PLANAR (Y444_12LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_12BIS), MAKE_FORMAT_YUV_PLANAR (Y444_12LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_12BIS),
@ -1730,6 +1744,9 @@ gst_cuda_converter_setup (GstCudaConverter * self)
case GST_VIDEO_FORMAT_I420_10LE: case GST_VIDEO_FORMAT_I420_10LE:
write_func = WRITE_I420_10; write_func = WRITE_I420_10;
break; break;
case GST_VIDEO_FORMAT_I420_12LE:
write_func = WRITE_I420_12;
break;
case GST_VIDEO_FORMAT_Y444: case GST_VIDEO_FORMAT_Y444:
write_func = WRITE_Y444; write_func = WRITE_Y444;
break; break;

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

@ -30,7 +30,7 @@ GST_DEBUG_CATEGORY_STATIC (gst_cuda_base_convert_debug);
#define GST_CAT_DEFAULT gst_cuda_base_convert_debug #define GST_CAT_DEFAULT gst_cuda_base_convert_debug
#define GST_CUDA_CONVET_FORMATS \ #define GST_CUDA_CONVET_FORMATS \
"{ I420, YV12, NV12, NV21, P010_10LE, P012_LE, P016_LE, I420_10LE, Y444, " \ "{ 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, " \ "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, " \ "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 }"

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

@ -24,7 +24,7 @@
G_BEGIN_DECLS G_BEGIN_DECLS
#define GST_CUDA_FORMATS \ #define GST_CUDA_FORMATS \
"{ I420, YV12, NV12, NV21, P010_10LE, P012_LE, P016_LE, I420_10LE, Y444, " \ "{ 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, " \ "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, " \ "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 }"
@ -35,7 +35,7 @@ G_BEGIN_DECLS
"YUY2, UYVY, RGBP, BGRP, GBR, GBRA }" "YUY2, UYVY, RGBP, BGRP, GBR, GBRA }"
#define GST_CUDA_D3D11_FORMATS \ #define GST_CUDA_D3D11_FORMATS \
"{ I420, YV12, I420_10LE, Y444, Y444_10LE, Y444_12LE, Y444_16LE, " \ "{ I420, YV12, I420_10LE, I420_12LE, Y444, Y444_10LE, Y444_12LE, Y444_16LE, " \
"BGRA, RGBA, BGRx, RGBx, Y42B, I422_10LE, I422_12LE, GBR, GBR, GBR_10LE, " \ "BGRA, RGBA, BGRx, RGBx, Y42B, I422_10LE, I422_12LE, GBR, GBR, GBR_10LE, " \
"GBR_12LE, GBR_16LE }" "GBR_12LE, GBR_16LE }"