2 · Mettre en place l'environnement

Pilote, toolkit, PyTorch, Triton, profileurs. La partie ingrate, et celle qui fait abandonner le plus de gens.

---

2.1 La pile, et qui dépend de qui

┌─────────────────────────────────────────────┐
│ Votre code                                  │
├─────────────────────────────────────────────┤
│ PyTorch / Triton / CuPy / Numba             │
├─────────────────────────────────────────────┤
│ CUDA Toolkit  (nvcc, cuBLAS, cuDNN, Nsight) │
├─────────────────────────────────────────────┤
│ Pilote NVIDIA  (nvidia-smi, libcuda.so)     │
├─────────────────────────────────────────────┤
│ Noyau Linux / Windows                       │
└─────────────────────────────────────────────┘

La règle de compatibilité : le pilote doit être au moins aussi récent que le toolkit. Un pilote récent exécute du code compilé avec un toolkit plus ancien (compatibilité descendante) ; l'inverse est faux.

nvidia-smi          # ligne "CUDA Version" = version MAXIMALE supportée par le pilote
nvcc --version      # version du toolkit installé

La confusion la plus fréquente

nvidia-smi affiche CUDA Version: 13.1. Cela ne signifie pas que CUDA 13.1 est installé : cela signifie que le pilote peut exécuter jusqu'à CUDA 13.1.

La version réellement installée est celle de nvcc --version. Les deux peuvent différer, et c'est normal.

---

2.2 Installer

Le pilote

Linux (Ubuntu/Debian) :

sudo apt update
sudo apt install nvidia-driver-580        # numéro à adapter
sudo reboot
nvidia-smi                                 # vérifier

Attention : depuis CUDA 13.0, le toolkit Windows ne fournit plus de pilote. Il faut l'installer séparément.

Le toolkit

La méthode qui évite le plus de problèmes :

# Ubuntu 24.04, exemple
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt update
sudo apt install cuda-toolkit-13-1

Puis dans votre ~/.bashrc :

export CUDA_HOME=/usr/local/cuda
export PATH=$CUDA_HOME/bin:$PATH
export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH

La méthode conda (souvent plus simple)

conda create -n gpu python=3.12
conda activate gpu
conda install -c nvidia cuda-toolkit=13.1

Avantage : environnements isolés, plusieurs versions de CUDA cohabitent. C'est la méthode recommandée quand on travaille sur plusieurs projets.

PyTorch

Toujours prendre la commande sur pytorch.org, qui génère la ligne correcte pour votre combinaison.

# Exemple (à adapter à la version courante)
pip install torch torchvision --index-url https://download.pytorch.org/whl/cu130

PyTorch embarque ses propres bibliothèques CUDA (cuBLAS, cuDNN, NCCL). Vous n'avez besoin du toolkit complet que pour compiler vos propres noyaux.

import torch
print(torch.__version__)
print(torch.version.cuda)         # version CUDA compilée dans PyTorch
print(torch.cuda.is_available())

Triton

pip install triton

Livré avec PyTorch depuis la 2.x sur Linux. Vérification :

import triton
print(triton.__version__)

Les profileurs

Nsight Systems et Nsight Compute sont inclus dans le toolkit, mais souvent dans une version ancienne. Les versions autonomes sont préférables :

# Vérifier
nsys --version
ncu --version

Téléchargement des dernières versions sur le site développeur NVIDIA.

Autorisation des compteurs (sinon ncu échoue avec ERR_NVGPUCTRPERM) :

sudo sh -c 'echo "options nvidia NVreg_RestrictProfilingToAdminUsers=0" \
  > /etc/modprobe.d/nvidia-profiler.conf'
sudo update-initramfs -u
sudo reboot

---

2.3 Docker, la solution aux problèmes de version

Si vous voulez éviter tous les conflits :

# Installer le NVIDIA Container Toolkit
sudo apt install nvidia-container-toolkit
sudo systemctl restart docker

# Lancer un conteneur avec accès GPU
docker run --gpus all -it --rm \
  -v $(pwd):/workspace \
  nvcr.io/nvidia/pytorch:25.12-py3

Les images nvcr.io/nvidia/pytorch contiennent PyTorch, CUDA, cuDNN, NCCL, TensorRT, Nsight — tout cohérent, testé ensemble, et mis à jour mensuellement.

