FLASH INFORMATIQUE FI



Les GPU ne sont pas uniquement faits pour les consoles de jeux




Francis LAPIQUE


GPU vs CPU : to be, or not to be

JPEG - 9.9 ko
fig. 1

La figure 1 présente un comparatif, en performance par Watt, entre Nvidia et Intel qui s’affrontent par déclarations interposées quant à l’avenir respectif du GPU et du CPU.
Nvidia a ainsi récemment déclaré que les processeurs n’avaient plus aucun avenir et qu’ils seraient très prochainement supplantés par les puces graphiques !
Intel, répond en déclarant que Cuda ne sera qu’une intéressante note de bas de page dans les annales de l’histoire informatique et d’expliquer que malgré les gains exceptionnels promis par la technologie Cuda, un problème de taille subsiste. Les gains de performances sont bien présents, mais peu de développeurs sont en mesure de l’exploiter pleinement.
Sur ce dernier point, l’année dernière un cours pour les architectures multi-core sur la base de la bibliothèque Intel Threading Building Blocks (TBB) était annoncé en posant la question : que va-t-on en tirer au niveau applicatif ? Je dois avouer que les résultats peu encourageants, de mon point de vue, avec des double ou quadri coeurs m’ont décidé à reporter ce cours. Je ne sais pas si Cuda ne sera qu’une intéressante note de bas de page dans les annales de l’histoire informatique, mais aujourd’hui, nous le verrons dans la suite de cet article, cette technologie nous offre des gains de performance 50x à 100x pour de faibles coûts de développement en terme de temps et de prix.

Architecture


JPEG - 8.8 ko
fig. 2
la carte NVIDIA GeForce GTX 280

L’architecture des cartes GeForce compte un ensemble de multiprocesseurs (ou SM pour Streaming Multiprocessor) au nombre de 30 pour la carte GeForce GTX280 (voir fig. 2). Chacun de ces SM est équipé de 8 ALUs (Arithmetic Logic Unit).
Une unité SIMT (Single-Instruction Multiple-Thread) prend en charge la création, l’ordonnancement et l’exécution de warps (groupe de 32 threads, le terme vient des machines à tisser, il désigne un ensemble de fils de coton).
Un thread CUDA n’a pas tout à fait le même sens qu’un thread CPU, c’est un élément de base des données à traiter. A l’inverse des threads CPU, les threads CUDA sont extrêmement légers ce qui signifie qu’un changement de contexte est une opération peu coûteuse. Un warp est exécuté tous les 2 cycles ; par exemple, pendant les deux premiers cycles, warp 8 exécute l’instruction 11, pendant les deux cycles suivants, wrap 1 exécute l’instruction 42. Chaque multiprocesseur peut prendre en compte 32 warps soit 1024 threads.

Modèle de programmation

Au niveau programmation on ne manipule pas directement ces warps, on écrit des noyaux de programmation ou kernels.
Ces kernels se subdivisent en blocs composés de warps. Chaque SM dispose d’une mémoire partagée ou Shared Memory de 16’384 bytes et de 16’384 registres de 32 bits (fig. 3). La mémoire partagée n’est pas une mémoire cache, c’est un espace ouvert à la programmation.

JPEG - 12.7 ko
fig. 3
Architecture SM


GIF - 5 ko
fig. 4
Modèle de programmation

L’environnement de développement CUDA compte quelques extensions au langage C, des librairies, un compilateur (nvcc) et un pilote (fig. 5). Sans trop rentrer dans le détail de ces extensions on trouve des qualificateurs comme __global__ pour désigner un kernel, __device__ pour une fonction exécutée et appelée depuis le GPU, il désigne également une variable dans l’espace Global memory à savoir la DRAM du GPU (1Go pour la carte GTX280), __shared__ pour une variable Share memory.

JPEG - 8.8 ko
fig. 5
Architecture CUDA

Un kernel :

__global__ void Func(float* parameter)

