Coder pour le GPU : threads,
blocks, grids
Kernels, threads, warps, grilles de thread-blocks, mémoire unifiée : tout le modèle de programmation GPU illustré via SAXPY en CUDA. Mise en regard de PTX et SASS pour l'architecture NVIDIA.
Maxime J10 min de lecture
Un doute sur votre exposition ?
On audite votre surface d'attaque externe et on vous dit ce qui est exploitable.
Parler à un pentesterPrérequis
- Comprendre le fonctionnement d'un CPU
- Unités de calcul (ALU, IFU, SFU, FPU, ...)
- Pipelining et exécution spéculative
- Caches
- Introduction aux GPUs
Hello ! o/
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é 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 que ça soit 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. La relation est simple : Puissance (W) = Fréquence x Voltage². Et si aujourd'hui on descendait encore le voltage, la distinction des entrées binaires ne se ferait 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, on a exploité le principe de parallélisme. On a besoin d'un composant qui permette de faire plusieurs opérations en même temps avec un coût en puissance le plus faible possible : le GPU.
Par conséquent le CPU et le GPU sont techniquement proches 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.
- Les GPUs sont programmables, ils sont Turing-complete.
- Pour une puce performante sur des opérations parallèles, faire une seule opération sera lente, d'autant plus que la mémoire GPU n'est pas partageable.
- Les GPUs sont architecturalement plus simples que les CPUs : pas de prédiction de branche, pas d'exécution out-of-order.
Le GPU, une puce d'accélération
Mémoire unifiée et flow d'exécution
Une application dédiée au GPU commence à s'exécuter sur le CPU. Quand les GPUs partageaient le même cache que le CPU, les deux unités se partageaient la mémoire : c'était un GPU intégré. Aujourd'hui on a des GPU discrets avec leur propre mémoire GDDR. C'est l'architecture Pascal de Nvidia qui a introduit le support hardware et software pour transférer automatiquement les données CPU vers GPU via la mémoire virtuelle que Nvidia appelle 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 ou noyau) définit quelles portions de code devront s'exécuter sur le GPU, le nombre de threads à lancer et où sont les données à utiliser. Le driver fait donc beaucoup d'opérations pour donner au GPU toutes les informations nécessaires.
Les coeurs Nvidia sont appelés Streaming Multiprocessors (SM) et les coeurs AMD Compute Units. Chaque coeur 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 coeur GPU peut exécuter environ 1 000 threads.
Source : Wikipédia
SIMD, MIMD et SIMT
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 brut au programmeur, les APIs comme CUDA (Compute Unified Device Architecture) et OpenCL offrent un modèle de programmation MIMD (Multiple Instructions, Multiple Data) pour exécuter un grand nombre de threads scalaires sur le GPU.
Source : Parallel Programming, Concepts and Practice
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 scalaires est nommé warp chez Nvidia et wavefront chez AMD. Ces warps sont exécutés en même temps sur le système SIMD, d'où le terme SIMT.
Source : Rendered Insecure: GPU Side Channel Attacks are Practical
SAXPY : exemple concret CPU vs GPU
Chaque kernel est composé de milliers de threads, et chaque thread exécute le même programme, mais certains threads vont suivre un flow de contrôle différent selon les dépendances du programme. Un exemple basique : SAXPY (Single precision scalar value A times vector value X Plus vector value Y), soit grossièrement A * X + Y. SAXPY fait partie de la bibliothèque BLAS (Basic Linear Algebra Subprograms). Voici le code 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, avec l'API CUDA de Nvidia, le code est 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
}On divise notre tableau en chunks de taille 256 appelés thread-blocks. Chaque thread-block est exécuté sur un coeur GPU (un SM). Un groupe de thread-blocks est appelé une grid de thread-blocks. Lors de l'appel saxpy<<<nblocks, 256>>>, nblocks représente la taille de la grid et 256 le nombre de threads par thread-block.
Mémoire partagée et synchronisation
Il y a plusieurs types de mémoire dans la programmation GPU/parallèle :
| Type | Emplacement | Portée | Vitesse |
|---|---|---|---|
| Registres | SM (on-chip) | 1 thread | Tres rapide |
| Mémoire partagée | SM (on-chip) | 1 thread-block | Rapide |
| Mémoire locale | GDDR/HBM (off-chip) | 1 thread | Lente |
| Mémoire globale | GDDR/HBM (off-chip) | Tous les threads | Lente |
La mémoire partagée est plus rapide que la mémoire locale et globale (stockées off-chip dans la GDDR ou en cache CPU). Elle est directement sur le GPU (on-chip). Tous les threads d'un thread-block ont accès à la même mémoire partagée. Les accès concurrents ne posent pas de problème grâce à un accès sérialisé : les requêtes concurrentes sont séparées en plusieurs requêtes sans conflit, ce qui réduit l'efficacité de la bande passante GPU. L'exception : quand tous les threads d'un warp ont la même adresse de mémoire partagée, ça résulte en un broadcast, très utile pour la synchronisation.
La synchronisation de threads est supportée par des unités hardware qui font barrière à l'exécution d'instructions.
Source : ArmorAll: Compiler-based Resilience Targeting GPU Applications
ISA GPU : PTX et SASS
Le GPU, comme le CPU, contient son propre ISA (Instruction Set Architecture). Depuis l'introduction de CUDA, l'ISA est haut niveau et virtuel, aussi appelé PTX (Parallel Thread Execution ISA). Le PTX est assez similaire aux instructions d'une architecture processeur RISC (ARM, MIPS, SPARC, etc). Voici le PTX du code CUDA SAXPY ci-dessus :
.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;
}Comme tout code haut niveau, il faut le compiler pour avoir l'ISA compatible avec le GPU ciblé. On transforme donc le PTX en langage assembleur intermédiaire RISC appelé SASS (Streaming ASSembler). Ce process est effectué par un programme nommé ptxas pour les GPU Nvidia avec CUDA.
Source : General-Purpose Graphics Processor Architecture
SASS : peu documenté, difficile à reverse
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 ensuite introduit l'outil cuobjdump pour la documentation du SASS. Aujourd'hui cette documentation correspond juste à la liste des opcodes, sans détail 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érents. Pour chaque nouvelle architecture Nvidia (Tesla, Fermi, Pascal, ...) le SASS est différent.
Nvidia vs AMD : transparence ISA
s_) et vectorielles (v_).Source : General-Purpose Graphics Processor Architecture
Veille sécurité
La suite, une fois par mois
Ce qui bouge vraiment sur la surface d'attaque externe, NIS2 et la sécurité applicative, écrit par le pentester qui signe ces articles. Un envoi par mois, désinscription en un clic.
Votre adresse ne sert qu'à cet envoi. Voir la politique de confidentialité.