我正在学习NVIDIA NVENCAPI。SDK提供了一个名为“NvEncoderCudaInterop”的示例。其中有大量代码将YUV平面数组从CPU复制到GPU缓冲区。
这是代码:
// copy luma
CUDA_MEMCPY2D copyParam;
memset(©Param, 0, sizeof(copyParam));
copyParam.dstMemoryType = CU_MEMORYTYPE_DEVICE;
copyParam.dstDevice = pEncodeBuffer->stInputBfr.pNV12devPtr;
copyParam.dstPitch = pEncodeBuffer->stInputBfr.uNV12Stride;
copyParam.srcMemoryType = CU_MEMORYTYPE_HOST;
copyParam.srcHost = yuv[0];
copyParam.srcPitch = width;
copyParam.WidthInBytes = width;
copyParam.Height = height;
__cu(cuMemcpy2D(©Param));
// copy chroma
__cu(cuMemcpyHtoD(m_ChromaDevPtr[0], yuv[1], width*height / 4));
__cu(cuMemcpyHtoD(m_ChromaDevPtr[1], yuv[2], width*height / 4));
我确实了解该过程的原理。将内存复制到GPU以供内核处理。我不明白的原因是,为了复制Y平面,使用cuMemcpy2D并将其用于UV cuMemcpyHtoD吗?也可以使用cuMemcpyHtoD复制吗?据我了解,YUV平面具有相同的线性内存布局。唯一的区别是它们的大小。
PS:我最初在Computer Graphics网站上问了这个问题,但没有得到答案。
最佳答案
在主机上,YUV缓冲区数据(假定为)存储为未分隔的YUV 4:2:0数据,存储在单独的平面中。这意味着Y有它自己的平面(yuv[0]
),然后是U(yuv[1]
),然后是V(yuv[2]
)。
设备上的预期存储目标是(NV12)缓冲区格式,定义为NV_ENC_BUFFER_FORMAT_NV12_PL
,文档(NvEncodeAPI_v.5.0.pdf,第12页)定义为:
NV_ENC_BUFFER_FORMAT_NV12_PL半平面YUV [UV交错]分配为串行2D缓冲区。
请注意,这是为了:
pEncodeBuffer->stInputBfr.pNV12devPtr
先前已使用cuMemAllocPitch
在该文件中分配了)U0V0 U1V1 U2V2 ...
因此,只需一次2D memcpy调用即可轻松复制Y数据。但是,UV平面需要从单独的缓冲区进行一些组装。该代码的编写者选择按以下方式进行汇编:
m_ChromaDevPtr[0]
和m_ChromaDevPtr[1]
(U然后V,分开,不加音调)。 从非交错的,未插入间距的
m_ChromaDevPtr[0]
和m_ChromaDevPtr[1]
缓冲区在设备上完成UV交错的,倾斜的平面的“组装”的内核称为m_cuInterleaveUVFunction
,它在此处启动(紧随显示的代码之后,从您显示的代码的结尾): __cu(cuMemcpyHtoD(m_ChromaDevPtr[0], yuv[1], width*height / 4));
__cu(cuMemcpyHtoD(m_ChromaDevPtr[1], yuv[2], width*height / 4));
#define BLOCK_X 32
#define BLOCK_Y 16
int chromaHeight = height / 2;
int chromaWidth = width / 2;
dim3 block(BLOCK_X, BLOCK_Y, 1);
dim3 grid((chromaWidth + BLOCK_X - 1) / BLOCK_X, (chromaHeight + BLOCK_Y - 1) / BLOCK_Y, 1);
#undef BLOCK_Y
#undef BLOCK_X
CUdeviceptr dNV12Chroma = (CUdeviceptr)((unsigned char*)pEncodeBuffer->stInputBfr.pNV12devPtr + pEncodeBuffer->stInputBfr.uNV12Stride*height);
void *args[8] = { &m_ChromaDevPtr[0], &m_ChromaDevPtr[1], &dNV12Chroma, &chromaWidth, &chromaHeight, &chromaWidth, &chromaWidth, &pEncodeBuffer->stInputBfr.uNV12Stride};
__cu(cuLaunchKernel(m_cuInterleaveUVFunction, grid.x, grid.y, grid.z,
block.x, block.y, block.z,
0,
NULL, args, NULL));
CUresult cuResult = cuStreamQuery(NULL);
if (!((cuResult == CUDA_SUCCESS) || (cuResult == CUDA_ERROR_NOT_READY)))
{
return NV_ENC_ERR_GENERIC;
}
return NV_ENC_SUCCESS;
}
请注意,传递给此“UV Assembly”内核的一些参数是:
&m_ChromaDevPtr[0]
等)&dNV12Chroma
)&pEncodeBuffer->stInputBfr.uNV12Stride
)就像您将要编写自己的内核来进行该汇编一样。如果要查看程序集内核中的实际内容,则在该示例项目的preproc.cu文件中。
编辑:
在评论中回答问题。在主机上,Y数据是这样存储的(假设每行仅包含4个元素。这对于YUV 4:2:0数据并不真正正确,但是这里的重点是复制操作,而不是行长) ):
Y0 Y1 Y2 Y3
Y4 Y5 Y6 Y7
....
在设备上,该缓冲区的组织方式如下:
Y0 Y1 Y2 Y3 X X X X
Y4 Y5 Y6 Y7 X X X X
...
其中
X
值被填充以使每行等于音高。要从上方的主机缓冲区复制到上方的设备缓冲区,我们必须使用带间距的副本,即cuMemcpy2D
。在主机上,U数据的组织方式如下:
U0 U1 U2 U3
U4 U5 U6 U7
....
V数据的组织方式类似:
V0 V1 V2 V3
V4 V5 V6 V7
....
在设备上,以上的U和V数据最终将组合到一个单独的UV平面中,该平面也将像这样倾斜:
U0V0 U1V1 U2V2 U3V3 X X X X
U4V4 U5V5 U6V6 U7V7 X X X X
...
没有一个单独的memcpy操作可以正确地从未分配主机的仅U和V缓冲区中获取数据,并根据上述模式存储该数据。它需要将U和V缓冲区组装在一起,然后将数据存放在倾斜的目标缓冲区中。首先,通过将U和V数据复制到单独的设备缓冲区中来进行处理,这些缓冲区的组织方式与在主机上完全相同:
U0 U1 U2 U3
U4 U5 U6 U7
....
这种类型的副本使用普通的,无节距的
cuMemcpyHtoD
处理这是操作图:
笔记:
cuMemcpyHtoD
进行Y数据的复制,因为目标数据已发送。 cuMemcpyHtoD
。 关于cuda - 为什么NVENC示例同时使用cuMemcpyHtoD和cuMemcpy2D复制YUV数据?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/33012998/