Как работают CUDA streams и как они помогают оверлапить compute и communication?

Краткий тезис

CUDA streams — это упорядоченные очереди операций (kernel, memory copy, events), выполняемые на GPU. Использование нескольких concurrent streams позволяет перекрывать (overlap) вычисления (kernel execution) с передачей данных (host-to-device / device-to-host), что значительно повышает утилизацию GPU и снижает общее время выполнения. Типичный прирост производительности — 20–40% для задач с интенсивным обменом данными.

1. Термин: CUDA stream

CUDA stream — это последовательность операций, которые выполняются на GPU в порядке их постановки. Все операции внутри одного stream гарантированно выполняются последовательно (в порядке FIFO). Разные streams могут выполняться конкурентно (concurrently) на одном GPU, если ресурсы (SM, memory bandwidth) позволяют.

Default stream (stream 0]]) — синхронный поток, который блокирует все остальные операции на GPU. Для оверлапа нужно использовать non-default streams (созданные явно через cudaStreamCreate).

Термин «Overlap» (перекрытие) — одновременное выполнение разных типов операций (например, kernel и memory copy) на GPU за счёт разделения аппаратных ресурсов (DMA engine, compute units).

2. Как работают CUDA streams: модель выполнения

GPU имеет несколько аппаратных очередей (hardware queues) для разных типов операций:

• DMA engine (для memory copy H2D/D2H)
• Kernel launch unit (для запуска ядер)
• Copy engine (для peer-to-peer transfers)

Когда вы запускаете операции в разных streams, драйвер может планировать их на разные аппаратные очереди, что позволяет выполнять их параллельно.

Пример без streams (последовательное выполнение):

# Псевдокод на Python с PyCUDA или cupy
cuda.memcpy_htod(d_a, h_a)   # H2D copy
cuda.memcpy_htod(d_b, h_b)   # H2D copy
kernel<<<grid, block>>>(d_a, d_b, d_c)  # kernel
cuda.memcpy_dtoh(h_c, d_c)   # D2H copy

Здесь всё выполняется последовательно: сначала копирование, потом kernel, потом обратное копирование. GPU простаивает во время копирования.

Пример с двумя streams (overlap):

stream1 = cuda.Stream()
stream2 = cuda.Stream()

# Stream 1: копирование данных A и запуск kernel A
cuda.memcpy_htod_async(d_a, h_a, stream=stream1)
kernel_a<<<grid, block, 0, stream1>>>(d_a, ...)

# Stream 2: копирование данных B и запуск kernel B
cuda.memcpy_htod_async(d_b, h_b, stream=stream2)
kernel_b<<<grid, block, 0, stream2>>>(d_b, ...)

# Дожидаемся завершения обоих streams
cuda.Device.synchronize()

Если DMA engine и compute units независимы, то H2D copy из stream1 может выполняться одновременно с kernel из stream2 (если kernel не требует данных, которые ещё копируются).

3. Типы оверлапа

Ключевое условие для оверлапа операции должны использовать разные аппаратные ресурсы. Если все операции — только memory copy, оверлап может быть ограничен одним DMA engine.

4. Аппаратные ограничения

• Количество DMA engines: на большинстве GPU (например, NVIDIA A100) есть два DMA engine (один для H2D, один для D2H). Это позволяет одновременно копировать в обе стороны.
• Количество concurrent kernels: зависит от архитектуры (Maxwell: 32, Pascal: 128, Volta: 128). Но реальное параллельное выполнение ограничено числом SM и ресурсами.
• Pinned memory: для асинхронных memory copy (async) необходимо использовать pinned (page-locked) host memory, иначе драйвер будет вынужден делать синхронное копирование через промежуточный буфер.

Термин «Pinned memory» — память на host, которая не может быть выгружена на диск, и её физические адреса фиксированы. Это позволяет GPU обращаться к ней напрямую через DMA без участия CPU.

5. Пример на C++/CUDA (реальный код)

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, N * sizeof(float));
cudaMalloc(&d_B, N * sizeof(float));
cudaMalloc(&d_C, N * sizeof(float));

// Асинхронное копирование в stream1
cudaMemcpyAsync(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice, stream1);
// Kernel в stream2 (не зависит от d_A)
myKernel<<<grid, block, 0, stream2>>>(d_B, d_C, N);
// Копирование результата из stream2 обратно
cudaMemcpyAsync(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost, stream2);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

Важно cudaMemcpyAsync требует pinned memory для host-указателей. Если h_A не pinned, вызов будет синхронным (драйвер сделает неявное копирование через pinned staging buffer, что может сломать оверлап).

6. Инструменты профилирования оверлапа

• NVIDIA Nsight Systems (nsys) — показывает временные диаграммы (timeline) операций на GPU. Можно увидеть, перекрываются ли kernel и memory copy.
• NVIDIA Nsight Compute (ncu) — детальный анализ kernel, но не timeline.
• CUDA Events — программные маркеры для измерения времени между операциями в разных streams.

Пример использования CUDA events для измерения оверлапа:

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream1);
// ... операции ...
cudaEventRecord(stop, stream2);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);

7. Best practices для оверлапа

1. Используйте pinned memory для всех host-буферов, участвующих в асинхронных копиях.
2. Разделяйте независимые данные на разные streams: например, данные для разных частей модели или разные батчи.
3. Избегайте синхронизаций между streams без необходимости (cudaStreamSynchronize, cudaDeviceSynchronize).
4. Используйте cudaStreamWaitEvent для создания зависимостей между streams (например, stream2 ждёт, пока stream1 закончит копирование).
5. Проверяйте оверлап через профилировщик — не все комбинации дают выигрыш из-за аппаратных ограничений.
6. Для pipeline parallelism (например, в Deep Learning) streams — основа: один stream для forward, другой для backward, третий для передачи данных на следующий ускоритель.

8. Типичные ошибки

• Использование default stream — все операции идут последовательно, оверлапа нет.
• Не-pinned память — cudaMemcpyAsync ведёт себя как синхронный, оверлап не работает.
• Слишком много streams — может привести к конкуренции за ресурсы и деградации производительности (обычно 2–4 streams оптимально).
• Зависимости данных — если kernel в stream2 использует данные, которые ещё копируются в stream1, будет race condition (нужно синхронизировать через events).

9. Связь с другими концепциями CUDA

Пет-проект для закрепления

Задача Реализовать простой векторный конвейер (vector addition pipeline), который перекрывает копирование данных с вычислениями.

Инструменты CUDA C++ (или PyCUDA), Nsight Systems для профилирования.

Шаги:

1. Создайте два больших массива на host (pinned) и device.
2. Разбейте массивы на 4 части (chunks).
3. Для каждого chunk: скопируйте H2D в stream1, запустите kernel сложения в stream2 (после завершения копирования через cudaStreamWaitEvent), скопируйте результат D2H в stream3.
4. Запустите конвейер: пока stream1 копирует chunk2, stream2 считает chunk1, stream3 копирует результат chunk0.
5. Сравните время с последовательной версией (один stream).
6. Постройте timeline в Nsight Systems, чтобы увидеть перекрытие.

Ожидаемый результат Ускорение в 1.5–2 раза по сравнению с последовательным выполнением. В профилировщике будут видны перекрывающиеся полосы kernel и memory copy.

Связь с другими вопросами

Навигация

• Предыдущий: 307
• Следующий: 309
• Индекс: 00. Индекс разборов