为什么NVENC样本同时使用cuMemcpyHtoD和cuMemcpy2D来复制YUV数据？

Question

我正在学习NVIDIA NVENC API.SDK提供了一个名为"NvEncoderCudaInterop"的样本.有一大块代码可以将YUV平面阵列从CPU复制到GPU缓冲区.这是代码:

 // copy luma
 CUDA_MEMCPY2D copyParam;
memset(&copyParam, 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(&copyParam));

// 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？为什么Y可以'也可以使用cuMemcpyHtoD复制？据我所知,YUV平面具有相同的线性内存布局.唯一的区别是它们的大小.

PS:我最初在计算机图形网站上问过这个问题,但没有得到答案.

Answer 1

在主机上,YUV缓冲器数据(假设)存储为存储在不同平面中的无节点YUV 4:2:0数据.这意味着Y拥有它自己的plane(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缓冲区.

请注意,这是为了:

1. 倾斜存储(这很明显,因为主缓冲区指针pEncodeBuffer->stInputBfr.pNV12devPtr先前已在该文件中分配cuMemAllocPitch)

2. "半平面"存储.主机上的(未切换的)平面存储器意味着Y后跟U,后跟V.设备上的"半平面"存储器意味着Y平面,后面是U和V交错的特殊平面:

   U0V0  U1V1  U2V2 ...
   

因此,使用单个2D memcpy调用复制Y数据非常容易.但UV平面需要从单独的缓冲区进行一些组装.这段代码的编写者选择按如下方式进行汇编:

1. 将U和V平面独立地复制到设备,连接到独立的无节点缓冲区.这是已显示的代码,并在设备上的独立的缓冲器是m_ChromaDevPtr[0]和m_ChromaDevPtr[1]分别(U然后V,单独的,未俯仰).

2. 使用CUDA内核在器件上组装 UV交错的倾斜平面.这是有道理的,因为存在相当数量的数据移动,并且具有更高存储器带宽的设备可以比在主机上更有效地执行此操作.另请注意,单个2D memcpy调用无法处理此情况,因为我们实际上有两个步幅.一个是从元素到元素的(短)步幅,例如在上面的例子中从U0到U1的短步幅.另一个步幅是每条线末端的"更长"步幅,与投球分配相关的"正常"步幅.

完成从非交错,无需切换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"内核的一些参数是:

1. 指向设备上单独的U和V缓冲区的指针(例如,&m_ChromaDevPtr[0]等等)
2. 指向UV交错平面所在主缓冲区中起始位置的指针(&dNV12Chroma)
3. 指向目标缓冲区间距的指针(&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-only和V-only缓冲区中获取数据,并根据上述模式存储该数据.它需要将U和V缓冲区组装在一起,然后将该数据存放在斜距目标缓冲区中.首先通过将U和V数据复制到单独的设备缓冲区来处理,这些缓冲区的组织方式与主机上的完全相同:

U0  U1  U2  U3
U4  U5  U6  U7
....

这种类型的副本是用普通的,无需处理的 cuMemcpyHtoD

这是一个操作图:

笔记:

1. Y数据的副本不能用普通数据完成cuMemcpyHtoD,因为目标数据是倾斜的.
2. U和V数据的副本是从无节点缓冲区到无节点缓冲区,因此可以使用cuMemcpyHtoD.
3. U和V数据的主机到设备副本无法直接进入NV12缓冲区,因为没有可以处理特定目标存储模式的cuMemcpy操作(2D或其他).