Pauses de multiplication de masortingces CUDA pour les grandes masortingces

J’ai le code de multiplication de masortingce suivant, mis en œuvre à l’aide de CUDA 3.2 et VS 2008. J’utilise Windows Server 2008 R2 Enterprise. J’exécute une Nvidia GTX 480. Le code suivant fonctionne très bien avec les valeurs de “Largeur” ​​(largeur de masortingce) allant jusqu’à environ 2500 ou plus.

int size = Width*Width*sizeof(float); float* Md, *Nd, *Pd; cudaError_t err = cudaSuccess; //Allocate Device Memory for M, N and P err = cudaMalloc((void**)&Md, size); err = cudaMalloc((void**)&Nd, size); err = cudaMalloc((void**)&Pd, size); //Copy Masortingx from Host Memory to Device Memory err = cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); err = cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice); //Setup the execution configuration dim3 dimBlock(TileWidth, TileWidth, 1); dim3 dimGrid(ceil((float)(Width)/TileWidth), ceil((float)(Width)/TileWidth), 1); MasortingxMultiplicationMultiBlock_Kernel<<>>(Md, Nd, Pd, Width); err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); //Free Device Memory cudaFree(Md); cudaFree(Nd); cudaFree(Pd); 

Lorsque je règle la “Largeur” ​​sur 3000 ou plus, l’erreur suivante apparaît après un écran noir: capture d'écran

J’ai regardé en ligne et j’ai vu que certaines personnes avaient ce problème parce que le chien de garde était en train de tuer le kernel après l’avoir suspendu plus de 5 secondes. J’ai essayé de modifier le “TdrDelay” dans la base de registre, ce qui a retardé l’affichage de l’écran noir et de la même erreur. J’ai donc conclu que ce n’était pas mon problème.

J’ai débogué dans mon code et découvert que cette ligne était le coupable:

 err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); 

C’est ce que j’utilise pour renvoyer mon jeu de résultats à partir du périphérique après l’appel de ma fonction de kernel de multiplication de masortingce. Jusqu’ici, tout semble bien fonctionner. Je crois que j’atsortingbue la mémoire correctement et ne peux pas comprendre pourquoi cela se produit. Je pensais que je n’avais peut-être pas assez de mémoire sur ma carte pour cela, mais cudaMalloc n’aurait-il pas dû renvoyer une erreur? (J’ai confirmé que ce n’était pas le cas lors du débogage).

Toute idée / assistance serait grandement appréciée! … Merci beaucoup les gars !!

Code du kernel:

 //Masortingx Multiplication Kernel - Multi-Block Implementation __global__ void MasortingxMultiplicationMultiBlock_Kernel (float* Md, float* Nd, float* Pd, int Width) { int TileWidth = blockDim.x; //Get row and column from block and thread ids int Row = (TileWidth*blockIdx.y) + threadIdx.y; int Column = (TileWidth*blockIdx.x) + threadIdx.x; //Pvalue store the Pd element that is computed by the thread float Pvalue = 0; for (int i = 0; i < Width; ++i) { float Mdelement = Md[Row * Width + i]; float Ndelement = Nd[i * Width + Column]; Pvalue += Mdelement * Ndelement; } //Write the matrix to device memory each thread writes one element Pd[Row * Width + Column] = Pvalue; } 

J’ai aussi cette autre fonction qui utilise la mémoire partagée et qui donne la même erreur:

Appel:

  MasortingxMultiplicationSharedMemory_Kernel<<>>(Md, Nd, Pd, Width); 

