2

入力行列（A [x / num_of_streams * y] B [x y] = C [x / num_of_streams * y]）を細分化することにより、単一のGPU（Tesla C2050）上のさまざまなストリームでCUBLAS v2.0を実行していますが、どういうわけかCUDAストリームを使用すると時間がかかります。コードスニペットは次のとおりです。

             //plan is a struct containing the matrix dimensions and stream numbers
             //parallel in nstreams - should be! MAX 16 streams could run concurrently
            //Copy A - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyA_in_streams (&plan[i]);
            //Copy B - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyB_in_streams (&plan[i]);

            //Create handles - serial
            for(i = 0; i < nstreams; i++)
                    handle[i] = create_handle();

            //Run kernels - first doing a cublasSetStream(handle, plan->stream) before running cublasDgemm... 
            for(i = 0; i < nstreams; i++)
                    cudgemm_kernel_in_streams (&plan[i], handle[i], 1.0f, 1.0f);

            //Destroy handles - serial
            for(i = 0; i < nstreams; i++)
                    destroy_handle (handle[i]);

            //Copy C - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyC_in_streams (&plan[i]);

            //EDIT: Function body

            //The other two copy functions are exactly the same as this
            void cudgemm_copyA_in_streams(TGPUplan *plan)
           {
                 cudasafe(cudaMemcpyAsync(plan->Ad_Data, plan->Ah_Data, (plan->Acols * plan->Arows * sizeof(double)), cudaMemcpyHostToDevice, plan->stream) );

            }

            //Create handle
            cublasHandle_t create_handle ()
            {
                   cublasHandle_t handle;
                   checkError(cublasCreate(&handle), "cublasCreate() error!\n");
                   return handle;
             }

             //Destroy handle
             void destroy_handle (cublasHandle_t handle)
             {
                  checkError(cublasDestroy(handle), "cublasDestroy() error!\n");
             }

             //Kernel
             void cudgemm_kernel_in_streams(TGPUplan *plan, cublasHandle_t handle, const double alpha, const double beta)
             {
                   cublasStatus_t ret;
                   cublasSetStream(handle, plan->stream);

                   ret = cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, plan->Arows, plan->Ccols, plan->Acols, &alpha, plan->Ad_Data, plan->Arows, plan->Bd_Data, plan->Brows, &beta, plan->Cd_Data, plan->Crows);
                   checkError(ret, "cublas Dgemm returned an error!\n");
              }

そのため、ストリーム間を行き来して作業を割り当て、実行時間が短縮されることを期待していますが、ストリームの数が多いほど、ストリームを使用しないバージョンと比較して、プログラムに時間がかかることに気付きます。どこが間違っているのですか？Nvidiaフォーラムへのクロスポスト-http://forums.nvidia.com/index.php?showtopic= 209420

編集：

プログラムを次のように変更しました。

            //copy data
            for(i = 0; i < nstreams; i++)
            {
                    cudgemm_copyA_in_streams (&plan[i]);
                    cudgemm_copyB_in_streams (&plan[i]);
            }

            //Run kernel and copy back
            for(i = 0; i < nstreams; i++)
            {
                    cudgemm_kernel_in_streams (&plan[i], handle[i], 1.0f, 1.0f);
                    cudgemm_copyC_in_streams (&plan[i]);
            }

プログラムを6144の行列次数でプロファイリングすると、次の情報が得られます。

Kernel time = 42.75 % of total GPU time 
Memory copy time = 28.9 % of total GPU time
Kernel taking maximum time = fermiDgemm_v2_kernel_val (42.8% of total GPU time)
Memory copy taking maximum time = memcpyHtoDasync (21.7% of total GPU time)
Total overlap time in GPU = 65268.3 micro sec. (3.6% of total GPU time)

青=カーネル、緑=2ストリームのcudaMemCpyAsync

上記のループの時間を計ると、0.000284秒の時間が得られますが、ストリームを使用しないバージョンの場合は1.703289秒です（このバージョンでも、2つのシーケンシャルメモリコピー、カーネル呼び出しと残りのmemCpyの時間を計ります）。私は同期構造を使用していないので、計算が実際に終了するまでの時間を印刷しているのではないかと思います（100％の改善があるとは信じがたいです）。

4

2 に答える 2

2

私は2つの変更を提案します：

1）cuBLASハンドルの作成/破棄をコピーとカーネル呼び出しの外部に移動します。不要なコンテキスト同期を実行することにより、同時実行性を壊している可能性があります。

2）ストリーム上で1つのループでmemcpyを一緒に実行します。このように、ストリーム0のBコピーは、Amemcpyが完了するまで待機するための追加の同期を行いません。つまり、これを行います：

        for(i = 0; i < nstreams; i++) {
                cudgemm_copyA_in_streams (&plan[i]);
                cudgemm_copyB_in_streams (&plan[i]);
        }

これではない：

        for(i = 0; i < nstreams; i++)
                cudgemm_copyA_in_streams (&plan[i]);
        for(i = 0; i < nstreams; i++)
                cudgemm_copyB_in_streams (&plan[i]);

転送と計算の重複によって40％以上のスピードアップが得られなくても、驚かないでください。ストリームは、データの転送と処理に同じ時間を費やすワークロードに最大のメリットをもたらし、そのカテゴリに分類されるワークロードはごくわずかです。

于 2011-09-07T13:24:06.740 に答える
1

また、コピーのサイズを確認することをお勧めします。メモリの1つのブロックを転送する時間が、その計算に必要な時間と比較できる場合にのみ、異なるストリームの使用を開始する必要があります。転送にかかる時間が計算時間に比べて短い場合は、ストリームを追加すると、管理のオーバーヘッドが増えます。Visual Profilerを使用して、さまざまな手順にかかる時間を確認します。さまざまなメモリ入力でグラフを作成します。

于 2011-09-06T17:12:05.717 に答える