Certaines grilles enfants ne sont pas exécutées avec le parallélisme dynamique CUDA

Question

J'expérimente la nouvelle fonctionnalité de parallélisme dynamique dans CUDA 5.0 (GTK 110). Je suis confronté au comportement étrange que mon programme ne renvoie pas le résultat attendu pour certaines configurations, non seulement inattendu, mais aussi un résultat différent à chaque lancement.

Maintenant, je pense avoir trouvé la source de mon problème: Il semble que certains enfants (noyaux enserre lancés par d'autres noyaux) sont parfois exécutés lorsque trop de grilles d'enfants sont donné naissance à en même temps.

j'ai écrit ce comportement un petit programme de test pour illustrer:

#include <stdio.h> 

__global__ void out_kernel(char* d_out, int index) 
{ 
    d_out[index] = 1; 
} 

__global__ void kernel(char* d_out) 
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    out_kernel<<<1, 1>>>(d_out, index); 
} 

int main(int argc, char** argv) { 

    int griddim = 10, blockdim = 210; 
    // optional: read griddim and blockdim from command line 
    if(argc > 1) griddim = atoi(argv[1]); 
    if(argc > 2) blockdim = atoi(argv[2]); 

    const int numLaunches = griddim * blockdim; 
    const int memsize = numLaunches * sizeof(char); 

    // allocate device memory, set to 0 
    char* d_out; cudaMalloc(&d_out, memsize); 
    cudaMemset(d_out, 0, memsize); 

    // launch outer kernel 
    kernel<<<griddim, blockdim>>>(d_out); 
    cudaDeviceSynchronize(); 

    // dowload results 
    char* h_out = new char[numLaunches]; 
    cudaMemcpy(h_out, d_out, memsize, cudaMemcpyDeviceToHost); 

    // check results, reduce output to 10 errors 
    int maxErrors = 10; 
    for (int i = 0; i < numLaunches; ++i) { 
     if (h_out[i] != 1) { 
      printf("Value at index %d is %d, should be 1.\n", i, h_out[i]); 
      if(maxErrors-- == 0) break; 
     } 
    } 

    // clean up 
    delete[] h_out; 
    cudaFree(d_out); 
    cudaDeviceReset(); 
    return maxErrors < 10 ? 1 : 0; 
} 

Le programme lance un noyau dans un certain nombre de blocs donné (1er paramètre) avec un nombre donné de fils chacun (2ème paramètre). Chaque thread dans ce noyau lancera ensuite un autre noyau avec un seul thread. Ce noyau enfant écrira un 1 dans sa partie d'un tableau de sortie (qui a été initialisé avec 0s).

À la fin de l'exécution, toutes les valeurs du tableau de sortie doivent être égales à 1. Mais étrangement, pour certaines tailles de blocs et de grilles, certaines des valeurs de tableau sont toujours égales à zéro. Cela signifie que certaines grilles ne sont pas exécutées.

Cela se produit uniquement si plusieurs des grilles enfants sont générées en même temps. Sur mon système de test (un Tesla K20x) c'est le cas pour 10 blocs contenant 210 threads chacun. 10 blocs avec 200 threads fournissent le résultat correct, cependant. Mais aussi 3 blocs avec 1024 threads causent chacun l'erreur. Étrangement, aucune erreur n'est signalée par le moteur d'exécution. Les grilles d'enfant semblent simplement être ignorées par le planificateur.

Est-ce que quelqu'un d'autre fait face au même problème? Ce comportement est-il documenté quelque part (je n'ai rien trouvé), ou est-ce vraiment un bug dans l'exécution du périphérique?

Answer 1

Vous ne faites aucun error checking d'aucune sorte que je peux voir. Vous pouvez et devriez faire une vérification d'erreur similaire sur les lancements de noyau de périphérique. Reportez-vous à la documentation Ces erreurs ne seront pas nécessairement barboter à l'hôte:

Les erreurs sont enregistrées par thread, de sorte que chaque fil peut identifier l'erreur la plus récente qu'elle a généré.

Vous devez les piéger dans l'appareil. Il existe de nombreux exemples de ce type de vérification d'erreur de périphérique dans la documentation.

Si vous effectuez une vérification d'erreur correcte, vous découvrirez que dans chaque cas où un noyau n'a pas pu démarrer, l'API d'exécution de périphérique cuda renvoyait l'erreur 69, cudaErrorLaunchPendingCountExceeded.

Si vous analysez le documentation pour cette erreur, vous trouverez ceci:

cudaLimitDevRuntimePendingLaunchCount

Contrôle la quantité de mémoire réservée pour le lancement du noyau tampon qui n'ont pas encore commencé à exécuter, en raison soit à des dépendances non résolues ou à un manque de ressources d'exécution. Lorsque le tampon est plein, les lancements définissent la dernière erreur du thread sur cudaErrorLaunchPendingCountExceeded. Le nombre de lancements en attente par défaut est de 2048 lancements.

À 10 blocs * 200 threads, vous lancez 2000 noyaux, et tout semble fonctionner.

À 10 blocs * 210 threads, vous lancez 2100 noyaux, ce qui dépasse la limite de 2048 mentionnée ci-dessus.

Notez que c'est quelque peu dynamique dans la nature; Selon la façon dont votre application lance les noyaux enfants, vous pouvez lancer plus de 2048 noyaux facilement sans atteindre cette limite. Mais puisque votre application lance tous les noyaux approximativement simultanément, vous atteignez la limite.

Une vérification correcte des erreurs cuda est recommandée chaque fois que votre code CUDA ne se comporte pas comme prévu.

Si vous souhaitez obtenir une confirmation de ce qui précède, dans votre code, vous pouvez modifier votre noyau principal comme ceci:

__global__ void kernel(char* d_out) 
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    out_kernel<<<1, 1>>>(d_out, index); 
// cudaDeviceSynchronize(); // not necessary since error 69 is returned immediately 
    cudaError_t err = cudaGetLastError(); 
    if (err != cudaSuccess) d_out[index] = (char)err; 
} 

La limite de comptage de lancement en attente est modifiable. Reportez-vous à la documentation pour cudaLimitDevRuntimePendingLaunchCount