J'ai le code de multiplication matricielle suivant, mis en œuvre à l'aide de CUDA 3.2 et VS 2008. Je suis en cours d'exécution sur Windows Server 2008 R2 Enterprise. Je suis en cours d'exécution d'un NVIDIA GTX 480. Le code suivant fonctionne bien avec des valeurs de "largeur" (largeur de matrice) jusqu'à environ 2500 environ.
//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;
}
3 Réponses :
contrôler le délai d'attente WDDM Le problème est en réalité le noyau non le noyau n'est pas le Vous pouvez double vérifier cela en appelant Après avoir changé l'heure du TDR, avez-vous Redémarrez votre machine? Malheureusement, Windows doit être redémarré pour modifier les paramètres TDR. Ce document Microsoft a une assez bonne description des paramètres complets disponibles. < / p> problèmes de noyau Dans ce cas, le problème n'est pas réellement le délai d'attente WDDM. Il y a des erreurs dans le noyau que vous auriez besoin de résoudre (par exemple, vous devriez être en mesure d'augmenter Le problème le plus critique du code est que vous utilisez la mémoire partagée sans réellement allouer. Dans votre noyau, vous avez: mais lorsque vous lancez le noyau, vous ne spécifiez pas la quantité de mémoire partagée à allouer pour chaque bloc: La syntaxe <<<> prend en réalité quatre arguments où les troisième et quatrième sont facultatifs. Le quatrième est l'index de flux utilisé pour obtenir un chevauchement entre le virement de calcul et le transfert de données (et pour l'exécution simultanée du noyau), mais l'argument Le problème principal Comme vous le mentionnez dans votre commentaire, le problème réel était que votre largeur de matrice n'était pas un multiple de la largeur de bloc (et la hauteur car elle est carrée, ce qui signifie que les threads au-delà de l'extrémité n'auraient accès au-delà de la fin de la matrice. Le code doit soit gérer le boîtier non multiple, soit il convient de garantir que la largeur est un multiple de la taille du bloc. J'aurais dû a suggéré cela plus tôt, mais il est souvent utile d'exécuter cudamemcpy () . Lorsque vous lancez le noyau, le GPU s'éteint et fait le travail de manière asynchrone avec la CPU, ce n'est donc que lorsque vous synchronisez avec le GPU que vous devez attendre que le travail se termine. cudamemcpy () implique une synchronisation implicite, c'est donc là où vous voyez le problème. cudathreadsynchroniser () après le noyau et le problème sembleront figurera sur le cudathreadsynchroniser () au lieu du cudamemcpy () . i par plus d'un sur chaque itération) et de vérifier le Matrixmul Sample dans le SDK peut être utile. Incidemment, j'espère que c'est un exercice d'apprentissage depuis en réalité, vous seriez mieux désactivé (pour la performance) à l'aide de Cubla pour effectuer une multiplication de matrice. TILEWIDTH * TILEWIDTH flotte dans la mémoire partagée, de sorte que vous utiliseriez: cuda-memcheck pour vérifier les violations de l'accès des memes comme celle-ci.
J'ai fait le redémarrage et les changements ont pris effet, car l'écran noir a pris plus de temps à apparaître .. Cependant, il s'est toujours écrasé ....
Je vais essayer le cudathreadsynchroniser et republier quand je rentre à la maison!
Ok, tu avais raison. Je viens de faire cela et j'ai eu la même erreur ... J'ai ajouté TdrDelay en tant que REG_DWORD à HKEY_LOCAL_MACHINE \ SYSTEM \ CurrentControlset \ Contol \ GraphicsD Rivières. J'ai redémarré ma machine et j'ai remarqué qu'il a fallu plus longtemps pour que l'écran soit noir et que l'erreur apparaisse .. à peu près aussi longtemps que je fixe le retard pour .. mais cela ne fonctionne toujours pas. Je ne suis pas entièrement convaincu que c'est un retard, car cela traite une largeur de 2500, mais quelque chose de beaucoup plus que cela et il se bloque .. même 2800 ... Je manque quelque chose?
Je remarque également que lorsque j'ajoute que j'ajoute Cudathreadsynchroniser (), il se casse pour des matrices de tailles qu'elle travaillait avant ... 2500 le casse avec cette ligne ajoutée. (Je ne sais pas si cela vous donne des indices ou non)
IIUC jusqu'à ~ 2500 Il se termine rapidement (disons moins de quelques secondes), mais à ce sujet, il ne finit jamais, peu importe la sommet que vous définissez le TDR. Vous pouvez vérifier cela en définissant le TDR assez haut, juste être patient! Cela étant le cas, la prochaine étape serait de regarder le noyau.
Merci Tom, en fait j'ai posté la mauvaise signature .. J'ai deux fonctions matricielles .. L'une optimisation de la mémoire partagée et celle qui n'est pas ... Les deux casseront pour de grandes valeurs ... Je vais éditer pour mettre à jour l'appel de la fonction de droite. . Mais en fait j'ai alloué la mémoire.
Désolé pour la confusion = / ... J'ai mis à jour le poste, et j'ai réellement attribué TileWidth * TILEWIDTH * TIZEOFOF * 2 ... depuis que j'ai deux matrices NDS et MDS que j'utilise pour stocker mes valeurs. être ok, non?
Oui, mais concentrez-vous sur la version mémoire non partagée pour l'instant depuis que cela échoue de la même manière. Mieux ne pas surcharger!
Vous devez modifier les paramètres du timeau du pilote, est la fonction Windows pour empêcher les pilotes défectueux de rendre le système insensible. Consultez le Microsoft Page Décrivez comment faire cela.
Devrais-je essayer autre chose que TdrDelay?
Vous devez également vérifier le paramètre d'indicateur "Timeout" sur votre périphérique GPU. Si vous avez installé le CUDA SDK, je pense que l'application "DeviceQuery" signalera cette propriété.
Merci pour votre réponse! Où dois-je modifier cette propriété?
Je ne sais pas comment le modifier - c'est quelque chose que le conducteur traite. Il peut être associé à si vous avez un affichage connecté à l'appareil.
Pouvez-vous poster le code du noyau s'il vous plaît?
EDIT: Ajout des fonctions de code du noyau
Pour complétude, pouvez-vous poster l'assidation Dimgrid / DimBlock s'il vous plaît? Je suppose que dimgrid.x = dimgrid.y et dimblock.x = dimblock.y, et cette largeur est un multiple de dimblock.x?
@Tom, j'ai ajouté ma définition Dimgrid and DimBlock et là-bas dans le problème. Je n'utilisais pas de multiple de DimBlock ... mon code de noyau n'était pas manipulant cela correctement .. merci tellement !!!!, s'il vous plaît ajouter Ce commentaire à votre réponse ci-dessous pour que je puisse le choisir. Merci encore!!!
Je t'en prie. J'aurais suggéré d'exécuter
cuda-memcheck plus tôt, désolé à ce sujet. Il devrait détecter immédiatement les accès hors limites.