2013-03-17 4 views
1

j'ai mesuré la bande passante lors de l'exécution d'un noyau faisant une addition entre deux vecteurs:CUDA: vecteurs addition et vecteurs taille

__global__ void add(float *a, float *b, float *c, int n) 
{ 
    int tid = blockIdx.x*blockDim.x + threadIdx.x; 

    while (tid < n) 
    { 
     c[tid] = a[tid] + b [tid]; 
     tid += blockDim.x * gridDim.x; 
    } 
} 

je tout d'abord lance une fois le noyau, de sorte qu'il est chargé sur l'appareil, et à côté , Je mesure 10 itérations de l'exécution du noyau.

La bande passante est de loin mieux quand mon vectors'length est 1000000 à 1000.

Pourquoi?

Merci.

+2

Le GPU peut masquer la latence de la mémoire lorsqu'il a plus de travail à faire. –

Répondre

5

Les GPU ont une bande passante de mémoire très élevée (bonne) et très haute latence de la mémoire (mauvaise). Il faut plusieurs centaines de cycles pour qu'une demande de mémoire soit satisfaite, mais il peut y avoir beaucoup de requêtes en vol qui sont pipelinées.

Si vous ajoutez seulement 1000 éléments, ce que vous mesurez réellement est la latence de la mémoire. Avec des vecteurs beaucoup plus gros, il y aura suffisamment de blocs de threads qui, quand on est bloqué en mémoire, le GPU peut passer à un autre bloc de threads dont la requête de mémoire est en cours de saturation. Avec eux, vous mesurez la bande passante.

+0

Mais d'après ce que j'ai lu, les données devraient être lues par blocs de 128 octets (warps, coalescing ...). Si oui, il devrait y avoir une latence pour chaque bloc, et comme je compare les éléments 1000 et 1000000, la différence devrait être linéaire, non?
Il y aurait une différence si les données étaient lues une fois dans les deux cas. Que je comprendrais pourquoi la latence est relativement plus importante avec 1000 elemnts. – Oodini

+0

Ce n'est pas linéaire en raison du pipelining. Supposons que chaque bloc lit 128 octets et qu'il nécessite 300 cycles de latence pour une demande de mémoire. Supposons que vous avez 300 blocs. Sur l'horloge tick 0, le premier bloc effectue une requête mémoire qui sera satisfaite par l'horloge tick 300. A l'horloge tick 1, un autre bloc s'exécute (les commutateurs de contexte sont libres sur les GPU), et il calera jusqu'à tick 301. Le dernier bloc émettra sa demande de mémoire à la coche 299 et il sera satisfait à la coche 599. Ainsi, tous les 300 blocs se terminera en 600 cycles au total, pas 300 * 300 = 90 000. La latence réelle dépend de votre carte particulière. –