GPU - Modèle de programmation

Dans cet article on verra comment est fait le modèle de programmation GPU. C’est-à-dire comment on programme dessus et comment le comprendre grossièrement. Le GPU peut être considérée comme une extension du CPU qui permet de l’aider dans ses calculs. Le CPU envoie les instructions au GPU et attend ses résultats sans pour autant être une action bloquante. Aujourd’hui il y a plusieurs choses qui font que les CPUs ne sont plus aussi évolutifs qu’avant : l’augmentation de la fréquence entrainait forcément une réduction du voltage pour une consommation modérée, mais aujourd’hui le voltage ne peut plus descendre plus bas : Puissance (W) = Fréquence x Voltage². Et si aujourd’hui on descendait encore le voltage, la distinction des entrées binaires ne se feraient plus correctement : il serait complexe de faire la différence entre un 0 et un 1. Du coup, au lieu d’augmenter la fréquence, nous avons exploité le principe de parallélisme, cela veut dire que nous avons besoin d’un composant qui permette de faire plusieurs opérations en même temps, avec un coup en puissance la plus faible possible : le GPU. Par conséquent le CPU et le GPU sont techniquement proche et loin en même temps : la fréquence du GPU est inférieure à celle du CPU, le GPU utilise de la mémoire GDDR ou HBM2 alors que le CPU utilise du DDR, la latence du GPU est camouflée par son haut niveau de multithreading, le CPU peut partager sa mémoire avec un autre CPU, contrairement aux GPUs, et comme nous l’avons dit dans l’introduction de ce chapitre, les GPU sont programmables, ils sont Turing-Complete, une autre différence se situe au niveau des entrées et sorties, en effet pour une puce qui est performante sur des opérations parallèle, le fait de faire une seule opération sera longue et fastidieuse, d’autant plus que la mémoire du GPU n’est pas partageable, l’opération n’en sera que plus longue, et dans l’idée, si nous voulons paralléliser cette opération, cela semble impossible : mettre une centaine de threads qui vont rentrer en collision sur une seule requête, les GPU sont aussi architecturalement plus simple que les CPU, ils n’ont pas de prédiction de branche, pas d’exécution out-of-order (dans le désordre). Mais il ne faut pas oublier que les CPU ne peuvent pas être remplacés par le GPU, le GPU reste une puce d’accélération.

Une application dédiée au GPU commence à s’exécuter sur le CPU. Quand on avait des GPUs qui partageaient le même cache que le CPU les 2 unités se partageaient la mémoire, c’était un GPU intégré, aujourd’hui nous avons des GPU discrets : ils ont leur mémoire GDDR, et c’est grâce à l’introduction de l’architecture Pascal de Nvidia qui implémente un support hardware et software pour automatiquement transférer les données du CPU vers le GPU grâce à la mémoire virtuelle que Nvidia nomme « la mémoire unifiée ».

Avec l’aide d’un driver qui s’exécute sur le CPU, le code de l’application GPU (aussi appelé « kernel »/noyau) défini quelles portions de code devront s’exécuter sur le GPU, le nombre de threads qui vont être lancés et où sont les données qui vont devoir être utilisées. Donc le driver devra faire beaucoup d’opérations pour donner au GPU les informations nécessaires pour exécuter le kernel. Les cœurs Nvidia sont appelés « Streaming Multiprocessors » et les cœurs AMD « Compute Units ». Chaque cœur GPU exécute un programme SIMT (Single Instruction, Multiple Thread), par exemple une opération d’addition sur un millier de données, le kernel. Chaque cœur GPU peut exécuter environ 1000 threads.

Selon la Taxonomie de Flynn, les GPUs modernes sont basés sur un système SIMD (Single Instruction, Multiple Data), ce qui leur permet d’exploiter le parallélisme au niveau des données. Mais au lieu d’exposer un système SIMD au programmeur, les APIs, comme CUDA et OpenCL, offrent des options qui s’apparente à un modèle de programmation MIMD (Multiple Instructions, Multiple Data) afin d’exécuter un grand nombre de threads scalaire sur le GPU.

Chacun de ces threads scalaires peut suivre son propre chemin d’exécution et peut accéder arbitrairement à des emplacements mémoire. A l’exécution, un groupe de threads scalaire sont nommés par Nvidia et AMD respectivement warps et wavefronts, et ces warps sont exécutés en même temps sur le système SIMD, c’est ce qu’on appelle donc le SIMT.

Chaque kernel est composé de milliers de threads, et chaque thread exécute le même programme, mais des threads vont suivre un flow de contrôle différent suivant les dépendances du programme. Un exemple basique que nous pouvons prendre est SAXPY (Single precision scalar value A times vector value X Plus vector value Y) (grossièrement A * X + Y) faisant parti de la bibliothèque BLAS (Basic Linear Algebra Subprograms). Voici un exemple de code pouvant s’exécuter sur le CPU :

void 
saxpy (int n, 
       float a, 
       float * x, 
       float * y)
{
    for (int i = 0; i < n; ++i) 
    {
        y[i] = a * x[i] + y[i];
    }
}

void
main (void) 
{
	float * x, * y;
	int n;
	// malloc x et y et mettre les données dedans
	saxpy(n, 2.0, x, y) ;
	// free x et y
}

Sur un code dédié au GPU, prenons par exemple l’API CUDA de Nvidia, le code sera totalement différent :

