1

複数のストリームを使用して、CUDAコードの次のビットを高速化しようとしています。

#define N (4096 * 4096)
#define blockDimX  16
#define blockDimY  16

float domain1 [N];
float domain2 [N];

__global__ updateDomain1_kernel(const int dimX, const int dimY) {
    // update mechanism here for domain1
    // ...
}

__global__ updateDomain2_kernel(const int dimX, const int dimY) {
    // update mechanism here for domain2, which is nearly the same
    // ...
}

__global__ addDomainsTogether_kernel(float* domainOut, 
                                     const int dimX, 
                                     const int dimY) 
{
    // add domain1 and domain2 together and fill domainOut
}

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    updateDomain1_kernel<<<blocks, threads>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads>>> (dimX, dimY);
    addDomainsTogether_kernel<<<block, threads>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);
}

正確な実装は実際には重要ではありません。重要なのは、それぞれのドメインの更新が2つの完全に独立した操作であり、その後、両方が3番目のカーネル呼び出しで使用されることです。したがって、同時に実行したい各更新カーネルを独自のストリームに入れて、高速化することをお勧めします。だから私はそれを次のように変更しました：

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    cudaStream_t stream0, stream1;
    cudaStreamCreate(&stream0);
    cudaStreamCreate(&stream1);

    updateDomain1_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads, 0, stream1>>> (dimX, dimY);
    cudaDeviceSynchronize();

    addDomainsTogether_kernel<<<block, threads>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaStreamDestroy(stream0);
    cudaStreamDestroy(stream1);
}

パフォーマンス速度に違いがあると思いましたが、目立った違いは全くありません。したがって、コンパイラーが更新呼び出しを同時に自動的にスケジュールすることによって初めて賢くなったのではないかと考えて、次のようにパフォーマンスが低下するはずだと思いました。

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    cudaStream_t stream0;
    cudaStreamCreate(&stream0);

    updateDomain1_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);

    addDomainsTogether_kernel<<<block, threads0, stream0>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaStreamDestroy(stream0);
}

ただし、ここでもパフォーマンス速度にほとんど違いはありません。どちらかといえば、最後のものが最速のようです。それは私が理解していないCUDAストリームについて何かがあると私に思わせます。誰かがこのコードをどのように加速するかについて私に教えてもらえますか？

4

1 に答える 1

1

並列処理を増やすと、使用可能なすべてのコアをまだ使用していない場合にのみ、計算スループットが向上します。すでに十分な並列処理が行われている場合は、同期のオーバーヘッドを増やす以外は役に立ちません。

于 2013-03-04T19:24:05.600 に答える