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 7096b7935f..d108f82b14 100644 --- a/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp +++ b/subprojects/gst-plugins-bad/gst-libs/gst/cuda/gstcudamemory.cpp @@ -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; diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c index ca336c0af6..5f8bf5781e 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconverter.c @@ -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(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; } diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c index 272b074f6b..a374a50f69 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaconvertscale.c @@ -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, diff --git a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h index 252a1ab979..875fb47656 100644 --- a/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h +++ b/subprojects/gst-plugins-bad/sys/nvcodec/gstcudaformat.h @@ -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, " \