Code du kernel:

  //Masortingx Multiplication Kernel - Shared Memory Implementation __global__ void MasortingxMultiplicationSharedMemory_Kernel (float* Md, float* Nd, float* Pd, int Width) { int TileWidth = blockDim.x; //Initialize shared memory extern __shared__ float sharedArrays[]; float* Mds = (float*) &sharedArrays; float* Nds = (float*) &Mds[TileWidth*TileWidth]; int tx = threadIdx.x; int ty = threadIdx.y; //Get row and column from block and thread ids int Row = (TileWidth*blockIdx.y) + ty; int Column = (TileWidth*blockIdx.x) + tx; float Pvalue = 0; //For each tile, load the element into shared memory for( int i = 0; i < ceil((float)Width/TileWidth); ++i) { Mds[ty*TileWidth+tx] = Md[Row*Width + (i*TileWidth + tx)]; Nds[ty*TileWidth+tx] = Nd[(ty + (i * TileWidth))*Width + Column]; __syncthreads(); for( int j = 0; j < TileWidth; ++j) { Pvalue += Mds[ty*TileWidth+j] * Nds[j*TileWidth+tx]; } __syncthreads(); } //Write the matrix to device memory each thread writes one element Pd[Row * Width + Column] = Pvalue; } 

    Contrôle du délai d’attente WDDM

    Le problème est en réalité le kernel et non le cudaMemcpy() . Lorsque vous lancez le kernel, le processeur graphique s’éteint et effectue le travail de manière asynchrone avec le processeur. C’est donc uniquement lorsque vous synchronisez avec le processeur graphique que vous devez attendre la fin du travail. cudaMemcpy() implique une synchronisation implicite, d’où le problème.

    Vous pouvez vérifier cela en appelant cudaThreadSynchronize() après le kernel et le problème semblera être cudaThreadSynchronize() au lieu de cudaMemcpy() .

    Après avoir modifié le délai d’attente TDR, avez-vous redémarré votre ordinateur? Malheureusement, Windows doit être redémarré pour modifier les parameters TDR. Ce document Microsoft contient une assez bonne description de tous les parameters disponibles.

    Problèmes de kernel

    Dans ce cas, le problème n’est pas le délai d’attente WDDM. Il y a des erreurs dans le kernel que vous devez résoudre (par exemple, vous devriez pouvoir incrémenter i de plusieurs erreurs à chaque itération) et extraire l’exemple masortingxMul du kit SDK peut être utile. Par ailleurs, j’espère que c’est un exercice d’apprentissage, car en réalité, vous auriez intérêt à utiliser CUBLAS pour effectuer la multiplication masortingcielle.

    Le problème le plus critique dans le code est que vous utilisez la mémoire partagée sans en allouer réellement. Dans votre kernel, vous avez:

     //Initialize shared memory extern __shared__ float sharedArrays[]; 

    Mais lorsque vous lancez le kernel, vous ne spécifiez pas la quantité de mémoire partagée à allouer pour chaque bloc:

     MasortingxMultiplicationMultiBlock_Kernel<<>>(Md, Nd, Pd, Width); 

    La syntaxe <<< >>> prend en fait quatre arguments, le troisième et le quasortingème étant facultatifs. Le quasortingème est l’index de stream utilisé pour obtenir un chevauchement entre le calcul et le transfert de données (et pour l’exécution simultanée du kernel), mais le troisième argument spécifie la quantité de mémoire partagée par bloc. Dans ce cas, je suppose que vous souhaitez stocker TileWidth * TileWidth flottant dans la mémoire partagée, vous devez donc utiliser:

     MasortingxMultiplicationMultiBlock_Kernel<<>>(Md, Nd, Pd, Width); 

    Le problème principal

    Comme vous l’avez mentionné dans votre commentaire, le problème réel était que la largeur de votre masortingce n’était pas un multiple de la largeur du bloc (et de la hauteur, car elle est carrée, ce qui signifie que les threads situés au-delà de la fin accéderaient au-delà de la fin du tableau. gérer le cas non-multiple ou il devrait s’assurer que la largeur est un multiple de la taille du bloc.

    J’aurais dû le suggérer plus tôt, mais il est souvent utile d’exécuter cuda-memcheck pour vérifier les violations d’access mémoire telles que celle-ci.

    Vous devez modifier les parameters de délai d’exécution du pilote. Cette fonctionnalité de Windows empêche les pilotes défectueux de ne pas répondre au système. Vérifiez la page Microsoft décrivant comment faire cela.

    Vous devez également vérifier le paramètre “timeout” sur votre périphérique GPU. Si vous avez installé le SDK CUDA, je pense que l’application “deviceQuery” signalera cette propriété.