Coder pour le GPU : threads,
blocks, grids

Pré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.

  CPU - Exécution out-of-order (OoO)
  ─────────────────────────────────────────────
  Programme (ordre logique)     Pipeline CPU
  ─────────────────────────     ─────────────────────────────────────
  1. LOAD  R1, [addr_A]    →    [Fetch] [Decode] [Issue queue]
  2. ADD   R2, R1, R3           ↓ dépend de R1 (attente)
  3. LOAD  R4, [addr_B]    →    exécutée EN AVANCE (indépendante)
  4. MUL   R5, R4, R6      →    exécutée EN AVANCE
  5. ADD   R7, R2, R5           ↓ attend R2 et R5

  Résultat : instructions 3 et 4 s'exécutent avant 2
             sans changer le résultat final.
  GPU : SIMT, pas d'OoO. Chaque warp suit un chemin linéaire.

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.

  Avant Pascal (GPU discret classique)      Depuis Pascal (mémoire unifiée)
  ──────────────────────────────────        ───────────────────────────────────
  CPU RAM  ←─ PCIe ─→  GPU VRAM            CPU RAM  ─┐
  cudaMemcpy() explicite                   GPU VRAM ─┤  espace d'adressage
  h_ptr ≠ d_ptr                            cuDMA     │  virtuel UNIFIE
                                           ptr unique┘
  Code :                                   Code :
  cudaMalloc(&d_x, size);                  cudaMallocManaged(&x, size);
  cudaMemcpy(d_x, h_x, ...);              saxpy<<<...>>>(n, a, x, y);
  saxpy<<<...>>>(n, a, d_x, d_y);         // pas de cudaMemcpy !
  cudaMemcpy(h_x, d_x, ...);

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.

  GPU
  ┌──────────────────────────────────────────────────────┐
  │  SM 0          SM 1          SM 2         SM N-1     │
  │ ┌──────┐      ┌──────┐      ┌──────┐     ┌──────┐   │
  │ │Block │      │Block │      │Block │     │Block │   │
  │ │  0   │      │  1   │      │  2   │     │ ...  │   │
  │ │256 th│      │256 th│      │256 th│     │256 th│   │
  │ └──────┘      └──────┘      └──────┘     └──────┘   │
  │  SIMT          SIMT          SIMT          SIMT      │
  │                                                      │
  │         Mémoire GDDR / HBM2 (globale)                │
  └──────────────────────────────────────────────────────┘
  Chaque SM = ~1000 threads actifs en même temps

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.

  ┌──────────────────────────────────────────────────────────────┐
  │           Taxonomie de Flynn                                 │
  │                                                              │
  │  SISD                        SIMD                           │
  │  Single Instr / Single Data  Single Instr / Multiple Data   │
  │  → CPU classique (1 coeur)   → GPU interne, SSE/AVX CPU     │
  │                                                              │
  │  MISD                        MIMD                           │
  │  Multiple Instr / Single D.  Multiple Instr / Multiple D.   │
  │  → rare (tolérance de faute) → CUDA/OpenCL vu du prog.      │
  │                              → multi-coeurs CPU, clusters   │
  └──────────────────────────────────────────────────────────────┘

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.

  Thread-block (256 threads)
  ─────────────────────────────────────────────────────
  Warp 0  : threads  0 -  31  ──→  SIMD unit (32 lanes)
  Warp 1  : threads 32 -  63  ──→  SIMD unit
  Warp 2  : threads 64 -  95  ──→  SIMD unit
  ...
  Warp 7  : threads 224 - 255 ──→  SIMD unit

  Nvidia : warp     = 32 threads
  AMD    : wavefront = 64 threads

  Tous les threads d'un warp exécutent la MEME instruction
  au même cycle (lock-step). Si des threads divergent
  (if/else), les branches sont sérialisées : SIMT divergence.

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
}

Ce qu'il faut comprendre ici : 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.

  Grid (nblocks thread-blocks)
  ┌─────────────────────────────────────────────────────┐
  │  Block 0          Block 1          Block nblocks-1  │
  │ ┌───────────┐    ┌───────────┐    ┌───────────┐     │
  │ │ t0  t1    │    │ t0  t1    │    │ t0  t1    │     │
  │ │ t2  t3    │    │ t2  t3    │    │ t2  t3    │     │
  │ │ ...       │    │ ...       │    │ ...       │     │
  │ │ t254 t255 │    │ t254 t255 │    │ t254 t255 │     │
  │ └───────────┘    └───────────┘    └───────────┘     │
  │   → SM 0           → SM 1          → SM k           │
  └─────────────────────────────────────────────────────┘

  Chaque thread calcule son index global :
  i = blockIdx.x * blockDim.x + threadIdx.x
      (numéro du bloc × taille du bloc + position dans le bloc)

Mémoire partagée et synchronisation

Il y a plusieurs types de mémoire dans la programmation GPU/parallèle :

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.

  SM (Streaming Multiprocessor)
  ┌────────────────────────────────────────────────────────┐
  │                                                        │
  │  Thread-Block 0                                        │
  │  ┌──────────────────────────────────────────────┐     │
  │  │  t0    t1    t2    ...    t255                │     │
  │  │   │     │     │            │                  │     │
  │  │   └─────┴──┬──┴────────────┘                  │     │
  │  │            ▼                                  │     │
  │  │    Shared Memory (on-chip, ~48KB)             │     │
  │  │    acces rapide, portee = 1 block             │     │
  │  └──────────────────────────────────────────────┘     │
  │             │                                          │
  └─────────────│──────────────────────────────────────────┘
                │
                ▼
  GDDR / HBM - Mémoire Globale (off-chip, GBs)
  acces lent, portee = tous les threads de la grid
  ─────────────────────────────────────────────────────────
  Broadcast : si tous les threads d'un warp lisent
  la meme adresse → 1 seule lecture + diffusion

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.

  Code CUDA (.cu)
       │
       ▼  nvcc (compilateur Nvidia)
  PTX - ISA virtuel haut niveau (portable)
       │
       ▼  ptxas (compilateur PTX -> SASS)
  SASS - ISA machine (specifique a l'architecture)
       │
       ├── Tesla  → SASS v1.x
       ├── Fermi  → SASS v2.x
       ├── Pascal → SASS v6.x
       └── Ampere → SASS v8.x

  Le PTX est stable entre architectures.
  Le SASS change a chaque generation de GPU.

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                             AMD
  ────────────────────────────       ────────────────────────────────
  Source CUDA (.cu)                  Source OpenCL / HIP (.cl)
       │                                  │
       ▼ nvcc                             ▼ clang / hipcc
  PTX (ISA virtuel)                  HSAIL (ISA virtuel)
       │                                  │
       ▼ ptxas                            ▼ LC codegen
  SASS (ISA machine)                 GCN ISA (ISA machine)
       │                                  │
       ├── opcodes listés seulement       ├── spec publique complète
       └── sémantique non documentée      ├── s_ = instructions scalaires
                                          └── v_ = instructions vectorielles

Source : General-Purpose Graphics Processor Architecture