est appelé d’une façon un peu particulière ;

Func<<< Dg, Db, Ns >>>(parameter);

Dg représente la taille de la grille en nombre de blocs, Db la taille du bloc en nombre de threads et Ns un espace mémoire optionnel dynamiquement alloué par bloc dans la mémoire partagée et entre parenthèses la liste des paramètres.
Un jeu de variables intégrées ou built-in permet d’identifier un thread. La ligne de commande qui suit est un exemple de calcul d’index de thread à travers ces variables built-in :

int idx = blockIdx.x * blockDim.x + threadIdx.x;

blockIdx contient l’index du bloc dans la grille, threadIdx celui du thread dans le bloc et blockDim le nombre de thread par bloc.
La fonction __synchthreads() permet de synchroniser les threads à l’intérieur d’un bloc.
Pour résumer les limitations matérielles de la carte GTX280 qui nous sont données par le programme device Query :

Total amount of global memory         1'073'479'680 bytes
Number of multiprocessors         30
Number of cores        240
Total amount of constant memory        65'536 bytes
Total amount of shared memory per block        16'384 bytes
Total number of registers available per block        16384
Warp size        32
Maximum number of threads per block        512
Maximum sizes of each dimension of a block        512 x 512 x 64
Maximum sizes of each dimension of a grid        65'535 x 65'535 x 1
Maximum memory pitch        262'144 bytes

Le développement de votre application au niveau kernel (grille et bloc) ne vous affranchit pas de penser aux niveaux multiprocesseurs d’exécution de warps et au niveau ALU (Unité Arithmétique et Logique, sous-partie d’un processeur capable d’effectuer des opérations définies par un jeux d’instructions) d’identifiant de bloc et de thread. Deux points très importants qui concernent l’optimisation : le conflit de banques de la mémoire partagée et la coalescence de la mémoire globale. Vous pouvez approfondir ces deux notions dans le Programming Guide CUDA [1].

Shared memory et conflit de banques

JPEG - 11.7 ko
fig. 6
Conflit de banques

JPEG - 28.3 ko

L’espace de communication Shared memory, espace de 16 Ko, est organisé en 16 banques (fig. 7).

GIF - 9.2 ko
fig. 7

La figure 6.a présente un exemple de conflit : les thread 0 et 8 accèdent à la même case 0 de la banque. C’est quelque chose qu’il faut éviter, car dans ce cas l’accès est sérialisé avec une chute de la bande passante (fig. 6.b). En cas de non-conflit, l’accès est presque aussi rapide que les registres contrairement aux 400 cycles d’horloge de la mémoire globale.

Global memory et coalescence

La mémoire globale souffre d’une latence importante (400 à 600 cycles d’horloge pour un accès). Pour remédier à ce problème, CUDA donne la possibilité d’accéder à un bloc de plusieurs cases mémoires ou coalescence. La coalescence est réalisée au niveau du demi-wrap (un cycle d’horloge du SIMT soit 16 threads) si on accède à des régions mémoires de 32, 64 ou 128 octets.
On obtient également des gains de vitesse si l’on accède à des cases mémoires voisines et dans l’ordre des indices de threads. Les figures 6a et 6c illustrent ces propos.
Il existe deux types d’accès non coalescents. Le premier est dû au fait que les threads n’accèdent pas dans l’ordre à des cases voisines (fig. 6b) ; le second est dû à un problème d’alignement (fig. 6c). Le premier thread d’un warp doit accéder à une case mémoire multiple de 64.

JPEG - 10 ko
fig. 6b - 6c
coalescence

Une dernière remarque lorsque différents threads d’un même wrap travaillent dans deux branches différentes d’un même if, l’impact sur les performances peut être important, car le GPU est obligé d’exécuter le warp sur plusieurs cycles d’horloge.

Hello World !

Pour découvrir le monde CUDA voici un exemple de somme de deux vecteurs.

