Repérer les accès CUDA hors limites sur DGX Spark

Le noyau CUDA a bien été lancé, mais l’erreur n’apparaît qu’au moment de rapatrier les résultats ; le même programme fonctionne à nouveau lorsque l’on réduit la taille de l’entrée. Face à ces symptômes, il faut d’abord vérifier les accès à la mémoire du GPU, plutôt que de modifier immédiatement le pilote. Sur DGX Spark, un petit programme autonome permet de construire une démarche de diagnostic complète : provoquer volontairement un accès hors limites, retrouver la ligne de code concernée, corriger l’indice, puis vérifier que la détection des erreurs n’a pas été désactivée. Cette procédure ne nécessite aucun poids de modèle et n’implique aucun service en production.

1. Fixer l’environnement et délimiter le périmètre du test

Le DGX Spark de TcpNode repose sur NVIDIA GB10 Grace Blackwell, une architecture ARM64 et NVIDIA DGX OS, avec 128 Go de mémoire unifiée LPDDR5x. Le fait que cette mémoire soit partagée entre le CPU et le GPU n’autorise pas un pointeur à dépasser les limites de son allocation. Nous utiliserons ici cudaMalloc pour allouer la mémoire du GPU et définir ainsi une plage d’adresses valide sans ambiguïté.

Après le paiement, consultez les identifiants de connexion dans la console, puis connectez-vous en SSH. Effectuez les tests dans un répertoire de travail distinct, avec un compte utilisateur ordinaire. N’injectez pas d’erreur dans un processus d’inférence qui traite des requêtes.

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

L’architecture affichée doit être aarch64. Consignez les versions des outils et les informations du pilote, puis conservez le même environnement pour les vérifications ultérieures. Une pile CUDA préinstallée ne signifie pas que tous ses outils sont accessibles dans le PATH actuel. Si une commande est introuvable, vérifiez d’abord la chaîne d’outils installée et ses chemins d’installation, sans remplacer d’emblée l’ensemble des pilotes.

L’exemple d’accès hors limites doit être exécuté uniquement dans un processus isolé que l’on peut abandonner sans conséquence. Après un accès mémoire illégal, arrêtez ce processus et relancez-le : les résultats obtenus ensuite dans le même contexte CUDA ne constituent pas des éléments de preuve valides.

2. Construire un cas de reproduction en ne faisant varier qu’un seul paramètre

Enregistrez le contenu suivant dans repro.cu. Le programme n’alloue qu’un seul entier : le seul indice valide est donc zéro. La position d’écriture est choisie à la compilation avec TEST_INDEX. La version fautive et la version corrigée partagent ainsi le même code source, sans modifier simultanément l’entrée, la taille de l’allocation et le chemin d’exécution.

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

static void checked(cudaError_t status) {
    if (status != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(status));
        std::exit(EXIT_FAILURE);
    }
}

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

Compilez les deux programmes séparément :

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

L’option -arch=native cible le GPU actuellement visible. Elle convient à une vérification locale, mais pas à la production d’un binaire destiné à différents GPU. Si le compilateur ne prend pas en charge le GPU présent, vérifiez d’abord la compatibilité de la chaîne d’outils, sans choisir arbitrairement des paramètres correspondant à une autre architecture. L’option -lineinfo conserve les informations de lignes du code source exécuté sur le GPU. Pour cette première localisation, il n’est pas nécessaire d’activer le mode complet de débogage du code GPU, qui modifierait sensiblement les caractéristiques d’exécution.

Pourquoi vérifier à la fois le lancement et la synchronisation

cudaGetLastError vérifie les erreurs liées au lancement, mais l’exécution du noyau est généralement asynchrone. cudaDeviceSynchronize attend la fin des travaux précédemment soumis au GPU ; les erreurs survenues pendant leur exécution peuvent alors être signalées à cet endroit. Il faut vérifier séparément ces deux appels, et non se contenter de voir si le programme affiche un résultat.

3. Rendre le résultat de la détection exploitable par le code de sortie

Exécutez d’abord la version fautive en enregistrant séparément la sortie standard et le journal de détection :

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