La recommandation

Pour apprendre : conda ou Docker.

Pour un déploiement : Docker, sans hésiter. Les conflits de versions CUDA en production sont une source de pannes classique.

---

2.4 Un squelette de projet

mon-projet/
├── CMakeLists.txt
├── src/
│   ├── main.cu
│   └── noyaux.cuh
├── tests/
│   └── test_noyaux.py
├── bench/
│   └── bench.py
└── README.md

CMakeLists.txt :

cmake_minimum_required(VERSION 3.24)
project(mon_projet LANGUAGES CXX CUDA)

set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_ARCHITECTURES native)   # ou "90a" pour Hopper

add_executable(mon_projet src/main.cu)

target_compile_options(mon_projet PRIVATE
    $<$<COMPILE_LANGUAGE:CUDA>:
        -O3
        -lineinfo                       # pour le profileur
        --expt-relaxed-constexpr
        -Xptxas=-v                      # afficher registres et spills
    >)

cmake -B build -S .
cmake --build build -j

Compilation rapide sans CMake

nvcc -O3 -arch=native -lineinfo -Xptxas -v mon_noyau.cu -o mon_noyau

Les options à connaître :

Option	Effet
-arch=native	détecte la carte présente
-arch=sm_90a	Hopper avec TMA/wgmma
-O3	optimisations
-lineinfo	métadonnées pour le profileur
-Xptxas -v	registres, mémoire partagée, spills
-G	débogage complet — jamais pour mesurer
-use_fast_math	intrinsèques rapides
--ptxas-options=-warn-spills	avertir sur les spills

---

2.5 Le premier test

Un fichier verif.cu qui valide toute la chaîne :

#include <cstdio>

__global__ void saluer() {
    printf("Bloc %d, thread %d\n", blockIdx.x, threadIdx.x);
}

int main() {
    int n = 0;
    cudaGetDeviceCount(&n);
    printf("GPU détectés : %d\n", n);

    for (int i = 0; i < n; ++i) {
        cudaDeviceProp p;
        cudaGetDeviceProperties(&p, i);
        printf("[%d] %s  cc %d.%d  SM=%d  mem=%.1f Go  smem/bloc=%zu Ko\n",
               i, p.name, p.major, p.minor, p.multiProcessorCount,
               p.totalGlobalMem / 1e9, p.sharedMemPerBlock / 1024);
    }

    saluer<<<2, 4>>>();
    cudaError_t err = cudaDeviceSynchronize();
    printf("Statut : %s\n", cudaGetErrorString(err));
    return 0;
}

nvcc -arch=native verif.cu -o verif && ./verif

Si cela affiche votre carte et huit lignes de salutation, la chaîne fonctionne.

---

2.6 Les pièges de version

Les six problèmes classiques

1. CUDA error: no kernel image is available for execution on the device → Le binaire n'a pas été compilé pour votre architecture. Ajoutez -arch=native ou le bon sm_XX.

2. libcudart.so.12: cannot open shared object file → LD_LIBRARY_PATH ne contient pas $CUDA_HOME/lib64.

3. PyTorch dit torch.cuda.is_available() == False → Version de PyTorch compilée pour une autre version de CUDA, ou pilote trop ancien. Réinstallez depuis pytorch.org.

4. ERR_NVGPUCTRPERM avec ncu → Droits de profilage non accordés. Voir §2.2.

5. Compilation extrêmement lente → Trop d'architectures cibles. -arch=native en développement, la liste complète seulement pour distribuer.

6. Résultats différents entre deux machines → TF32 activé par défaut, versions de cuDNN différentes, ou algorithme sélectionné différemment par cudnn.benchmark. Voir Fondations 6.

---

2.7 L'environnement de mesure

Pour des mesures reproductibles :

# Verrouiller les fréquences (nécessite les droits root)
sudo nvidia-smi -pm 1                    # mode persistant
sudo nvidia-smi -lgc 1410,1410           # horloge graphique
sudo nvidia-smi --lock-memory-clocks=1593

# Vérifier
nvidia-smi -q -d CLOCK

# Restaurer
sudo nvidia-smi -rgc
sudo nvidia-smi -rmc

Les valeurs d'horloge dépendent de la carte ; nvidia-smi -q -d SUPPORTED_CLOCKS les liste.

