cuda - 为什么NVENC示例同时使用cuMemcpyHtoD和cuMemcpy2D复制YUV数据？

Question

我正在学习NVIDIA NVENCAPI。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吗？也可以使用cuMemcpyHtoD复制吗？据我了解，YUV平面具有相同的线性内存布局。唯一的区别是它们的大小。

PS:我最初在 Computer Graphics网站上问了这个问题，但没有得到答案。

Answer 1

在主机上，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在该文件中分配了)

“半平面”存储。主机上的(无间距)平面存储意味着Y，然后是U，然后是V。设备上的“半平面”存储意味着Y平面，然后是具有U和V交错的特殊平面:

U0V0  U1V1  U2V2 ...

因此，只需一次2D memcpy调用即可轻松复制Y数据。但是，UV平面需要从单独的缓冲区进行一些组装。该代码的编写者选择按以下方式进行汇编:

将U平面和V平面分别复制到设备，并复制到独立的未插入缓冲区中。那就是您所显示的代码，并且设备上的独立缓冲区分别是m_ChromaDevPtr[0]和m_ChromaDevPtr[1](U然后V，分开，不加音调)。

使用CUDA内核在设备上组装UV交错的倾斜平面。这是有道理的，因为存在大量数据移动，并且具有更高内存带宽的设备可以比主机上更有效地执行此操作。还要注意，单个2D memcpy调用无法处理这种情况，因为我们实际上有2个大步。一个是从一个元素到另一个元素的(短)步幅，例如，在上面的示例中，从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”内核的一些参数是:

指向设备上单独的U和V缓冲区的指针(例如&m_ChromaDevPtr[0]等)

指向主缓冲区中UV交错平面将为(&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数据的复制，因为目标数据已发送。

U和V数据的副本是从未分配缓冲区到未分配缓冲区的，因此可以使用cuMemcpyHtoD。

U和V数据的主机到设备副本无法直接进入NV12缓冲区，因为没有cuMemcpy操作(2D或其他方式)可以处理该特定目标存储模式。