Introduction à la programmation GPU - Part 4
Rechercher plus de performances avec CUDA
Retour sur l’organisation de la mémoire
Dans l’article précédent, nous avons esquissé l’organisation de la mémoire en acceptant quelques simplifications didactiques. Notre première implémentation mérite un peu plus d’explications sur l’organisation mémoire sur une architecture CUDA. Pour faciliter notre compréhension, nous avons représenté sur le schéma ci-dessous (figure 1) une grille contenant deux blocs de threads (0, 0) et (1,0).
Figure 1 – Organisation de la mémoire vis-à-vis des threads GPUs
Au bas du schéma de la figure 1, deux morceaux de mémoire sont représentés: la mémoire globale (Global Memory) et la mémoire constante (Constant Memory) qui n’apparaissaient pas sur le schéma du précèdent article. On remarque que la partie host (code CPU) communique directement via l’interfacematérielle PCI Express avec ces deux types de mémoire. La mémoire constante est visible par tous les threads de la grille, mais elle est en lecture seule, offrant de ce fait des performances meilleures que la mémoire globale. La mémoire constante permet de stocker les variables « constantes » pour toute la durée d’exécution de l’application. Les variables constantes sont souvent représentées par les valeurs fournies en entrée des méthodes« noyau ». Les variables constantes sont stockées dans la mémoire globale, mais aussi cachées dans la mémoire constante pour offrir un accès plus rapide. Avec ce mécanisme de mise en cache, l’accès à la mémoire constante peut se révéler très intéressant sur le plan des performances lorsque plusieurs lectures se déclenchent en parallèle. La mémoire globale est visible par tous les threads de la grille en lecture et en écriture, mais elle est relativement lente. Son principal avantage est sa grande taille, dans la majorité des cas, vous ne vous poserez pas de question sur la taille des données que vous souhaitez copier. Les échanges entre l’host et la mémoire globale sont simples d’accès via les méthodes réalisant des copies entre l’host et la mémoire globale côté GPU. Pour les threads, l’accès à la mémoire est naturel (peut-être trop) et ne diffère pas du code traditionnel sur CPU.
En haut du schéma, deux blocs de threads sont représentés. Ces blocs possèdent chacun deux threads CUDA avec leurs registres respectifs, ainsi qu’une petite ligne de cache, appelée Shared Memory en dialecte CUDA. Cette mémoire partagée pour tous les threads du bloc courant est extrêmement rapide (de l’ordre de 10 fois plus rapide que la mémoire globale), mais relativement petite (de l’ordre de quelques kilos octets). Les accès à la mémoire partagée sont explicites et réclament de la part du développeur un développement spécifique. Enfin, notez que cette mémoire est partagée par tous les threads d’un bloc, permettant aux threads de collaborer pour accéder en lecture et en écriture aux données stockées en mémoire partagée avec d’excellentes performances. Cependant, nous pouvons déjà imaginer que les threads d’un même block devront se synchroniser pour ne pas aboutir à des traitements incohérents. Enfin, vous pouvez remarquer que les autres threads de l’autre bloc ne peuvent pas accéder à la mémoire partagée d’un autre bloc. La taille de la mémoire partagée étant très faible, le développeur devra généralement procéder à un découpage des données pour profiter des performances offertes par ce type de mémoire.
Retour sur notre algorithme de multiplication matricielle
Si nous analysons le code de la multiplication matricielle de précédenst articles, nous pouvons observer un défaut majeur. A votre avis, dans le code ci-dessous (Fragment de code 1) y a-t des recouvrements de valeurs dans le calcul matriciel ? En d’autres mots, les indices des matrices A et B sont-ils réutilisés plusieurs fois au cours du calcul ?
__global__
void MatrixMultiply(float* A, float* B, float* C, int size)
{
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
if (j < size && i < size)
{
float sum = 0;
for (int k = 0; k < size; ++k) {
sum += A[i * size + k] * B[k * size + j];
}
C[i * size + j] = sum;
}
}
Fragment de code 1 : code du calcul matriciel simple
Pour simplifier, prenons le cas de quatre threads traitant en parallèle le calcul du produit matriciel.
Figure 2 : analyse de calcul matriciel vis-à-vis de l’accès mémoire
Sur une ligne horizontale sur la figure 2, que remarquez-vous ?
Figure 3 : analyse de calcul matriciel montre un chargement multiple des valeurs
On remarque que chacune des valeurs sont systématiquement chargées deux fois !
Les threads chargent plusieurs fois la même valeur depuis la mémoire globale, ce qui provoque une surcharge substantielle en termes de trafic avec la mémoire globale, ce qui ralentit forcément notre traitement. Comment éviter ce problème ? Naturellement, l’idée serait d’utiliser la mémoire partagée qui est bien plus rapide, mais malheureusement très petite.
Imagions que toutes les données soient déjà présentes en mémoire partagée, les quatre threads pourraient charger plusieurs fois les mêmes valeurs sans problème, car le coût d’accès à la mémoire partagée est très rapide (10 fois moins couteux).
Si nous résumons nos idées :
-
La mémoire partagée est réduite en taille, il va donc falloir fractionner nos données en petit rectangles de données adaptés à la capacité de la mémoire partagée.
-
Pour chaque rectangle, nous devons recopier ses données dans la mémoire partagée.
-
Une fois que toutes les données du rectangle sont recopiées, nous pouvons traiter notre calcul réparti sur tous les threads GPU correspondants à notre rectangle qui puiseront leurs données depuis la mémoire partagée.
-
Les calculs intermédiaires d’un rectangle doivent être placés dans une variable temporaire. Au final, la variable temporaire est recopiée dans la cellule correspondante de la matrice C.
Il est usuel en programmation GPU de comparer ce découpage en rectangles au recouvrementd’une surface avec des tuiles. L’élément clef de cette répartition est de permettre aux calculs de s’exécuter sans surcout. Nous pouvons remarquer que l’usage de la mémoire partagée est conditionné par l’algorithme. Ici le produit matriciel utilise plusieurs fois les mêmes valeurs au cours du traitement c’est donc un excellent candidat pour utiliser la mémoire partagée.
Technique de découpage en tuiles n’est pas toujours la bonne solution
Il n’est pas toujours possible d’utiliser la technique des tuiles pour exploiter la mémoire partagée. En effet, si votre algorithme ne comporte d’utilisation multiple des mêmes données, il n’y a pas de gain de performance possible dans cette direction.
Par exemple dans le cadre d’une simple addition de vecteurs, cela n’aurait aucun sens, car les données ne sont jamais utilisées plusieurs fois durant le calcul. Mais dans le cas d’addition de deux vecteurs, une autre technologie basée sur les techniques vectorisation comme SSE (Streaming SIMD Extensions) ou dernièrement AVX (Advanced Vector Extensions), serait sans doute une meilleure piste sur le plan des performances.
Lorsquevous décidez de paralléliser un morceau de code, vous devez rester vigilant surl’algorithme afin de déterminer comment le paralléliser efficacement.
Révision de l’algorithme du code noyau
Après avoir découvertl’organisation de la mémoire, vous vous demandez peut-être comment exprimer qu’une variable est définie dans la mémoire partagée. Le compilateur CUDA C expose un mot clef __shared__ permettant de définir de la mémoire partagée sur des variables définies dans la pile. Dans notre réflexion, nous considérons que les traitements sont marqués en étapes : la copie des données dans la mémoire partagée, puis le calcule matriciel et enfin la copie de la valeur dans la matrice résultat. Pour s’assurer que toutes les étapes sont bien orchestrées, nous utiliserons une nouvelle fonction CUDA C, __syncthreads() . La fonction __synchthread() permet d’attendre que tous les threads au sein d’une tuile aient terminé leurs accèsmémoire.
Comment définir la taille de vos tuiles ?
Nous n’avons pas expliqué comment définir la taille des tuiles (TILE_WIDTH dans le code). La réponse est simple, vous devez expérimenter vous-même les performances de votre programme en fonction des tailles utilisées. En général pour le produit matriciel et vis-à-vis de la génération des cartes actuelles, une taille de tuile de 16 x 16 est souvent optimum. Mais ce chiffre pourra évoluer en fonction des matériels.
Performances de calcul matriciel en fonction de la taille des tuiles sur une carte NVIDIA Quadro 2000M.
La ligne « Performance »est exprimée en milliseconde.
Taille 2 4 8 16
Performance 3 464 523 92 52
On remarque que la taille à une conséquence importante vis-à-vis des performances du programme.
Je vous encourage à tester les performances de vos traitements reposant sur la technique des tuiles afin d’évaluer la taille adéquate des tuiles vis-à-vis de votre programme.
Le code ci-dessous (Fragment de code 2) illustre nos idées précédentes.
__global__
void MatrixMultiplySquareCachedTiled(float* A, float* B, float* C, int size)
{
__shared__ float As[TILE_WIDTH][TILE_WIDTH];
__shared__ float Bs[TILE_WIDTH][TILE_WIDTH];
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
int i = by * TILE_WIDTH + ty;
int j = bx *TILE_WIDTH + tx;
float sum = 0;
for (int m = 0; m < size/TILE_WIDTH; ++m)
{
// Phase 1
As[ty][tx] = A[i* size + (m * TILE_WIDTH + tx)];
Bs[ty][tx] = B[(m * TILE_WIDTH + ty) * size + j];
__syncthreads();
// Phase 2
for (int k = 0; k < TILE_WIDTH; ++k) {
sum += As[ty][k] * Bs[k][tx];
}
__syncthreads();
// Phase 3
C[i * size + j] = sum;
}
}
Fragment de code 2 : code du calcul matricieltuilé
On retrouve le fractionnement des données et les trois phases dans le traitement d’une tuile :
1. Le chargement des données depuis la mémoire globale vers la mémoire partagée.
2. Calcul matriciel pour la tuile courante.
3. Recopie de la variable temporaire dans la cellule correspondante à matrice C.
Exécution du produit matriciel pour chaque mode d’exécution
En programmation parallèle la performance est sans aucun doute un des buts premier. Une comparaison avec toutes formes de calculs matriciels étudiées dans cette série d’articles permet de nous montrer les gains obtenus pour chaque mode d’exécution.
Figure 4- Exécution de tous nos calculs matriciels exprimés en millisecondes
Pour éviter d’avoir un effet réducteur, nous ne comparons ici que les implémentations CUDA.
Figure 5- Comparatif CUDA avec et sans cache exprimé en millisecondes
Cette nouvelle implémentation nous apporte un gain de 38% de performance au prix d’une complexité substantielle, mais qui ne devrait pas rebuter le passionné de performance qui caractérise les programmeurs GPU en général.
En conclusion
La programmation CUDA C exige une bonne compréhension de l’architecture CUDA. En effet, les capacités de votre carte graphique seront toujours à prendre en compte vis-à-vis de votre appétit à accélérer vos calculs. La maitrise du nombre de blocs et du nombre de threads, des indices et du nombre de tuiles de votre algorithme en mode noyau constitue un élément essentiel pour réussir d’excellents programmes CUDA. Ces calculs peuvent sembler compliqués pour le néophyte en programmation GPU, mais une pratique régulière de ces calculs vous permettra de ne plus être pénalisés par cette petite gymnastique intellectuelle. Sur le plan de la programmation, le Framework CUDA C impose une bonne organisation du code afin de ne pas sombrer dans des mélanges douteux entre le code CPU et le code GPU. Si votre algorithme s’y prête, l’utilisation de la mémoire cache permetd’obtenir des performances extrêmes, mais au prix d’un code un peu plus complexe. Il est sans doute déconseillé de se lancer dans la programmation CUDA C trop rapidement. Ce type de programmation reste un peu rugueux, mais les performances sont littéralement exceptionnelles lorsque vos algorithmes sont orientés données (traitements orientés boucles). Enfin, j’aimerais ajouter que la démarche et les principes décrits dans cet article sont communs à toutes les technologies GPU. Par exemple Open CL et C++ AMP respectent ces principes.
Ressources
Si vous souhaitez aller plus loin sur le développement CUDA, je vous conseille les ouvrages suivants. Le premier est sans doute le plus didactique pour une première approche. Le second est beaucoup plus complet et offre des explications très détaillées, c’est donc un ouvrage de référence.
A bientôt
Bruno