Проверка CUDA-сборки на DGX Spark: архитектура GPU, SASS и PTX

Если CUDA-проект на DGX Spark компилируется, но при запуске ядра возникает ошибка no kernel image is available for execution on the device, первым делом проверьте, для каких архитектур GPU собраны артефакты, а не переустанавливайте всё окружение. Есть и менее очевидная проблема: приложение всё время работает за счёт JIT-компиляции PTX, хотя в документации к релизу ошибочно указано, что сборка содержит нативный код устройства. В этой статье небольшая автономная программа связывает определение целевой архитектуры, проверку артефактов и испытание двух путей выполнения в единую процедуру приёмки.

1. Разделяйте архитектуру хоста и целевую архитектуру GPU

DGX Spark построен на NVIDIA GB10 Grace Blackwell. Архитектура хоста — ARM64, операционная система — NVIDIA DGX OS на базе Ubuntu 24.04 с предустановленным стеком CUDA. После оплаты данные для подключения можно посмотреть в консоли, а дальнейшую работу выполнять по SSH.

ARM64 описывает архитектуру исполняемых файлов на стороне CPU, а вычислительная способность CUDA — целевую архитектуру кода для GPU. Это разные условия совместимости. Если file показывает AArch64, это подтверждает лишь правильный формат хостовой части, но не наличие пригодных для выполнения ядер внутри файла.

SASS — машинный код для целевого GPU, а PTX — промежуточное представление, которое драйвер может скомпилировать во время выполнения. Релизный пакет может содержать оба варианта, но проверять их нужно по отдельности. Успешный запуск в обычном режиме не означает, что проверены оба пути.

В ходе эксперимента исходные и исполняемые файлы создаются только в новом каталоге. Драйвер, системные инструменты сборки и код приложения не изменяются. Результаты примера предназначены для проверки совместимости, а не для оценки производительности.

2. Запрашивайте фактическую вычислительную способность, а не угадывайте номер архитектуры

Сначала создайте отдельный рабочий каталог и сохраните сведения о системе и инструментах сборки. Версия CUDA в выводе nvidia-smi отражает возможности драйвера, а не версию инструментария, к которому относится текущий nvcc. Поэтому нужны оба вывода.

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

Сохраните следующий код в probe.cu. Программа через CUDA Runtime запрашивает вычислительную способность устройства с индексом ноль. При ошибке она возвращает ненулевой код завершения, а не выводит идентификатор архитектуры, который можно было бы ошибочно использовать дальше.

#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

Идентификатор целевой архитектуры GB10 здесь намеренно не задан жёстко. Если в списке инструментария нет полученной цели sm_, сначала уточните, какие архитектуры он поддерживает. Нельзя просто подставить более старый номер и продолжить выпуск. Если программа проверки завершается с ошибкой, сначала устраните проблемы доступа к устройству или библиотек среды выполнения и лишь затем переходите к компиляции.

3. Явно включайте SASS и PTX в сборку и проверяйте результат

Сохраните следующий код в check.cu. Программа запускает всего один поток, который записывает целое число, затем копирует его на хост и проверяет результат. Ошибки проверяются после каждого вызова CUDA, чтобы сбой запуска ядра не сопровождался ложным сообщением об успехе.

#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;
}

В том же рабочем каталоге выполните команды, явно указав генерацию и нативного кода, и 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

Убедитесь, что список ELF и вывод SASS содержат нужную целевую архитектуру и ядро, а вывод PTX — директивы .version, .target и точку входа ядра. Пустые файлы, несовпадение архитектуры или отсутствие ожидаемого ядра должны блокировать приёмку. Одной проверки кодов завершения команд недостаточно.

Проверяйте итоговый файл, а не промежуточные объектные файлы

В реальном проекте могут быть этапы компоновки со статическими библиотеками, компоновки кода устройства или упаковки. Наличие PTX в объектном файле ещё не доказывает, что он сохранился в итоговом поставляемом файле. Повторите проверку для конечного исполняемого файла и фактически загружаемых разделяемых библиотек, а также зафиксируйте их хеши.

4. Проверяйте нативное выполнение и JIT-компиляцию по отдельности

Сначала запретите JIT-компиляцию PTX, чтобы текущий артефакт выполнялся только за счёт нативного кода устройства. Затем принудительно включите путь через PTX и отключите дисковый кеш компиляции. Для каждого испытания запускайте новый процесс, чтобы уже загруженные в процесс модули не влияли на результат.

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

В обоих журналах должно быть PASS, а оба кода завершения должны равняться нулю. Тестовые переменные действуют только для соответствующей команды. Не добавляйте их в глобальное окружение, оставляя рабочие процессы приложения в режиме принудительной JIT-компиляции.

Если нативное выполнение завершается с ошибкой, а PTX работает, в первую очередь перепроверьте целевую архитектуру нативного кода и финальную компоновку. Если нативный код работает, а PTX — нет, проверьте наличие PTX и поддержку его версии драйвером. Если не работают оба пути, исследуйте инструменты сборки, библиотеки среды выполнения и состояние устройства с учётом конкретных ошибок. Постоянная смена параметров компиляции не заменяет поиск причины.

5. Сделайте результаты проверки обязательным условием выпуска

Успешный запуск небольшой тестовой программы доказывает лишь то, что базовая цепочка сборки и это конкретное ядро работают в текущем окружении. Приложение может дополнительно загружать библиотеки расширений, выбирать разные ядра в зависимости от входных данных или генерировать код во время выполнения. Поэтому необходимы и тесты через реальные точки входа приложения, охватывающие основные типы данных и ветви выполнения.

Для каждого релиза рекомендуется сохранять четыре группы подтверждающих данных:

• Архитектура хоста, вывод драйвера, версия инструментария и фактические команды компиляции.
• Вычислительная способность GPU и результаты проверки SASS и PTX в итоговых файлах.
• Коды завершения и журналы для обоих путей выполнения, а также записи о проверке результатов приложения.
• Версия исходного кода, SHA-256 итоговых файлов и список динамических зависимостей.

Команда sha256sum ./check позволяет однозначно зафиксировать файл, прошедший эту приёмку. После повторной компоновки приёмку нужно провести заново, даже если имя файла не изменилось. При использовании контейнеров проверяйте инструменты сборки и итоговые файлы внутри фактически поставляемого образа, а испытания выполнения проводите в реальном окружении запуска.

PTX — это путь обеспечения совместимости, а не гарантия работы со всеми драйверами. Тем более один успешный тест не позволяет делать выводы о работе на других устройствах. Критерий допуска к выпуску следует формулировать явно: «Проверка пройдена с текущим драйвером, текущими артефактами и для охваченных тестами точек входа». При смене архитектуры или инструментария подтверждающие данные нужно собирать заново, а не использовать старые журналы.

Часто задаваемые вопросы