Проверка зависимостей потоков CUDA и времени жизни буферов на DGX Spark

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

1. Сначала определите зависимости данных, затем создавайте потоки

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

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

Входные данные готовы → копирование завершено → умножение завершено → начинается сложение → обратное копирование завершено → проверка на хосте.

В одном потоке операции выполняются в порядке этого потока; между двумя потоками зависимость нужно задать явно. Событие следует записать в потоке-производителе после ядра умножения. В потоке-потребителе сначала нужно поставить ожидание этого события, а затем отправить на выполнение ядро сложения.

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

2. Создайте минимальную программу без зависимости от потока по умолчанию

После оплаты на TcpNode откройте в панели управления данные для подключения и войдите в NVIDIA DGX OS по SSH. Сначала выполните uname -m, nvidia-smi и nvcc --version, сохранив соответственно архитектуру хоста, сведения об устройстве и драйвере, а также информацию об инструментарии компиляции. Архитектура должна быть aarch64, что соответствует ARM64. Версию компилятора нельзя определять по выводу драйвера.

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

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

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

__global__ void scale(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] *= 2.0f;
}

__global__ void add(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] += 1.0f;
}

int main() {
    constexpr int count = 4096;
    const size_t bytes = count * sizeof(float);
    float *input = nullptr, *output = nullptr, *device = nullptr;
    cudaStream_t producer, consumer;
    cudaEvent_t ready;

    check(cudaMallocHost(reinterpret_cast<void**>(&input), bytes));
    check(cudaMallocHost(reinterpret_cast<void**>(&output), bytes));
    check(cudaMalloc(reinterpret_cast<void**>(&device), bytes));
    check(cudaStreamCreateWithFlags(&producer, cudaStreamNonBlocking));
    check(cudaStreamCreateWithFlags(&consumer, cudaStreamNonBlocking));
    check(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));

    for (int i = 0; i < count; ++i) {
        input[i] = static_cast<float>(i);
        output[i] = -1.0f;
    }

    check(cudaMemcpyAsync(device, input, bytes,
                          cudaMemcpyHostToDevice, producer));
    scale<<<(count + 255) / 256, 256, 0, producer>>>(device, count);
    check(cudaGetLastError());
    check(cudaEventRecord(ready, producer));

    check(cudaStreamWaitEvent(consumer, ready, 0));
    add<<<(count + 255) / 256, 256, 0, consumer>>>(device, count);
    check(cudaGetLastError());
    check(cudaMemcpyAsync(output, device, bytes,
                          cudaMemcpyDeviceToHost, consumer));
    check(cudaStreamSynchronize(consumer));

    int failures = 0;
    for (int i = 0; i < count; ++i) {
        float expected = 2.0f * static_cast<float>(i) + 1.0f;
        if (output[i] != expected) ++failures;
    }
    std::printf("elements=%d mismatches=%d\n", count, failures);

    check(cudaEventDestroy(ready));
    check(cudaStreamDestroy(consumer));
    check(cudaStreamDestroy(producer));
    check(cudaFree(device));
    check(cudaFreeHost(output));
    check(cudaFreeHost(input));
    return failures == 0 ? 0 : 2;
}

3. Проверяйте и вывод, и код завершения

Скомпилируйте и запустите программу в том же каталоге:

nvcc -O2 stream_contract.cu -o stream_contract
./stream_contract
printf 'exit=%s\n' "$?"

Успешный результат — elements=4096 mismatches=0 и нулевой код завершения. При ошибке вызова API или синхронизации программа возвращает единицу, при несовпадении численных результатов — двойку. Если проверять только отсутствие аварийного завершения процесса, можно пропустить вычисления, которые завершились, но дали неверные данные.

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

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

Убедитесь, что обработка ошибок действительно работает

Временно замените последнее 1.0f в выражении для ожидаемого значения на 2.0f. Программа должна сообщить о несовпадениях и вернуть двойку. Затем восстановите исходное выражение. Такая проверка подтверждает, что механизм контроля результата работает. Не используйте в качестве негативного теста требование «после удаления ожидания обязательно должна возникнуть ошибка»: проявление гонки зависит от планирования, и программа с отсутствующей зависимостью тоже может какое-то время проходить проверку.

4. Увяжите границы событий со временем жизни буферов

Возврат из cudaMemcpyAsync обычно означает лишь, что вызов отправлен на выполнение, а не что входные данные уже можно перезаписывать. В примере входной буфер остаётся неизменным с момента отправки копирования до завершения синхронизации потока-потребителя. Выходной буфер также читается только после синхронизации потока, в котором выполняется обратное копирование. Это консервативная стратегия, но у неё чёткие границы, поэтому она подходит как исходный вариант для проверки корректности.

Закреплённая память помогает выполнить условия асинхронной передачи, но не служит уведомлением о её завершении. Буфер на устройстве точно так же нельзя освобождать, пока другой поток продолжает к нему обращаться. Наличие 128 ГБ унифицированной памяти не означает, что нужно закреплять большие объёмы памяти хоста целиком. Размер пакета следует ограничивать фактическим рабочим набором данных.

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

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

5. При диагностике постепенно сужайте область синхронизации

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

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

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

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

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