Aller au contenu

7 · Les bibliothèques NVIDIA

Ce qu'il ne faut jamais réécrire. Ce chapitre est un catalogue raisonné : pour chaque bibliothèque, ce qu'elle fait, quand l'utiliser, et ses pièges.

7.1 Le principe

La règle de base

Avant d'écrire un noyau, vérifiez qu'il n'existe pas. Les bibliothèques NVIDIA représentent des années-personnes d'optimisation, sont réglées pour chaque architecture, et sont mises à jour à chaque génération.

Le corollaire : si votre opération est standard et que vous êtes 5× plus lent qu'une bibliothèque, ce n'est pas que votre code est mauvais — c'est que vous n'auriez pas dû l'écrire.

7.2 Algèbre linéaire

cuBLAS

L'implémentation BLAS de référence.

cublasHandle_t handle;
cublasCreate(&handle);
cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N,
             m, n, k, &alpha,
             A, CUDA_R_16BF, lda,
             B, CUDA_R_16BF, ldb, &beta,
             C, CUDA_R_32F, ldc,
             CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP);

Points d'attention :

  • cublasLt est l'API moderne, plus flexible : elle permet de choisir un algorithme, de fusionner un épilogue (biais, GELU, ReLU), et de gérer les formats FP8 avec échelle ;
  • cuBLAS est column-major par héritage Fortran. Avec des tenseurs row-major (PyTorch, NumPy), on utilise l'identité \((\mathbf{A}\mathbf{B})^\top = \mathbf{B}^\top\mathbf{A}^\top\) pour éviter une transposition ;
  • CUDA 13.1 apporte 2 à 6× sur Blackwell en BF16, FP8 et formats à échelle par blocs, et ajoute le support de la GEMM groupée — décisif pour les MoE.

cuSOLVER, cuSPARSE

Factorisations (LU, QR, Cholesky), valeurs propres, systèmes linéaires denses (cuSOLVER) ou creux (cuSPARSE).

Point d'attention : les performances sur matrices creuses dépendent énormément du format de stockage (CSR, COO, BSR, ELL) et de la structure de la matrice. Un mauvais choix de format coûte un ordre de grandeur.

cuTENSOR

Contractions de tenseurs généralisées, permutations, réductions. Pertinent en chimie quantique, en physique et pour les réseaux de tenseurs.

7.3 Réseaux de neurones

cuDNN

Convolutions, normalisations, activations, pooling, RNN, et l'attention (fused multi-head attention).

Le mécanisme important : cuDNN dispose de plusieurs algorithmes par opération et choisit selon les dimensions. En PyTorch :

torch.backends.cudnn.benchmark = True   # mesure et choisit au premier appel

À activer si vos formes sont stables. Sinon, chaque nouvelle forme déclenche une phase de mesure coûteuse.

cuDNN contre FlashAttention

Les deux implémentent l'attention. FlashAttention-4 mesure 1,3× plus rapide que cuDNN 9.13 sur B200 (1 605 TFLOPS contre ~1 235). L'écart n'est pas énorme, ce qui témoigne du niveau de cuDNN.

En pratique : essayez les deux, mesurez sur vos formes.

TensorRT et TensorRT-LLM

Compilateurs d'inférence. TensorRT prend un graphe (ONNX ou autre), fusionne, sélectionne les noyaux, quantifie, et produit un moteur optimisé.

TensorRT-LLM se spécialise sur les modèles de langage : paged attention, in-flight batching, quantification, décodage spéculatif.

Point notable pour la partie 8 : Ada-MK s'intègre à TensorRT-LLM sous forme de greffon, ce qui permet de combiner un moteur classique haut débit pour le prefill et un megakernel basse latence pour le décodage.

7.4 Primitives parallèles : CCCL

CCCL (CUDA Core Compute Libraries) regroupe trois bibliothèques historiquement séparées.

Composant Niveau Contenu
Thrust haut algorithmes à la STL : sort, reduce, transform, scan
CUB moyen primitives par niveau : device, block, warp
libcu++ bas cuda::std::, cuda::atomic, cuda::barrier, cuda::pipeline 
// Thrust : simple
thrust::device_vector<float> v(n);
float s = thrust::reduce(v.begin(), v.end());

// CUB : contrôle du tampon temporaire, plus rapide
void* d_temp = nullptr; size_t temp_bytes = 0;
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, n);
cudaMalloc(&d_temp, temp_bytes);
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, n);

// CUB au niveau du bloc, dans votre propre noyau
using BlockReduce = cub::BlockReduce<float, 256>;
__shared__ typename BlockReduce::TempStorage stockage;
float somme = BlockReduce(stockage).Sum(ma_valeur);

Ce que CUB apporte

cub::DeviceRadixSort trie des milliards de clés à la bande passante mémoire. cub::DeviceScan fait un scan au prix d'une copie grâce au decoupled look-back. Ce sont des algorithmes non triviaux que personne ne devrait réimplémenter.

CUDA 13.1 simplifie les APIs CUB en une seule phase (plus besoin du double appel pour dimensionner le tampon) et ajoute des réductions flottantes déterministes.

7.5 Transformées et signal

Bibliothèque Contenu
cuFFT transformées de Fourier 1D à 3D, réelles et complexes
cuFFTDx FFT à l'intérieur d'un noyau (device extension), fusionnable
NPP primitives de traitement d'image (filtres, transformations, couleurs)
nvJPEG décodage JPEG accéléré
CV-CUDA opérations de vision par ordinateur pour les pipelines d'IA

cuFFTDx mérite l'attention : il permet d'exécuter une FFT dans votre noyau, sans repasser par la mémoire globale. C'est la même logique de fusion que partout ailleurs, appliquée aux transformées.

7.6 Données et bases de données

