Probleme mit CUDA -Streams

Question

Ich betreibe Cublas v2.0 auf verschiedenen Streams auf einer einzelnen GPU (Tesla C2050), indem ich die Eingangsmatrizen (a [x/num_of_streams*y] unterteilt.B [xy] = c [x/num_of_streams*y]), aber irgendwie dauert es mehr Zeit, wenn ich CUDA -Streams verwende. Hier ist der Code -Snippet:

             //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");
              }

Daher hüpfe ich zwischen Streams hin und her und zugewiesene Arbeiten, erwarte, eine bessere Ausführungszeit zu bekommen, aber ich stelle fest, dass das Programm mehr Zeit in Anspruch nimmt, im Vergleich zu der Version, die keinen Stream verwendet. Wo gehe ich falsch? Cross Post zu Nvidia -Foren - http://forums.nvidia.com/index.php?showtopic=209420

BEARBEITEN:

Ich habe mein Programm wie folgt geändert:

            //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]);
            }

Wenn ich mein Programm für eine Matrixreihenfolge von 6144 profiliere, erhalte ich die folgenden Informationen:

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)

Wenn ich die obige Schleife habe, bekomme ich eine Zeit von 0,000284s, gegenüber 1,703289s für die Version, die keine Streams verwendet (in dieser Version habe ich auch die beiden sequentiellen Speicherkopien, Kernel -Aufruf und die verbleibenden Memcpy). Ich denke, da ich keine Synchronisationskonstrukte verwende, drucke ich die Zeit, bevor die Berechnung tatsächlich abgeschlossen ist (ich finde es schwer zu glauben, dass es eine 100% ige Verbesserung gibt).

Answer 1

Ich schlage zwei Änderungen vor:

1) Bewegen Sie Ihre Krublas, die die Erstellung/Zerstörung auf die Kopien und Kernelaufrufe bearbeiten. Es ist möglich, dass es die Parallelität bricht, indem sie einen nicht benötigten Kontext synchronisieren.

2) Führen Sie die Memcpy in einer Schleife über den Strömen zusammen. Auf diese Weise führt die B -Kopie von Stream 0 keine zusätzliche Synchronisation durch, um zu warten, bis die A memcpy abgeschlossen ist. dh das tun:

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

nicht das:

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

Seien Sie nicht überrascht, wenn Sie keine Beschleunigung von mehr als 40% von überlappenden Übertragungen und Berechnungen erhalten können. Streams bieten die größten Vorteile bei Workloads, die gleichzeitig übertragen und verarbeitet werden, und nur sehr wenige Workloads fallen in diese Kategorie.

Answer 2

Ich würde auch empfehlen, die Größe der Kopien zu überprüfen. Sie sollten nur dann mit verschiedenen Streams beginnen, wenn der Zeitraum für die Übertragung eines Speicherblocks mit der Zeit verglichen werden kann, die für die Berechnung erforderlich ist. Wenn die Übertragungszeit im Vergleich zur Berechnungszeit wenig ist, fügen Sie mit ihrem Management mehr Overhead hinzu. Verwenden Sie den visuellen Profiler, um zu sehen, wie lange die verschiedenen Schritte dauern. Erstellen Sie ein Diagramm mit verschiedenen Speichereingängen.