Как найти выход за границы памяти CUDA на DGX Spark

Ядро CUDA запустилось, но ошибка появилась только при копировании результатов обратно на хост. При уменьшении входных данных та же программа работает. При таких симптомах сначала стоит проверить обращения к памяти устройства, а не менять драйвер. На DGX Spark небольшой автономный пример позволяет пройти всю цепочку диагностики: намеренно вызвать выход за границы, найти строку исходного кода, исправить индекс и убедиться, что обнаружение ошибок по-прежнему работает. Для этого не нужны веса модели, а работающий сервис в проверке не участвует.

1. Зафиксируйте окружение и изолируйте тест

DGX Spark от TcpNode использует NVIDIA GB10 Grace Blackwell, архитектуру ARM64 и NVIDIA DGX OS и оснащён 128 ГБ унифицированной памяти LPDDR5x. Общая память CPU и GPU не означает, что по указателю можно обращаться за границы выделенной области. В примере ниже память устройства выделяется через cudaMalloc, поэтому диапазон допустимых адресов определён явно.

После оплаты найдите данные для подключения в консоли и войдите по SSH. Создайте отдельный рабочий каталог и запускайте тест от обычного пользователя. Не внедряйте ошибки в процесс инференса, который обслуживает запросы.

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

В качестве архитектуры должно отображаться aarch64. Сохраните версии инструментов и сведения о драйвере, а повторные проверки выполняйте в том же окружении. Предустановленный стек CUDA не гарантирует, что все инструменты доступны через текущий PATH. Если команда не найдена, сначала проверьте установленный набор инструментов и пути установки, а не заменяйте сразу весь стек драйверов.

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

2. Создайте воспроизводимый пример с одной переменной

Сохраните следующий код в repro.cu. Он выделяет память только для одного целого числа, поэтому единственный допустимый индекс — ноль. Позиция записи задаётся при компиляции через TEST_INDEX. Так ошибочная и исправленная версии используют один исходный код, без одновременного изменения входных данных, размера выделенной памяти и пути выполнения.

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

Скомпилируйте две программы отдельно:

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

Параметр -arch=native компилирует код для GPU, доступного компилятору в данный момент. Он подходит для локальной проверки, но не для сборки универсального исполняемого файла для разных устройств. Если компилятор не поддерживает текущее устройство, сначала проверьте совместимость набора инструментов, а не подставляйте произвольные параметры другой архитектуры. -lineinfo сохраняет информацию о строках исходного кода устройства. Для первичного поиска ошибки не требуется полный режим отладки кода устройства, который может заметно изменить характер выполнения.

Зачем проверять и запуск, и синхронизацию

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

3. Получите код завершения, пригодный для автоматической проверки

Сначала запустите ошибочную версию, сохранив стандартный вывод и журнал диагностики в отдельные файлы:

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

Ожидается сообщение о недопустимой записи в глобальную память с указанием операции записи в write_value. Обычно отчёт также содержит размер обращения, координаты потока и положение адреса относительно выделенной области. Формат зависит от версии инструмента, поэтому конкретный фрагмент текста не должен быть единственным критерием проверки.

Параметр --error-exitcode 86 позволяет учитывать обнаруженные ошибки в результате автоматического запуска. Однако если сама проверяемая программа завершилась с ошибкой, итоговым может оказаться другой ненулевой код, возвращённый программой. Поэтому критерий приёмки — «ненулевой код завершения и отчёт, указывающий на ожидаемый выход за границы», а не обязательный возврат 86 при любом сбое.

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

4. Исправляйте индексацию по границам памяти, а не маскируйте симптомы

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

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

Для одномерной задачи обычно следует проверять index < count до обращения к памяти, а также убеждаться, что целочисленный тип, используемый при вычислении индекса, вмещает произведение. В ядрах с совместной работой через разделяемую память нельзя произвольно завершать часть потоков досрочно, если остальные затем доходят до барьера на уровне блока: это может создать новые проблемы синхронизации.

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

5. Подтвердите исправление положительным и отрицательным тестами

Запустите исправленную версию:

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

В этом примере должны одновременно выполняться три условия: нулевой код завершения, отсутствие ошибок обращения к памяти в отчёте и вывод result=7. Одного отсутствия сообщений об ошибках недостаточно: тест может выглядеть успешным, даже если ядро не выполнялось или результат на самом деле не был записан.

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

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

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