RAPIDS est une suite d'analyse de données sur GPU :

Composant Équivalent CPU
cuDF pandas
cuML scikit-learn
cuGraph NetworkX
cuSpatial GeoPandas
cuVS recherche vectorielle (index CAGRA, IVF-PQ) 
import cudf
df = cudf.read_parquet("gros_fichier.parquet")
resultat = df.groupby("cle").agg({"valeur": "sum"})

Accélérations typiques de 10 à 50× sur des jointures et agrégations de gros volumes, à condition que les données tiennent en VRAM — la contrainte dominante.

Depuis 2023, cudf.pandas permet d'accélérer du code pandas existant sans le modifier, avec repli automatique sur le CPU.

7.7 Communication multi-GPU

NCCL

La bibliothèque de collectives : AllReduce, AllGather, ReduceScatter, Broadcast, AllToAll. Elle détecte la topologie (NVLink, NVSwitch, PCIe, InfiniBand) et choisit les algorithmes (anneau, arbre) en conséquence.

C'est ce qu'utilisent PyTorch DDP/FSDP, DeepSpeed et Megatron.

NVSHMEM

Un modèle PGAS (Partitioned Global Address Space) : chaque GPU expose une mémoire dite symétrique que les autres peuvent lire et écrire par des opérations unilatérales (put, get, atomiques), depuis l'intérieur d'un noyau.

__global__ void noyau() {
    // Écrire directement dans la mémoire d'un GPU distant, depuis le noyau
    nvshmem_float_put(dest, source, nelems, pe_distant);
}

C'est ce qui rend possible le recouvrement calcul/communication à grain fin — et c'est la brique sur laquelle repose DeepEP, la bibliothèque de DeepSeek pour les MoE distribués. Notez toutefois que DeepEP V2 utilise par défaut un back-end NCCL (« Gin ») et ne dépend plus de NVSHMEM que pour des chemins hérités.

Mirage MPK utilise NVSHMEM pour la communication inter-GPU à l'intérieur de son megakernel.

7.8 Le tableau récapitulatif

Besoin Bibliothèque
GEMM cuBLAS / cuBLASLt
GEMM avec épilogue custom CUTLASS
GEMM groupée (MoE) cuBLAS (13.1+) ou CUTLASS
Convolution cuDNN
Attention FlashAttention, FlashInfer, cuDNN
Tri, scan, réduction CUB / Thrust
FFT cuFFT, cuFFTDx
Algèbre creuse cuSPARSE
Factorisations cuSOLVER
Traitement d'image NPP, CV-CUDA
DataFrames, ML classique RAPIDS
Recherche vectorielle cuVS
Collectives NCCL
Communication dans le noyau NVSHMEM
Inférence LLM TensorRT-LLM, vLLM, SGLang
Atomiques, barrières, pipelines libcu++

Résumé du chapitre

À retenir

  • Ne réécrivez pas ce qui existe. Ces bibliothèques sont réglées par architecture et mises à jour à chaque génération.
  • cuBLAS est column-major ; utilisez l'identité de transposition avec des tenseurs row-major. cuBLASLt permet l'épilogue fusionné et la GEMM groupée.
  • CCCL = Thrust (haut) + CUB (moyen) + libcu++ (bas). CUB au niveau du bloc s'utilise dans vos propres noyaux.
  • torch.backends.cudnn.benchmark = True seulement si vos formes sont stables.
  • NCCL pour les collectives, NVSHMEM pour la communication émise depuis l'intérieur d'un noyau — la brique des megakernels multi-GPU.
  • Écrivez un noyau quand votre opération est inhabituelle, pas quand elle est standard.

Vérifiez que vous avez compris

Pourquoi cuBLAS est-il column-major alors que tout le monde utilise du row-major ?

Héritage du BLAS Fortran des années 1970, dont cuBLAS reprend l'interface pour la compatibilité.

En pratique, on exploite l'identité :

\[(\mathbf{A}\mathbf{B})^\top = \mathbf{B}^\top \mathbf{A}^\top\]

Une matrice row-major \(M \times N\) est exactement une matrice column-major \(N \times M\) transposée. Il suffit donc d'échanger les opérandes et les dimensions dans l'appel, sans aucune copie. C'est ce que fait PyTorch en interne.

Votre pipeline RAPIDS est plus lent que pandas. Causes possibles ?

Quatre, par ordre de fréquence :

  1. Le volume est trop petit. En dessous de quelques centaines de mégaoctets, le transfert PCIe et le lancement des noyaux dominent.
  2. Les données font des allers-retours. Chaque .to_pandas() ou .to_arrow() coûte un transfert PCIe complet.
  3. L'opération est mal supportée et retombe silencieusement sur le CPU (fréquent avec cudf.pandas).
  4. Les chaînes de caractères. Le traitement de chaînes de longueur variable est le point faible structurel du GPU : accès irréguliers, divergence, indirections.
Quand NVSHMEM plutôt que NCCL ?

NCCL quand la communication est une phase distincte du calcul : un AllReduce après le backward, un AllGather avant une couche. C'est optimisé, robuste, et l'API est simple.

NVSHMEM quand la communication doit être entrelacée avec le calcul à l'intérieur d'un noyau : envoyer les premiers résultats pendant que les suivants se calculent. Cela permet un recouvrement impossible avec des collectives à granularité de noyau.

C'est exactement l'usage des megakernels multi-GPU : MPK, qui atteint 1,1-1,4× sur 8 H100 contre SGLang/vLLM, s'appuie sur NVSHMEM pour « maximiser le recouvrement du calcul, du chargement de données et de la communication inter-GPU entre les couches ».

Chapitre suivant : 8 · AMD, HIP et ROCm

Sources de ce chapitre