__global__ void 
saxpy (int n,
       float a,
       float * x,
       float * y) 
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < n)
		y[i] = a * x[i] + y[i];
}

void
main (void) 
{
	float * h_x, * h_y; // host memory pointer
	int n;
	// malloc h_x et h_y et initialiser les données
	float * d_x, * d_y ; // device memory pointer
	int nblocks = (n + 255) / 256 ;
	cudaMalloc(&d_x, n * sizeof(float));
	cudaMalloc(&d_y, n * sizeof(float));
	cudaMemcpy(d_x, h_x, n * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(d_y, h_y, n * sizeof(float), cudaMemcpyHostToDevice);
	saxpy<<<nblocks, 256>>>(n, 2.0, d_x, d_y);
	cudaMemcpy(h_x, d_x, n * sizeof(float), cudaMemcpyDeviceToHost);
	// free h_x, h_y, d_x, d_y
}

Ce qu’il faut comprendre ici c’est que on divise notre tableau en chunks de taille 256 appelés thread-blocks, et chaque thread-bloc est exécuté sur un cœur GPU, un SM (Streaming Multiprocessor). Un groupe de thread-blocks est appelée une grid de thread-blocks. Lors de la référence à saxpy<<<nblocks, 256>>>, nblocks représente la taille de la grid (qui contient la collection des threads-blocks), et 256 qui représente le nombre de threads par thread-block. Ici saxpy est un kernel, donc son code sera exécuté tel quel dans le GPU. Il y a aussi plusieurs fonctionnalités dans la programmation GPU/parallèle, comme le principe de mémoire partagée, plus rapide que la mémoire locale et globale (stockées soit sur le CPU dans le cache, soit dans la RAM (Random Memory Access)), la mémoire partagée est stockée directement sur le GPU, chaque thread-block alloué ont accès à la même mémoire partagée et il n’y a pas de problème d’accès concurrent, ceci s’explique par un accès sérialisé à la mémoire : les requêtes concurrentes sont séparées en plusieurs requêtes dénoué de conflits, donc évidemment ceci réduit l’efficacité de la bande passante GPU, la seule exception est lorsque tous les threads d’un warp ont la même adresse de mémoire partagée, ce qui résulte à un broadcast : ceci est très utile lorsque l’on parle de synchronisation dans les threads. La synchronisation de threads est supportée par des unités hardware qui font barrière à l’exécution d’instructions.

Le GPU, comme le CPU, contient son propre ISA (Instruction Set Architecture). Depuis l’introduction de CUDA, l’ISA est haut niveau/virtuel, aussi appelée Parallel Thead Execution ISA (PTX), le PTX est assez similaire aux instructions d’une architecture processeur RISC (donc comme ARM, MIPS, SPARC, etc). Par exemple le PTX du code CUDA au dessus ressemble à ceci:

.visible .entry _Z5saxpyifPfS_(
.param .u32 _Z5saxpyifPfS__param_0,
.param .f32 _Z5saxpyifPfS__param_1,
.param .u64 _Z5saxpyifPfS__param_2,
.param .u64 _Z5saxpyifPfS__param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<5>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;


ld.param.u32 %r2, [_Z5saxpyifPfS__param_0];
ld.param.f32 %f1, [_Z5saxpyifPfS__param_1];
ld.param.u64 %rd1, [_Z5saxpyifPfS__param_2];
ld.param.u64 %rd2, [_Z5saxpyifPfS__param_3];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r4, %r3, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra BB0_2;

cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f2, [%rd6];
add.s64 %rd7, %rd3, %rd5;
ld.global.f32 %f3, [%rd7];
fma.rn.f32 %f4, %f2, %f1, %f3;
st.global.f32 [%rd7], %f4;

BB0_2:
ret;
}

Donc comme tout code haut niveau, il faut le compiler (dans notre cas) pour avoir l’ISA compatible avec notre GPU, donc on transforme notre code haut niveau en langage assembleur/intermédiaire RISC, appelé SASS (Streaming ASSembler), ce process est effectué par un programme nommé ptxas pour les GPU Nvidia qui introduisent CUDA. Voici le code SASS du code CUDA précédent:

Alors le SASS est peu documenté, c’est grâce à la communauté informatique qui a développé des outils comme le projet decuda que de la documentation partielle est arrivée, Nvidia a introduit l’outil cuobjdump pour la documentation du SASS. Aujourd’hui cette documentation correspond juste à la liste des opcodes, mais aucun détail n’est fourni sur la sémantique, ce qui rend le travail de rétro-ingénierie assez complexe pour comprendre comment Nvidia optimise le code. Le code PTX et le code SASS sont très différent, sachant que pour chaque nouvelle architecture (pour Nvidia : Tesla, Fermi, Pascal, …) le SASS est différent. Contrairement à Nvidia, AMD fournit une documentation complète quant à leur ISA, et cela a aidé les chercheurs académiques pour leurs travaux de simulation bas-niveau. Le flow de compilation AMD fournit aussi un ISA virtuelle, appelée HSAIL (Heterogeneous System Architecture Intermediate Language). Une différence notoire entre les architectures Nvidia et AMD est la séparation des instructions scalaires et vectorielles pour AMD : les instructions scalaires commencent par le préfix s_, et les vectorielles par v_.

Prérequis

Hello ! o/