2

我有下面显示的这个简单代码，它什么都不做，只是使用流将一些数据从主机复制到设备。但是在运行 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 卡。

4

1 回答 1

3

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

您不能同时进行两个操作在同一方向上。一次一个，每个方向。

正如@JackOLantern 所说，流的主要好处是重叠内存副本和计算，或者允许多个内核同时执行。它还允许一个 DtoH 副本与一个 HtoD 副本同时运行。

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

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

我还应该指出，要启用并发行为，您应该使用cudaHostAlloc，而不是malloc，作为您的主机端缓冲区。

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

于 2013-09-19T12:19:14.587 回答