__global__ void sum(float* A, float *B,
float* C, int width)
{
....
3 exécution du noyau dans le GPU
.....
unsigned int idx = threadIdx.y * width + threadIdx.x;
C[idx] = A[idx] + B[idx];
}
int main()
{
...
Allocation des vecteurs de données à traiter par le GPU
....
unsigned int mem_size=10*10*sizeof(float);
float *a; cudaMalloc((void**)&a, mem_size);
float *b; cudaMalloc((void**)&b, mem_size);
float *c; cudaMalloc((void**)&c, mem_size);
float input_a[100]; // size = 10*10
float input_b[100];
float output_c[100];
...
1- Envoi des données vers le GPU
...
cudaMemcpy(a,input_a, mem_size, cudaMemcpyHostToDevice);
cudaMemcpy(b,input_b, mem_size), cudaMemcpyHostToDevice);
dim3 dimBlock(10, 10);

...
2 - demande d'exécution du noyau dans le GPU
...
sum<<<1, dimBlock>>>(A, B, C);

...
4 - retour du résultat vers le CPU
....
cudaMemcpy(output_c, C, mem_size,cudaMemcpyDeviceToHost);
...
}

Les numéros indiqués dans le code renvoient à ceux de la figure 8.

JPEG - 11.4 ko
fig. 8
CUDA processing flow - auteur : Tosaka

Smith-Waternam avec CUDA

Revenons à notre titre Les GPU ne sont pas uniquement faits pour les consoles de jeux. La technologie NVIDIA CUDA veut montrer au marché HPC (Calcul Haute Performance) qu’elle peut jouer un rôle important dans les années qui viennent. Il est vrai que son site présente une palette d’applications (fig. 9) qui vont bien au-delà de son métier de niche [2].
Pour nous faire notre propre opinion, nous avons acquis deux cartes GTX280 et décidé de porter un algorithme de recherche heuristique utilisée en bio-informatique permettant de trouver les régions similaires entre deux ou plusieurs séquences de nucléotides ou d’acides aminés.
La recherche d’alignements entre séquences génomiques est une des tâches fondamentales de la bio-informatique. L’objectif est de localiser des régions semblables dans des séquences d’ADN ou des séquences protéiques. Une application typique est l’interrogation d’une banque avec un gène dont la fonction est inconnue. Les résultats retournés correspondent à des segments similaires présentant un indice de ressemblance élevé. Plus exactement, l’information utile est un alignement, c’est-à-dire deux portions de séquence où sont précisément indiqués les appariements entre nucléotides (pour l’ADN) ou les appariements entre acides aminés (pour les protéines).
On doit parcourir systématiquement l’ensemble des banques, de la première à la dernière séquence. Il existe plusieurs algorithmes pour extraire des alignements. Les premiers, comme celui de Smith-Waterman élaboré en 1981 [3], utilisent des techniques de programmation dynamique et possèdent une complexité quadratique. Les seconds, apparus en 1990, comme le programme BLAST, se basent sur une heuristique très efficace (recherche de points d’ancrage) permettant de cibler directement de courtes zones identiques potentiellement intéressantes.
Nous avons porté notre choix sur celui de Smith-Waterman qui a fait l’objet ces derniers temps de plusieurs publications concernant son implémentation sur GPU [4].

JPEG - 16.3 ko
fig. 9

L’algorithme de Smith-Waterman

L’algorithme de Smith-Waterman évalue une ressemblance locale entre deux séquences A et B de taille m et n respectivement. Les éléments de A et B sont notés a1, … am et b1, … bn. La matrice de similitude entre éléments W(ai, bj) et les pénalités d’ouverture et d’extension de gap Ginit et Gext sont également données. On calcule les n x m matrices suivantes :



