Valider les artefacts CUDA sur DGX Spark : cibles GPU, SASS et PTX

Un projet CUDA peut compiler sans erreur sur DGX Spark, puis échouer au lancement d’un noyau avec le message no kernel image is available for execution on the device. Le premier réflexe doit être de vérifier les architectures GPU couvertes par le binaire, et non de réinstaller tout l’environnement. Un autre cas, plus discret, survient lorsque le programme fonctionne uniquement grâce à la compilation à la volée du PTX, alors que les notes de version indiquent à tort la présence de code natif pour le GPU. Cet article propose un petit programme autonome pour réunir la détection de la cible, l’inspection des artefacts et la vérification des deux chemins d’exécution dans une procédure de validation.

1. Distinguer l’architecture hôte de la cible GPU

DGX Spark utilise NVIDIA GB10 Grace Blackwell. Son architecture hôte est ARM64 et son système, NVIDIA DGX OS, repose sur Ubuntu 24.04 avec une pile CUDA préinstallée. Après paiement, les identifiants de connexion sont accessibles dans la console ; l’administration s’effectue ensuite par SSH.

ARM64 décrit le format des exécutables côté CPU ; la capacité de calcul CUDA définit la cible du code destiné au GPU. Ces deux critères de compatibilité sont distincts. Si file indique AArch64, cela confirme uniquement que le format côté hôte est correct, pas que le binaire contient des noyaux utilisables par le GPU.

Le SASS est le code machine destiné au GPU cible, tandis que le PTX est une représentation intermédiaire que le pilote peut compiler à la volée. Un paquet de distribution peut contenir les deux, mais ils doivent être vérifiés séparément : une exécution ordinaire réussie ne valide pas les deux chemins.

Les manipulations ci-dessous créent uniquement des fichiers source et des exécutables dans un nouveau répertoire. Elles ne modifient ni le pilote, ni la chaîne de compilation système, ni les applications en production. Les résultats servent à valider la compatibilité, pas à mesurer les performances.

2. Interroger la capacité de calcul réelle sans deviner la cible

Commencez par créer un répertoire de travail indépendant, puis enregistrez les informations sur le système et la chaîne de compilation. La version CUDA affichée par nvidia-smi correspond au niveau de prise en charge du pilote, et non à la version du toolkit utilisé par nvcc. Il faut donc conserver les deux sorties.

mkdir -p "$HOME/cuda-artifact-check"
cd "$HOME/cuda-artifact-check"
uname -m > host-arch.txt
nvidia-smi > driver.txt
nvcc --version > toolkit.txt
nvcc --list-gpu-code > supported-targets.txt

Enregistrez le contenu suivant dans probe.cu. Ce programme utilise CUDA Runtime pour interroger la capacité de calcul du périphérique zéro. En cas d’échec, il renvoie un code de sortie non nul sans afficher de numéro de cible qui risquerait d’être réutilisé à tort.

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

int main() {
    cudaDeviceProp p{};
    cudaError_t e = cudaGetDeviceProperties(&p, 0);
    if (e != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e));
        return 1;
    }
    std::printf("%d%d\n", p.major, p.minor);
    return 0;
}

nvcc probe.cu -o probe
CC=$(./probe) || exit 1
printf '%s\n' "$CC" > compute-capability.txt
grep -Fx "sm_${CC}" supported-targets.txt || exit 1

Le numéro de cible du GB10 n’est volontairement pas codé en dur. Si la liste du toolkit ne contient pas la cible sm_ détectée, vérifiez d’abord les architectures prises en charge. Ne remplacez pas arbitrairement cette cible par une plus ancienne pour poursuivre la publication. Si la sonde échoue, résolvez d’abord le problème d’accès au périphérique ou de bibliothèques d’exécution, sans passer à l’étape de compilation.

3. Générer et inspecter explicitement le SASS et le PTX

Enregistrez le contenu suivant dans check.cu. Le programme lance un seul thread qui écrit un entier, puis recopie cette valeur vers l’hôte pour la vérifier. Chaque appel CUDA fait l’objet d’un contrôle d’erreur, afin de ne pas annoncer un succès après l’échec du lancement du noyau.

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

#define CUDA_OK(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e)); \
        return 1; \
    } \
} while (0)

__global__ void write_value(int* p) {
    *p = 7;
}

int main() {
    int* device = nullptr;
    int host = 0;
    CUDA_OK(cudaMalloc(&device, sizeof(int)));
    write_value<<<1, 1>>>(device);
    CUDA_OK(cudaGetLastError());
    CUDA_OK(cudaDeviceSynchronize());
    CUDA_OK(cudaMemcpy(&host, device, sizeof(int),
                       cudaMemcpyDeviceToHost));
    CUDA_OK(cudaFree(device));
    if (host != 7) return 2;
    std::puts("PASS");
    return 0;
}