Le rapport doit signaler une écriture illégale en mémoire globale et désigner l’opération d’écriture dans write_value. Il indique généralement aussi la taille de l’accès, la position du thread et la relation entre l’adresse visée et l’allocation. Sa présentation varie selon la version de l’outil : une formulation précise ne doit donc pas être l’unique critère de détection.

L’option --error-exitcode 86 permet aux erreurs détectées de se répercuter sur le résultat d’une exécution automatisée. Toutefois, si le programme testé échoue lui-même, il est également possible d’obtenir un autre code non nul renvoyé par ce programme. Le critère de validation est donc « code de sortie non nul et rapport désignant l’accès hors limites attendu », et non l’obligation de renvoyer 86 pour toute défaillance.

Évitez de raccorder directement la commande à un simple pipeline de journalisation en ne consultant que le statut de la dernière commande du pipeline. Les redirections et la structure conditionnelle ci-dessus préservent le véritable résultat de l’exécution, y compris dans un script configuré pour s’arrêter dès qu’une commande échoue.

4. Corriger l’indexation selon la plage d’adresses, sans masquer les symptômes

Dans cet exemple, l’allocation contient un seul entier, mais l’écriture utilise l’indice un : elle dépasse donc la plage autorisée. Dans un programme réel, il faut relier l’adresse indiquée dans le rapport à l’allocation correspondante, puis examiner le calcul des indices, le type des éléments, les pas mémoire et les décalages entre lots.

Parmi les causes courantes : une taille d’allocation exprimée en octets alors que les indices sont interprétés en nombre d’éléments ; un dernier bloc de threads incomplet ; un pas entre lignes différent de la largeur logique dans un tableau à deux dimensions ; ou encore un pointeur réutilisé par une autre tâche après sa libération. Agrandir l’allocation peut masquer temporairement l’erreur, mais ne prouve pas que l’algorithme est correct.

Pour un traitement unidimensionnel, il faut généralement vérifier index < count avant l’accès et s’assurer que le type entier utilisé pour calculer l’indice peut représenter le produit obtenu. Dans un noyau où les threads coopèrent en mémoire partagée, on ne peut pas faire sortir prématurément certains threads sans précaution, puis laisser les autres atteindre une barrière de synchronisation à l’échelle du bloc : cela risque de créer un nouveau problème de synchronisation.

Si l’erreur n’apparaît que dans un traitement complexe, ajoutez temporairement une vérification avec synchronisation après une étape suspecte, puis réduisez progressivement le périmètre de recherche. N’ajoutez qu’un point de contrôle à la fois et conservez la première erreur. Ne continuez pas à soumettre des tâches après celle-ci, au risque de générer de nombreuses erreurs secondaires.

5. Valider le correctif avec un test positif et un test négatif

Exécutez la version corrigée :

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

Dans cet exemple, trois conditions doivent être remplies simultanément : un code de sortie nul, un rapport de détection ne signalant aucune erreur d’accès mémoire et la sortie result=7. L’absence de message d’erreur ne suffit pas : le test pourrait sembler réussi alors que le noyau ne s’est pas exécuté ou que le résultat n’a pas réellement été écrit.

Relancez ensuite la version fautive pour vérifier qu’elle échoue toujours. Ce test négatif permet de repérer une commande de détection supprimée, l’utilisation accidentelle d’un ancien binaire ou un pipeline qui masque le statut d’échec. Pour un noyau réel, les tests doivent aussi couvrir l’entrée valide minimale, les tailles correspondant à des blocs complets et les entrées dont la partie finale ne remplit pas un bloc entier. Si l’interface accepte une entrée vide, son traitement doit également être défini explicitement.

Archivez enfin le code source, les commandes de compilation, les informations de version, les sommes de contrôle et les journaux des deux exécutions. Le détecteur modifie le coût d’exécution : les temps mesurés sous son contrôle ne permettent pas de tirer directement des conclusions sur les performances. De même, un passage sans erreur sous memcheck ne garantit pas l’absence de courses de données ou de problèmes de synchronisation. Une validation fiable doit fournir des éléments vérifiables distincts pour la validité des accès, le résultat du calcul et la propagation des échecs, plutôt qu’une simple capture d’écran montrant un succès.

Questions fréquentes