Probleme mit CUDA Streams

Question

Ich lasse CUBLAS v2.0 auf verschiedenen Streams auf einer einzelnen GPU (Tesla C2050) laufen, indem ich die Eingabematrizen unterteile (A[x/num_of_streams*y] B[x y] = C[x/num_of_streams*y]), aber irgendwie braucht es mehr Zeit, wenn ich CUDA-Streams verwende. Hier ist der Codeschnipsel:

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

Ich springe also zwischen Streams hin und her und weise Arbeit zu, in der Erwartung, eine bessere Ausführungszeit zu erhalten, aber ich stelle fest, dass das Programm mit zunehmender Anzahl von Streams mehr Zeit benötigt als die Version, die keinen Stream verwendet. Wo mache ich etwas falsch? Querverweis auf Nvidia-Foren - http://forums.nvidia.com/index.php?showtopic=209420

EDITAR:

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 Matrixordnung von 6144 profiliere, erhalte ich die folgende Information:

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 Zeit für die obige Schleife messe, erhalte ich eine Zeit von 0,000284s, im Vergleich zu 1,703289s für die Version, die keine Streams verwendet (in dieser Version messe ich auch die beiden sequentiellen Speicherkopien, den Kernelaufruf und das verbleibende memCpy). Ich denke, da ich keine Synchronisationskonstrukte verwende, kann es sein, dass ich die Zeit ausdrucke, bevor die Berechnung tatsächlich abgeschlossen ist (ich finde es schwer zu glauben, dass es eine 100%ige Verbesserung gibt).