2

以下に示すこの単純なコードは、ストリームを使用してホストからデバイスにデータをコピーするだけで、何もしていません。しかし、cudamemcpyasync が実際には非同期であり、ストリームを理解しているため、nvprof を実行した後、私は混乱しています。

#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]

そのため、開始時間のために、ここでストリームを使用する意味がわかりませんでした。私にはシーケンシャルに見えます。ここで私が間違っていることを理解するのを手伝ってください。tesla K20c カードを使用しています。

4

1 に答える 1

3

GPU をシステムに接続する PCI Express リンクには、カードへのチャネルが 1 つと、カードからのチャネルが 1 つしかありません。つまり、最大で、方向ごとに、任意の時点で実際に実行される 1 つの cudaMemcpy(Async) 操作 (つまり、最大で 1 つの DtoH と 1 つの HtoD) を持つことができます。他のすべての cudaMemcpy(Async) 操作はキューに入れられ、先行する操作が完了するのを待ちます。

2 つの操作を同時に同じ方向に進めることはできません。方向ごとに 1 つずつ。

@JackOLantern が述べているように、ストリームの主な利点は、メモリコピーと計算をオーバーラップさせること、または複数のカーネルを同時に実行できるようにすることです。また、1 つの DtoH コピーを 1 つの HtoD コピーと同時に実行することもできます。

プログラムはすべての HtoD コピーを実行するため、それらはすべてシリアルに実行されます。各コピーは、その前のコピーが完了するまで待機する必要があります。

HtoD と DtoH の memcopy を同時に実行するには、複数のコピー エンジンを備えたデバイスが必要です。deviceQuery を使用して、デバイスに関するこれを検出できます。

また、同時動作を有効にするには、ホスト側のバッファーにcudaHostAllocではなくを使用する必要があることにも注意してください。malloc

編集:上記の回答には、最大 2 つのコピー エンジン (方向ごとに 1 つ) を持つ GPU が表示されており、そのような GPU に対しては依然として正しいです。ただし、いくつかの新しい Pascal および Volta ファミリー メンバーの GPU には、2 つ以上のコピー エンジンが搭載されています。その場合、方向ごとに 2 つ (またはそれ以上) のコピー エンジンを使用すると、理論的には、その方向で「飛行中」の 2 つ (またはそれ以上) の転送を行うことができます。ただし、これは PCIE (または NVLink) バス自体の特性を変更しません。利用可能な帯域幅にはまだ制限があり、正確な低レベルの動作 (そのような転送が「シリアル化」されているように見えるか、同時に実行されているように見えますが、帯域幅の共有のために時間がかかるかどうか) は、ほとんどの場合あまり重要ではありません。

于 2013-09-19T12:19:14.587 に答える