Problèmes avec les flux CUDA

Question

Je cours CUBLAS v2.0 sur différents flux sur un seul GPU (Tesla C2050) en subdivisant les matrices d'entrée (A [x/num_of_streams * y] B [x y] = C [x/num_of_streams * y]), mais en quelque sorte cela prend plus de temps quand j'utilise 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 l'attribution du travail, 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 que par rapport à la version qui n'utilise pas de flux. Où vais-je mal? 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 chronomètre la boucle ci-dessus, j'obtiens un temps de 0,000284s, vs 1,703289s pour la version qui n'utilise pas de flux (dans cette version aussi, je chronomètre les deux copies séquentielles, l'invocation du noyau et le memCpy restant). Je pense que puisque je n'utilise pas de constructions de synchronisation, il se peut que j'imprime le temps avant que le calcul ne finisse réellement (j'ai du mal à croire qu'il y a une amélioration de 100%).

Answer 1

Je propose deux changements:

1) déplacer votre cuBLAS gérer la création/destruction à l'extérieur des copies et des invocations du noyau. Il est possible qu'il casse la concurrence en effectuant une synchronisation de contexte inutile. 2) Faire les memcpy ensemble dans une boucle sur les flux. De cette façon, la copie B du flux 0 n'effectue aucune synchronisation supplémentaire pour attendre que le memcpy A soit terminé. à savoir faire:

 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. Les flux offrent les plus grands avantages sur les charges de travail qui passent un temps égal à transférer et traiter les données, et très peu de charges de travail tombent dans cette catégorie.

Answer 2

Je voudrais également suggérer de vérifier la taille des copies, vous devriez commencer à utiliser des flux différents seulement lorsque le temps de transférer un bloc de mémoire peut être comparé au temps nécessaire pour le calculer. Si le temps de transfert est faible par rapport au temps de calcul, l'ajout de flux ajoute une charge supplémentaire à leur gestion. Utilisez Visual Profiler pour voir combien de temps les différentes étapes sont nécessaires. Faire un graphique avec différentes entrées de mémoire.