pauses de multiplication CUDA matrice pour les grandes matrices

Question

J'ai le code de multiplication matricielle, mis en œuvre en utilisant CUDA 3.2 et VS 2008. Je suis en cours d'exécution sur Windows Server 2008 R2 Enterprise. J'exécute une Nvidia GTX 480. Le code suivant fonctionne très bien avec les valeurs de « largeur » (matrice largeur) 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 Matrix 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);

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

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

//Free Device Memory
cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);

Quand je mets la « largeur » à 3000 ou plus, je reçois l'erreur suivante après un écran noir:

J'ai regardé en ligne et j'ai vu que certaines personnes a cette question parce que le chien de garde a été en train de tuer le noyau après qu'il se bloque pendant plus de 5 secondes. J'essayé de modifier la « TdrDelay » dans le registre et ce qui a retardé le temps avant que l'écran noir et même erreur apparaissent. Donc, je concluais ce ne fut pas mon problème.

Je déboguée dans mon code et a trouvé cette ligne pour être le coupable:

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

Voici ce que j'utilise pour retourner mon jeu de résultats de l'appareil après ma fonction de noyau de multiplication de la matrice est appelée. Tout jusqu'à ce point semble fonctionner très bien. Je crois que je suis allouez de la mémoire correctement et ne peut pas comprendre pourquoi ce qui se passe. Je pensais que peut-être je ne l'ai pas assez de mémoire sur ma carte pour cela, mais alors ne devrait pas cudaMalloc suis retourné une erreur? (Je confirme qu'il n'a pas en cours de débogage).

Toutes les idées / aide serait grandement appréciée! ... Merci beaucoup les gars !!

Code du noyau:

//Matrix Multiplication Kernel - Multi-Block Implementation
__global__ void MatrixMultiplicationMultiBlock_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 il donne également la même erreur:

Appel:

            MatrixMultiplicationSharedMemory_Kernel<<<dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2>>>(Md, Nd, Pd, Width);

Code du noyau:

 //Matrix Multiplication Kernel - Shared Memory Implementation
 __global__ void MatrixMultiplicationSharedMemory_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;
}

Answer 1

Contrôle du délai d'attente WDDM

Le problème est en fait le noyau de ne pas le cudaMemcpy(). Lorsque vous lancez le noyau le GPU se déclenche et fait le travail de manière asynchrone avec l'unité centrale de traitement, il est donc que lorsque vous effectuez une synchronisation avec le GPU que vous devez attendre le travail à la fin. cudaMemcpy() implique une synchronisation implicite, donc qui est là où vous voyez le problème.

Vous pourriez bien vérifier en appelant cudaThreadSynchronize() après que le noyau et le problème semble être sur la cudaThreadSynchronize() au lieu du cudaMemcpy().

Après avoir modifié le délai d'attente du TDR, avez-vous redémarrer votre machine? Malheureusement, Windows doit être redémarré pour modifier les paramètres du TDR. Ce document Microsoft a une assez bonne description des paramètres complets sont disponibles.

problèmes du noyau

Dans ce cas, le problème est pas réellement le délai d'attente de WDDM. Il y a des erreurs dans le noyau que vous devez résoudre (par exemple, vous devriez être en mesure de incremement i par plus d'un à chaque itération) et vérifier l'échantillon matrixMul dans le SDK peut être utile. Soit dit en passant, j'espère que cela est un exercice d'apprentissage puisque, en réalité, vous seriez mieux (pour la performance) en utilisant CUBLAS pour effectuer la multiplication de matrices.

Le problème le plus critique dans le code est que vous utilisez la mémoire partagée sans allouer effectivement tout. Dans votre noyau, vous avez:

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

Mais quand vous lancez le noyau que vous ne spécifiez pas la quantité de mémoire à allouer pour chaque bloc:

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

La syntaxe <<< >>> prend en fait quatre arguments où les troisième et quatrième sont en option. Le quatrième est l'indice de flux qui est utilisé pour obtenir un chevauchement entre calcul et de transfert de données (et pour l'exécution du noyau en même temps), mais le troisième argument spécifie la quantité de mémoire partagée par bloc. Dans ce cas, je suppose que vous voulez flotte TileWidth * TileWidth magasin dans la mémoire partagée, de sorte que vous utilisez:

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock, dimBlock.x * dimBlock.x * sizeof(float)>>>(Md, Nd, Pd, Width);

Le principal problème

Comme vous le mentionnez dans votre commentaire, le problème réel est que votre largeur matrice était pas un multiple de la largeur du bloc (et la hauteur, car il est carré, ce qui signifie les fils au-delà de la fin accéderaient au-delà de la fin du tableau. La soit code doit traiter le cas de non-multiple ou elle doit veiller à ce que la largeur est un multiple de la taille de bloc.

Je aurais dû suggéré plus tôt, mais il est souvent utile d'exécuter cuda-memcheck pour vérifier les violations d'accès memeory comme celui-ci.

Answer 2

Vous devez modifier les paramètres de temporisation pilote, est caractéristique fenêtres pour empêcher les conducteurs défectueux pour rendre le système ne répond pas. Vérifiez la Microsoft décrivant comment faire.

Answer 3

Vous devriez également vérifier le réglage du drapeau « délai d'attente » sur votre appareil GPU. Si vous avez le SDK CUDA installé, je crois que l'application « deviceQuery » rapportera cette propriété.