Problèmes avec les flux CUDA

Question

Je cours v2.0 CUBLAS sur différents flux sur une seule GPU (Tesla C2050) en subdivisant les matrices d'entrée (A [X / num_of_streams * y] B [x y] = C [ x / y num_of_streams *]), mais en quelque sorte, il prend plus de temps lors de l'utilisation des flux CUDA. Voici l'extrait de code:

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

Je suis rebondissant entre les flux et le travail assignant, en attendant d'obtenir un meilleur temps d'exécution, mais je remarque que plus le nombre de flux, le programme prend plus de temps par rapport à la version qui n'utilise pas le flux. Où vais-je tort? poste de la Croix à des forums Nvidia - http://forums.nvidia.com/index.php? ShowTopic = 209420

EDIT:

Je modifié mon programme comme suit:

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

Quand je profil mon programme pour une commande de matrice de 6144, je reçois les informations suivantes:

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)

Quand je fois que la boucle au-dessus, je reçois un temps de 0.000284s, vs 1.703289s pour la version qui n'utilise pas les flux (dans cette version aussi, je fois que les deux copies de mémoire séquentielle, invocation du noyau et le Memcpy restant) . Je pense que je ne suis pas en utilisant des constructions de synchronisation, peut-être que je suis en train d'imprimer le temps avant que le calcul en fait finis (je trouve qu'il est difficile de croire qu'il ya une amélioration de 100%).

Answer 1

Je suggère deux changements:

1) déplacer votre cuBLAS création poignée / destruction à l'extérieur des copies et des invocations du noyau. Il est possible qu'il se brise en faisant une concurrence synchronisent contexte non nécessaire.

2) do ensemble du memcpy dans une boucle sur les cours d'eau. De cette façon, la copie B du flux 0 ne fait pas de synchronisation supplémentaire pour attendre jusqu'à ce que l'un memcpy est terminée. à savoir faire ceci:

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

pas:

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

Ne soyez pas surpris si vous ne parvenez pas à obtenir un gain de vitesse de plus de 40%, ou des transferts qui se chevauchent et le calcul. Cours d'eau offrent les plus grands avantages sur les charges de travail qui passent autant de temps de transfert et de traitement des données, et très peu de charges de travail tombent dans cette catégorie.

Answer 2

Je suggère également de vérifier la taille des copies, vous devriez commencer à utiliser différents flux seulement lorsque le temps de transférer un bloc de mémoire peut être comparé au temps nécessaire pour le calcul sur elle. Si le temps de transfert est peu par rapport au temps de calcul, puis en ajoutant les flux ajoutent plus de frais généraux à leur gestion. Utilisez le Visual Profiler pour voir combien de temps il prend les différentes étapes. Faire un graphique avec différentes entrées de mémoire.