Sans ce verrouillage, la variabilité thermique et le boost dynamique introduisent ±15 % de variation entre deux exécutions identiques.

---

2.8 Le côté AMD

# Vérifier
rocm-smi
hipcc --version

# Compiler
hipcc -O3 --offload-arch=gfx942 mon_noyau.hip -o mon_noyau

# Convertir du CUDA
hipify-clang --inplace mon_noyau.cu

# Profiler
rocprof --stats ./mon_binaire

PyTorch pour ROCm :

pip install torch --index-url https://download.pytorch.org/whl/rocm7.0

Les cibles courantes : gfx90a (MI250X), gfx942 (MI300X), gfx950 (MI350X/MI355X).

---

Résumé du chapitre

À retenir

• Le pilote doit être au moins aussi récent que le toolkit. La ligne « CUDA Version » de nvidia-smi est un maximum supporté, pas la version installée.
• conda ou Docker évitent la plupart des conflits. Les images nvcr.io/nvidia/pytorch sont cohérentes et testées.
• Toujours compiler avec -lineinfo (profileur) et -Xptxas -v (spills). Jamais -G pour mesurer.
• Autoriser les compteurs de profilage (NVreg_RestrictProfilingToAdminUsers=0), sinon ncu échoue.
• Verrouiller les fréquences pour des mesures reproductibles : sans cela, ±15 % de variation.
• Six erreurs classiques, toutes de version : mémorisez-les, elles reviendront.

---

Vérifiez que vous avez compris

nvidia-smi affiche CUDA 13.1, nvcc --version affiche 12.4. Est-ce un problème ?

Non, c'est normal et fonctionnel.

nvidia-smi indique que le pilote supporte jusqu'à CUDA 13.1. nvcc indique que le toolkit installé est en 12.4.

Un pilote récent exécute du code compilé avec un toolkit plus ancien : c'est la compatibilité descendante garantie par NVIDIA.

En revanche, l'inverse échoue : un toolkit 13.1 avec un pilote ne supportant que 12.4 produira CUDA driver version is insufficient for CUDA runtime version.

Conséquence pratique : le pilote se met à jour, le toolkit s'installe.

Pourquoi -arch=native en développement mais pas pour distribuer ?

-arch=native détecte la carte de la machine de compilation et ne génère que son code SASS. Le binaire ne tournera nulle part ailleurs.

En développement, c'est idéal : compilation rapide (une seule architecture), et on cible exactement sa carte.

Pour distribuer, il faut plusieurs cibles plus du PTX pour la compatibilité future :

nvcc -gencode arch=compute_80,code=sm_80 \
     -gencode arch=compute_86,code=sm_86 \
     -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     mon_noyau.cu

La dernière ligne embarque du PTX, que le pilote compilera en JIT pour une architecture plus récente inconnue à la compilation.

Le coût : le temps de compilation est multiplié par le nombre de cibles.

Votre noyau donne des résultats différents entre votre RTX 4090 et le H100 du cloud. Que vérifier ?

Dans l'ordre de probabilité :

1. TF32. Activé par défaut dans PyTorch, il tronque la mantisse à 10 bits sur les tensor cores. Les deux cartes peuvent le gérer différemment selon les versions. Testez avec torch.backends.cuda.matmul.allow_tf32 = False.
2. Ordre de réduction. Le nombre de SM diffère (128 contre 132), donc le nombre de blocs et l'ordre d'accumulation aussi. L'addition flottante n'étant pas associative, le résultat varie.
3. Algorithme cuDNN. Avec cudnn.benchmark = True, l'algorithme est choisi par mesure : il peut différer d'une carte à l'autre.
4. Instructions différentes. Une architecture peut utiliser tcgen05 là où l'autre utilise mma, avec des ordres d'accumulation internes différents.

Si l'écart est de l'ordre de \(10^{-6}\) en relatif, c'est normal. Au-delà de \(10^{-3}\), c'est TF32 ou un bug.

---

Chapitre suivant : 3 · Intégrer un noyau dans PyTorch

---

Sources de ce chapitre

• CUDA Installation Guide
• PyTorch — Get Started Locally
• NVIDIA NGC Containers
• Nsight Compute — permissions de profilage
• ROCm installation guide
• What's New in CUDA Toolkit 13.0 — fin des pilotes Windows fournis avec le toolkit.