La procédure qui permet de trouver l’alignement à partir de la matrice est la suivante : à partir de la cellule d’arrivée, remonter vers la(les) cellule(s) voisine(s) de score maximal ; itérer pour arriver à la cellule initiale.
Pour illustrer le propos, la figure 10 donne un exemple concret de la procédure de recherche locale entre deux séquences ADN avec une expression simplifiée de la matrice :
H(i,j) = max0, H(i,j-1)-α, H(i-1,j)-α, H(i-1,j-1) + W(S1i ,S2j ) , avec α=1, W(S1i ,S2j ) =+2
si les résidus i et j sont identiques et W(S1i,S2j)=-1 dans le cas contraire.

JPEG - 7.2 ko
fig. 10

Vous trouverez le pseudo-code d’implémentation à al référence [2]. Une des particularités de ce code est de charger dans la mémoire de texture une matrice acides aminés -requête dite Query-profile.
En partant du travail de ces deux auteurs, cet algorithme a été porté sur une machine de bureau (Intel 3.0 Ghz, cache 6MB, express Chipset. 4 GB DDR3, 2x Nvidia GTX 280, system Linux kernel 2.6.27).
Vous pouvez faire des tests en-ligne depuis la page gpu.epfl.ch/sw.html. Pour lancer la requête cliquer Query, car cette page contient par défaut un jeu de paramètres de test. La règle d’attribution du service est des plus simples premier arrivé, premier servi. La base de données (Swiss-Prot) compte un peu plus de 400 mille séquences et environ 145 millions de résidus. La réponse doit parvenir au bout de quelques secondes.

JPEG - 6.6 ko
fig. 11

À la lumière des résultats présentés à la figure 11, évitez le mode CPU pour des requêtes plus grandes qu’une dizaine de résidus, car n’oubliez pas que vous multipliez le temps par un facteur d’environ 50.

MATLAB

Les figures 12.a et 12.b donnent une idée du facteur d’accélération que l’on peut obtenir avec l’extension CUDA pour MATLAB que l’on peut obtenir depuis le site Nvidia [5]. Vous trouverez dans cette extension un script nvmex à l’image du script mex pour compiler l’application (code 2.0) contenant du code CUDA.

#include "cuda.h"
#include "mex.h"
/* Kernel to square elements of the array on the GPU */
__global__ void square_elements(float* in, float* out, int N)
{
int idx = blockIdx.x*blockDim.x+threadIdx.x;
if (idx < N) out[idx]=in[idx]*in[idx];
}
/* Gateway function */
void mexFunction(int nlhs, mxArray *plhs[], int nrhs, const mxArray *prhs[])
{

/* Call function on GPU */
square_elements<<<dimGrid,dimBlock>>>(data1f_gpu, data2f_gpu, n*m);

}



nvmex -f nvopts.sh square_me.cu -I/usr/local/cuda/include -L/usr/local/cuda/lib -lcufft -Wl,-rpath,-L/usr/local/cuda/lib -L../../common/lib/linux -lcudart -L../../common/lib/linux -lcutil
JPEG - 18.8 ko
fig. 12a
Advection of an elliptic vortex on a 256x256 mesh, stream function (left), vorticity (right) : MATLAB on Linux, 72.460358 seconds

JPEG - 18.6 ko
fig. 12 b
Advection of an elliptic vortex on a 256x256 mesh, stream function (left), vorticity (right):MATLAB with CUDA on Linux, 5.873686 seconds

À la conférence SC08 (voir article de Pascal Jermini et Jean-Claude Leballeur paru dans le FI1/9) qui a eu lieu à Austin en novembre dernier Wolfram Research a annoncé pour le premier trimestre de cette année une version CUDA accélérée pour Mathematica.

Conclusion

Ce premier contact avec CUDA est des plus encourageants, nous allons entreprendre d’autres tests critiques dans le monde du 64 bits flottant. Nvidia semble actuellement avoir une avance sur ses concurrents, mais la donne peut vite changer. Si vous voulez faire connaître ou partager vos expériences GPU vous pouvez me faire parvenir un lien pour le site gpu.epfl.ch. Vous pouvez également vous inscrire à un cours de sensibilisation à la programmation CUDA sur deux jours, les 28 et 29 avril prochain.