Dans le même répertoire de travail, exécutez les commandes suivantes pour demander explicitement la génération du code natif et du PTX :

CC=$(cat compute-capability.txt)
nvcc -O2 \
  -gencode "arch=compute_${CC},code=[sm_${CC},compute_${CC}]" \
  check.cu -o check

file check > binary-format.txt
cuobjdump --list-elf ./check > elf-list.txt
cuobjdump --dump-sass ./check > sass.txt
cuobjdump --dump-ptx ./check > ptx.txt

Vérifiez que la liste ELF et la sortie SASS contiennent la cible et le noyau attendus, et que la sortie PTX contient .version, .target et le point d’entrée du noyau. Un fichier vide, une cible incorrecte ou l’absence du noyau attendu doivent bloquer la validation. Contrôler uniquement le code de sortie des commandes ne suffit pas.

Inspecter le fichier final, pas les objets intermédiaires

Un projet réel peut comporter des étapes de liaison avec des bibliothèques statiques, de liaison du code pour le GPU ou de création de paquet. La présence de PTX dans un fichier objet ne prouve pas qu’il est conservé dans le fichier livré. Répétez les vérifications sur l’exécutable final et sur les bibliothèques partagées effectivement chargées, puis consignez leurs empreintes.

4. Vérifier séparément l’exécution native et la compilation à la volée

Désactivez d’abord la compilation à la volée du PTX pour obliger l’artefact à utiliser son code natif pour le GPU. Forcez ensuite le chemin PTX en désactivant le cache de compilation sur disque. Chaque test doit démarrer dans un nouveau processus, afin que des modules déjà chargés en mémoire ne faussent pas le résultat.

env -u CUDA_FORCE_PTX_JIT \
  CUDA_DISABLE_PTX_JIT=1 \
  ./check > native.log 2>&1
native_rc=$?

env -u CUDA_DISABLE_PTX_JIT \
  CUDA_FORCE_PTX_JIT=1 CUDA_CACHE_DISABLE=1 \
  ./check > ptx-jit.log 2>&1
ptx_rc=$?

printf 'native=%s\nptx_jit=%s\n' \
  "$native_rc" "$ptx_rc" > execution-status.txt

cat execution-status.txt native.log ptx-jit.log
test "$native_rc" -eq 0 && test "$ptx_rc" -eq 0

Les deux journaux doivent contenir PASS et les deux codes de sortie doivent être nuls. Les variables de test ne s’appliquent qu’à la commande concernée. Ne les ajoutez pas à l’environnement global, au risque de forcer durablement le JIT pour les processus applicatifs.

Si l’exécution native échoue mais que le PTX fonctionne, vérifiez en priorité la cible native et l’étape de liaison finale. Si l’exécution native réussit mais que le PTX échoue, contrôlez la présence du PTX et la prise en charge de sa version par le pilote. Si les deux échouent, appuyez-vous sur les erreurs précises pour examiner la chaîne de compilation, les bibliothèques d’exécution et l’état du périphérique. Modifier sans cesse les options de compilation ne remplace pas un diagnostic.

5. Faire de ces résultats un critère bloquant avant publication

La réussite de ce petit programme prouve seulement que la chaîne de compilation de base et ce noyau fonctionnent dans l’environnement actuel. L’application peut aussi charger des bibliothèques d’extension, sélectionner d’autres noyaux selon les données d’entrée ou générer du code à l’exécution. Il faut donc compléter cette vérification par des tests des points d’entrée réels, couvrant les principaux types de données et branches d’exécution.

Pour chaque publication, conservez quatre ensembles de preuves :

• L’architecture hôte, la sortie du pilote, la version du toolkit et la commande de compilation effectivement utilisée.
• La capacité de calcul du GPU et les résultats d’inspection du SASS et du PTX dans les fichiers finaux.
• Les codes de sortie et les journaux des deux chemins d’exécution, ainsi que les comptes rendus de vérification des résultats applicatifs.
• La version du code source, l’empreinte SHA-256 du fichier final et la liste des dépendances dynamiques.

La commande sha256sum ./check permet d’identifier précisément l’artefact validé. Toute nouvelle édition de liens impose une nouvelle validation, même si le nom du fichier reste identique. Si vous utilisez un conteneur, inspectez la chaîne de compilation et les fichiers finaux dans l’image réellement livrée, puis effectuez les tests d’exécution dans l’environnement où elle sera utilisée.

Le PTX offre un chemin de compatibilité, pas une garantie de fonctionnement avec tous les pilotes. Une exécution réussie ne permet pas non plus de conclure au bon fonctionnement sur d’autres périphériques. Le critère de publication doit être formulé explicitement ainsi : « validation réussie avec le pilote actuel, l’artefact actuel et les points d’entrée couverts ». Tout changement d’architecture ou de chaîne de compilation exige de nouvelles preuves, et non la réutilisation d’anciens journaux.

Questions fréquentes