cudamemcpyasync 和流行为理解

我有下面所示的简单代码，它什么也不做，只是使用流将一些数据从主机复制到设备。但在运行 nvprof 后我很困惑，因为 cudamemcpyasync 确实是异步的并且对流的理解。

#include <stdio.h>

#define NUM_STREAMS 4
cudaError_t memcpyUsingStreams (float           *fDest,
                                float           *fSrc,
                                int             iBytes,
                                cudaMemcpyKind  eDirection,
                                cudaStream_t    *pCuStream)
{
    int             iIndex = 0 ;
    cudaError_t     cuError = cudaSuccess ;
    int             iOffset = 0 ;

    iOffset = (iBytes / NUM_STREAMS) ;
    /*Creating streams if not present */
    if (NULL == pCuStream)
    {
            pCuStream = (cudaStream_t *) malloc(NUM_STREAMS * sizeof(cudaStream_t));
            for (iIndex = 0 ; iIndex < NUM_STREAMS; iIndex++)
            {
                    cuError = cudaStreamCreate (&pCuStream[iIndex]) ;
            }
    }

    if (cuError != cudaSuccess)
    {
            cuError = cudaMemcpy (fDest, fSrc, iBytes, eDirection) ;
    }
    else
    {
            for (iIndex = 0 ; iIndex < NUM_STREAMS; iIndex++)
            {
                    iOffset = iIndex * iOffset ;
                    cuError = cudaMemcpyAsync (fDest +  iOffset , fSrc + iOffset, iBytes / NUM_STREAMS , eDirection, pCuStream[iIndex]) ;
            }
    }

    if (NULL != pCuStream)
    {
            for (iIndex = 0 ; iIndex < NUM_STREAMS; iIndex++)
            {
                    cuError = cudaStreamDestroy (pCuStream[iIndex]) ;
            }
            free (pCuStream) ;
    }
    return cuError ;
}


int main()
{
    float *hdata = NULL ;
    float *ddata = NULL ;
    int i, j, k, index ;
    cudaStream_t *abc = NULL ;

    hdata = (float *) malloc (sizeof (float) * 256 * 256 * 256) ;

    cudaMalloc ((void **) &ddata, sizeof (float) * 256 * 256 * 256) ;

    for (i=0 ; i< 256 ; i++)
    {
        for (j=0; j< 256; j++)
        {
            for (k=0; k< 256 ; k++)
            {
                index = (((i * 256) + j) * 256) + k;
                hdata [index] = index ;
            }
        }
    }

    memcpyUsingStreams (ddata, hdata, sizeof (float) * 256 * 256 * 256,  cudaMemcpyHostToDevice, abc) ;

    cudaFree (ddata) ;
    free (hdata) ;

    return 0;
}

nvprof 结果如下。

    Start  Duration           Grid Size     Block Size     Regs*    SSMem*    DSMem*      Size  Throughput    Device   Context    Stream  Name
 104.35ms   10.38ms                   -              -         -         -         -   16.78MB    1.62GB/s         0         1         7  [CUDA memcpy HtoD]
 114.73ms   10.41ms                   -              -         -         -         -   16.78MB    1.61GB/s         0         1         8  [CUDA memcpy HtoD]
 125.14ms   10.46ms                   -              -         -         -         -   16.78MB    1.60GB/s         0         1         9  [CUDA memcpy HtoD]
 135.61ms   10.39ms                   -              -         -         -         -   16.78MB    1.61GB/s         0         1        10  [CUDA memcpy HtoD]

所以由于开始时间的原因，我不明白在这里使用流的意义。对我来说它看起来是连续的。请帮助我理解我在这里做错了什么。我正在使用特斯拉 K20c 卡。

将 GPU 连接到系统的 PCI Express 链路只有一个通往卡的通道和一个来自卡的通道。这意味着最多可以有一个 cudaMemcpy(Async) 操作，该操作实际上在每个方向的任何给定时间执行（即最多一个 DtoH 和一个 HtoD）。所有其他 cudaMemcpy(Async) 操作都将排队，等待前面的操作完成。

您不能同时进行同一方向的两项操作。每个方向一次一个。

正如 @JackOLantern 所说，流的主要好处是重叠内存副本和计算，或者允许多个内核同时执行。它也是允许一个 DtoH 副本与一个 HtoD 副本同时运行 http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#concurrent-data-transfers.

由于您的程序执行所有 HtoD 副本，因此它们都会串行执行。每个副本都必须等待其前面的副本完成。

即使同时执行 HtoD 和 DtoH memcopy 也需要具有多个复制引擎的设备；您可以使用 deviceQuery 发现有关您设备的信息。

我还应该指出，要启用并发行为，您应该使用cudaHostAlloc, not malloc，用于主机端缓冲区。

EDIT:上面的答案考虑的是最多有 2 个复制引擎（每个方向一个）的 GPU，并且对于此类 GPU 来说仍然是正确的。然而，存在一些较新的 Pascal 和 Volta 系列成员 GPU，它们具有 2 个以上的复制引擎。在这种情况下，每个方向有 2 个（或更多）复制引擎，理论上可以在该方向上有 2 个（或更多）“飞行中”传输。然而，这不会改变 PCIE（或 NVLink）总线本身的特性。您仍然受到可用带宽的限制，并且在大多数情况下，确切的低级行为（无论此类传输是否显示为“序列化”，或者是否显示为同时运行，但由于共享带宽而需要更长的时间）应该无关紧要。