使用ffmpeg、nvdia解码rtsp视频流，cuda做NV12-RGBA转换

    本章是在nvidia_video_sdk_6.0.1的基础之上做封装的，我研究了其中的NvDecodeGL工程；由于自己工作会遇到显示多路rtsp视频流及解码的情况，所以进行了研究。

    网上有其它的介绍ffmpeg和nvdia结合解码视频的文章，这里我将其实现了，并将官方的代码进行了精简和封装，封装后使用方法相当简单，示例如下

#include "NvDecode.h"
#include "opencv.hpp"
#include <iostream>

int main()
{
	NvDecode decod;//rtsp://admin:123ABCabc@192.168.2.236:554/Streaming/Channels/102?transportmode=unicast&profile=Profile_1
	//rtsp://184.72.239.149/vod/mp4://BigBuckBunny_175k.mov
	decod.start(std::string("rtsp://184.72.239.149/vod/mp4://BigBuckBunny_175k.mov"));
	unsigned char *rgbaPtr = nullptr;
	int width = 0, height = 0;
	unsigned long long timestamp = 0;
	while (!decod.m_pFrameQueue->isEndOfDecode()) //到了视频末尾退出循环
	{
		if (decod.deQueueFrame(&rgbaPtr, &width, &height, timestamp)) {
			cv::Mat frame(height, width, CV_8UC4, rgbaPtr);
			cv::imshow("video", frame);
			cv::waitKey(30);
		}
		else {
			cv::waitKey(20); //如果队列里面没有视频帧就等待一下
			continue;
		}
	}
	return 0;
}

     我是使用的opencv进行显示的，出来的结果已经是rgba了，具体是bgra还是rgba还有点懵，opencv能播放，应该是bgra吧。但是官方的cuda函数写的是rgba。先不管这个了。

    其中进行了大量的调试，官方的工程功能众多，主要说下流程吧

一、初始化设备、获取cuda的格式转换函数

	__cu(cuInit(0, __CUDA_API_VERSION, hHandleDriver));
	__cu(cuvidInit(0));

	__cu(cuDeviceGet(&device, 0)); //使用0号显卡
	__cu(cuCtxCreate(&cudaCtx, CU_CTX_SCHED_AUTO, device));
	__cu(cuvidCtxLockCreate(&ctxLock, cudaCtx));
	m_pFrameQueue = new CUVIDFrameQueue(ctxLock);

	CUresult oResult;
	//..\\3rd\\common\\kernels\\ptx\\NV12ToARGB_drvapi_Win32.ptx
	oResult = cuModuleLoad(&module, "..\\3rd\\common\\kernels\\ptx\\NV12ToARGB_drvapi_Win32.ptx");
	if (oResult != CUDA_SUCCESS) {
		std::cout << "load module failed error: " << oResult << std::endl;
		exit(-1);
	}
	oResult = cuModuleGetFunction(&g_kernelNV12toARGB, module, "NV12ToARGB_drvapi");
	if (oResult != CUDA_SUCCESS) {
		std::cout << "get cuda func NV12ToARGB_drvapi failed" << std::endl;
		exit(-1);
	}

    二、ffmpeg的初始化，并将视频信息和nvdia创建视频信息的数据进行匹配，这个网上多，主不说了

    三、在解码线程中循环的解码，将读取到的每一帧构造一个CUVIDSOURCEDATAPACKET并调用cuvidParseVideoData函数进行解码。

    四、将解码的结果使用cuda函数转换为rgba格式。我使用的是官方的FrameQueue做为解码和显示的桥梁，当显示线程暂停到某一帧时，队列为满，解码线程将会等待。官方代码如下

bool
FrameQueue::waitUntilFrameAvailable(int nPictureIndex)
{
    while (isInUse(nPictureIndex))
    {
        Sleep(1);   // Decoder is getting too far ahead from display
        if (isEndOfDecode())
            return false;
    }

    return true;
}

当显示线程取到一帧时，我是这样做转换的，此外研究了很久，删减了官方的大量封装，我只是要将图片转换为rgba格式，在GPU中完成，可以更大减少CPU的占用。

