diff --git a/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp b/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp index 62f6509d03..7bce1bdc49 100644 --- a/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp +++ b/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp @@ -195,6 +195,7 @@ gst_cuda_allocator_update_info (const GstVideoInfo * reference, case GST_VIDEO_FORMAT_I420: case GST_VIDEO_FORMAT_YV12: case GST_VIDEO_FORMAT_I420_10LE: + case GST_VIDEO_FORMAT_I420_12LE: /* we are wasting space yes, but required so that this memory * can be used in kernel function */ 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 (P016_LE, 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_10LE, 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_P016_LE: case GST_VIDEO_FORMAT_I420_10LE: + case GST_VIDEO_FORMAT_I420_12LE: alloc_height = GST_ROUND_UP_2 (alloc_height); break; default: @@ -1038,6 +1041,7 @@ gst_cuda_allocator_calculate_alloc_height (const GstVideoInfo * info) case GST_VIDEO_FORMAT_I420: case GST_VIDEO_FORMAT_YV12: case GST_VIDEO_FORMAT_I420_10LE: + case GST_VIDEO_FORMAT_I420_12LE: alloc_height *= 2; break; case GST_VIDEO_FORMAT_NV12: diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c index d0c4f27ed8..a6661fd692 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c @@ -645,6 +645,7 @@ typedef struct #define WRITE_NV21 "write_nv21" #define WRITE_P010 "write_p010" #define WRITE_I420_10 "write_i420_10" +#define WRITE_I420_12 "write_i420_12" #define WRITE_Y444 "write_y444" #define WRITE_Y444_10 "write_y444_10" #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" "__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" " unsigned char * dst3, float4 sample, int x, int y, int stride0, int stride1)\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 (P016_LE, UNSIGNED_INT16, SAMPLE_SEMI_PLANAR), 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_10LE, UNSIGNED_INT16, SAMPLE_YUV_PLANAR_10BIS), 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: write_func = WRITE_I420_10; break; + case GST_VIDEO_FORMAT_I420_12LE: + write_func = WRITE_I420_12; + break; case GST_VIDEO_FORMAT_Y444: write_func = WRITE_Y444; break; diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c index 2219be2ef9..f402afb521 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c @@ -30,7 +30,7 @@ GST_DEBUG_CATEGORY_STATIC (gst_cuda_base_convert_debug); #define GST_CAT_DEFAULT gst_cuda_base_convert_debug #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, " \ "BGR, BGR10A2_LE, RGB10A2_LE, Y42B, I422_10LE, I422_12LE, RGBP, BGRP, GBR, " \ "GBRA, GBR_10LE, GBR_12LE, GBR_16LE }" diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h index 233828329c..252a1ab979 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h @@ -24,7 +24,7 @@ G_BEGIN_DECLS #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, " \ "BGR, BGR10A2_LE, RGB10A2_LE, Y42B, I422_10LE, I422_12LE, YUY2, UYVY, RGBP, " \ "BGRP, GBR, GBR_10LE, GBR_12LE, GBR_16LE, GBRA }" @@ -35,7 +35,7 @@ G_BEGIN_DECLS "YUY2, UYVY, RGBP, BGRP, GBR, GBRA }" #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, " \ "GBR_12LE, GBR_16LE }"