Glossaire

Core : désigne l’équivalent d’un processeur généraliste semi-autonome qui partage un certain nombre de composants avec les autres cores.

Streaming Multiprocessor : grappe de processeurs ALU partageant une unité de contrôle et des mémoires

SIMD : Single instruction Multiple Data, type de processeur vectoriel utilisant une même instruction sur plusieurs données différentes

banques : désigne un segment sur une mémoire

threads : élément de calcul s’exécutant sur le processeur scalaire

kernel : fonctions appliquées à chaque ensemble de données nécessitant l’application du même traitement
warp : groupement de 32 threads

grille, blocs : le modèle de programmation considère une grille composée d’un ensemble de blocs qui exécutent simultanément des threads

BLAST (acronyme de basic local alignment search tool) est un algorithme utilisé en bio-informatique permettant de trouver les régions similaires entre deux ou plusieurs séquences de nucléotides ou d’acides aminés.

[1] Programming Guide CUDA

[2] Hervé Jourdren, Marc Pérache, Clément Koyesse. Retour d’expérience GPU cartes graphiques programmables en calcul haute performance
Michael Boyer, Kevin Skadron, Westley Weimer, Automated Dynamic Analysis of CUDA Programs

[3] Smith TF, Waterman MS (1981). Identification of Common Molecular Subsequences. Journal of Molecular Biology 147 : 195-197. doi:10.1016/0022-2836(81)90087-5.

[4] Manavski SA, Valle G (2008). [CUDA compatible GPU cards as efficient hardware accelerators for Smith-Waterman sequence alignment. BMC Bioinformatics 9 (Suppl 2:S10) : S10. doi:10.1186/1471-2105-9-S2-S10. www.biomedcentral.com/1471-2....
Rognes T, Seeberg E (2000). Six-fold speed-up of Smith-Waterman sequence database searches using parallel processing on common microprocessors. Bioinformatics 16 : 699-706. bioinformatics.oxfordjournals.org/cgi/reprint/16/8/699.pdf.
Farrar M S (2008). Optimizing Smith-Waterman for the Cell Broadband Engine. farrar.michael.googlepages.com/smith-watermanfortheibmcellbe.
Van Hoa NGUYEN, Dominique LAVENIER, Parallélisation de la recherche de similarités entre séquences protéiques sur GPU, Rencontres francophones du Parallélisme, Fribourg : Switzerland (2008)

[5] Massimiliano Fatica, NVIDIA, Won-Ki Jeong, University of Utah. Accelerating MATLAB with CUDA.



Cherchez ...

- dans tous les Flash informatique
(entre 1986 et 2001: seulement sur les titres et auteurs)
- par mot-clé

Avertissement

Cette page est un article d'une publication de l'EPFL.
Le contenu et certains liens ne sont peut-être plus d'actualité.

Responsabilité

Les articles n'engagent que leurs auteurs, sauf ceux qui concernent de façon évidente des prestations officielles (sous la responsabilité du DIT ou d'autres entités). Toute reproduction, même partielle, n'est autorisée qu'avec l'accord de la rédaction et des auteurs.


Archives sur clé USB

Le Flash informatique ne paraîtra plus. Le dernier numéro est daté de décembre 2013.

Taguage des articles

Depuis 2010, pour aider le lecteur, les articles sont taggués:
  •   tout public
    que vous soyiez utilisateur occasionnel du PC familial, ou bien simplement propriétaire d'un iPhone, lisez l'article marqué tout public, vous y apprendrez plein de choses qui vous permettront de mieux appréhender ces technologies qui envahissent votre quotidien
  •   public averti
    l'article parle de concepts techniques, mais à la portée de toute personne intéressée par les dessous des nouvelles technologies
  •   expert
    le sujet abordé n'intéresse que peu de lecteurs, mais ceux-là seront ravis d'approfondir un thème, d'en savoir plus sur un nouveau langage.