bool NvDecode::deQueueFrame(unsigned char ** ptr, int *width, int *height, unsigned long long *timestamp)
{
	CUVIDPARSERDISPINFO pInfo;
	if (!(m_pFrameQueue->isEndOfDecode() && m_pFrameQueue->isEmpty())) {
		if (m_pFrameQueue->dequeue(&pInfo)) {
			CCtxAutoLock lck(ctxLock);
			cuCtxPushCurrent(cudaCtx);
			CUdeviceptr pDecodedFrame[2] = { 0,0 };
			CUdeviceptr pInteropFrame[2] = { 0,0 };
			int distinct_fields = 1;
			if (!pInfo.progressive_frame && pInfo.repeat_first_field <= 1) {
				distinct_fields = 2;
			}

			for (int active_field = 0; active_field < distinct_fields; active_field++) {
				CUVIDPROCPARAMS oVPP = { 0 };
				oVPP.progressive_frame = pInfo.progressive_frame;
				oVPP.top_field_first = pInfo.top_field_first;
				oVPP.unpaired_field = (distinct_fields == 1);
				oVPP.second_field = active_field;
				unsigned int nDecodedPitch = 0; //将解码后的原始帧映射出来，nDecodePitch表示原来帧空间的每行所点字节，不一定是视频宽度，不知道的要去了解下cuda矩阵内存分配了
				if (cuvidMapVideoFrame(m_videoDecoder, pInfo.picture_index, &pDecodedFrame[active_field], &nDecodedPitch, &oVPP) != CUDA_SUCCESS) {
					m_pFrameQueue->releaseFrame(&pInfo);
					cuCtxPopCurrent(NULL);
					return false;
				}

				if (isFirstFrame) { //如果是第一帧需要初始化gpu上下文中的全局变量，实际是alapha透明度
					*ptr = rgbaBuf;
					*width = targetWidth;
					*height = targetHeight;
					float hueColorSpaceMat[9];
					setColorSpaceMatrix(ITU601, hueColorSpaceMat, 0.0f);
					updateConstantMemory_drvapi(module, hueColorSpaceMat);
					isFirstFrame = false;
				}

				pInteropFrame[active_field] = g_pInteropFrame; //设置转换为rgba格式的内存地址
				int dstPictch = targetWidth * 4;  //设置rgba格式的行宽，一个像素占4个字节，所以是targetWidth*4

				dim3 block(32, 16, 1); //blck和grid是从官方代码调试时获取的，我只要能解视频所以就直接拿底层的数据了
				dim3 grid((targetWidth + (2 * block.x - 1)) / (2 * block.x), (targetHeight + (block.y - 1)) / block.y, 1);
				void *args[] = { &pDecodedFrame[active_field], &nDecodedPitch, //传入的参数
					&pInteropFrame[active_field], &dstPictch,
					&targetWidth, &targetHeight
				};
			   CUresult oRes = cuLaunchKernel(g_kernelNV12toARGB, grid.x, grid.y, grid.z,//这里调用cuda的函数完成转换
					block.x, block.y, block.z,
					0, 0, args, NULL);
			   if (oRes != CUDA_SUCCESS) {
				   std::cout << "launchKernel failed,status" << oRes << std::endl;
				   return false;
			   }
                                //这里将转换的结果从显存拷贝到内存，这个技术点研究了很久。。。
				checkCudaErrors(cuMemcpyDtoH(rgbaBuf, pInteropFrame[active_field], dstPictch * targetHeight));
				cuvidUnmapVideoFrame(m_videoDecoder, pDecodedFrame[active_field]);
			}

			cuCtxPopCurrent(NULL);
			*timestamp = pInfo.timestamp;
			m_pFrameQueue->releaseFrame(&pInfo);
			return true;
		}
	}
	return false;
}

    将解码的结果进行格式转换，并取出来的核心代码就是上面了。经自己电脑验证vlc推送的udp形式的rtsp流不能调用cuvidDecodePicture进行解码，但公网上的可以，我试了一个rtsp://184.72.239.149/vod/mp4://BigBuckBunny_175k.mov；如果是视频文件的话也可以，用这个硬解码速度很快的，第二个公网上的偶而有些卡，估计是推送的网络速度吧。

    这次完成了用nvdia显示硬解码，并且在GPU中从nv12转换到rgba，上层应用就可以直接显示使用了，极大的减轻了cpu的负担；udp不能解的原因暂时还不知道。

整个工程缩略图如下

    需要工程代码的可以下载，需要积分罗，我搞了很久，赏个脸。到此下载