CUDA Streams: параллельное выполнение ядер и асинхронные операции

Введение

По умолчанию операции CUDA (запуск ядра, копирование памяти) выполняются в «default stream» — это означает, что операции в этом потоке исполнения GPU сериализуются друг за другом, даже если они логически независимы. CUDA Streams позволяют явно организовать несколько независимых последовательностей операций, которые могут выполняться параллельно или перекрываться (например, копирование данных одного блока, пока вычисляется ядро для другого) — это называется overlap вычислений и передачи данных. В статье разберём, как использовать потоки CUDA для повышения утилизации GPU.

Концепция

cudaStream_t представляет независимую последовательность операций на GPU. Операции внутри одного потока выполняются строго по порядку, но операции в разных потоках могут перекрываться, если для этого достаточно ресурсов (в первую очередь — отдельных DMA-каналов для асинхронного копирования памяти, доступных, если память на host выделена как pinned/page-locked через cudaMallocHost или cudaHostAlloc). Использование нескольких потоков особенно эффективно, когда нужно обработать данные по частям (chunks), перекрывая копирование следующей части с вычислением текущей.

Пример кода

#include <cuda_runtime.h>
#include <iostream>

__global__ void processChunk(float *data, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f + 1.0f;
    }
}

int main()
{
    const int numChunks = 4;
    const int chunkSize = 1 << 20; // 1M элементов на чанк
    const size_t chunkBytes = chunkSize * sizeof(float);

    // Pinned memory необходима для асинхронного копирования
    float *hostData;
    cudaMallocHost(&hostData, numChunks * chunkBytes);
    for (int i = 0; i < numChunks * chunkSize; ++i) {
        hostData[i] = static_cast<float>(i % 100);
    }

    float *deviceData[numChunks];
    cudaStream_t streams[numChunks];

    for (int i = 0; i < numChunks; ++i) {
        cudaMalloc(&deviceData[i], chunkBytes);
        cudaStreamCreate(&streams[i]);
    }

    const int threadsPerBlock = 256;
    const int blocksPerGrid = (chunkSize + threadsPerBlock - 1) / threadsPerBlock;

    // Каждый чанк обрабатывается в своём потоке: копирование на GPU,
    // вычисление и копирование обратно перекрываются между потоками
    for (int i = 0; i < numChunks; ++i) {
        float *hostChunk = hostData + i * chunkSize;
        cudaMemcpyAsync(deviceData[i], hostChunk, chunkBytes,
                         cudaMemcpyHostToDevice, streams[i]);
        processChunk<<<blocksPerGrid, threadsPerBlock, 0, streams[i]>>>(deviceData[i], chunkSize);
        cudaMemcpyAsync(hostChunk, deviceData[i], chunkBytes,
                         cudaMemcpyDeviceToHost, streams[i]);
    }

    for (int i = 0; i < numChunks; ++i) {
        cudaStreamSynchronize(streams[i]);
        cudaStreamDestroy(streams[i]);
        cudaFree(deviceData[i]);
    }

    std::cout << "Обработка завершена, пример значения: " << hostData[0] << std::endl;
    cudaFreeHost(hostData);
    return 0;
}

Пояснения к коду

Память на host выделяется через cudaMallocHost (pinned memory) вместо обычного malloc/new — это необходимое условие для использования действительно асинхронного копирования cudaMemcpyAsync: обычная (pageable) память сначала копируется драйвером во временный pinned-буфер, что устраняет выигрыш от асинхронности.

Каждый из четырёх чанков данных обрабатывается в своём собственном cudaStream_t: копирование на устройство, вычисление ядра и копирование обратно — все три операции для чанка i помещены в поток streams[i]. Поскольку операции в разных потоках не имеют зависимостей друг от друга, драйвер CUDA может перекрывать копирование данных одного чанка с вычислением другого, повышая общую утилизацию GPU по сравнению с полностью последовательной обработкой в default stream.

Подводные камни

• Использование default stream (поток 0) вперемешку с явными потоками может неявно сериализовать операции — в legacy-режиме default stream синхронизируется со всеми остальными потоками, что ломает ожидаемый параллелизм. Для полностью независимых потоков стоит использовать cudaStreamNonBlocking при создании или per-thread default stream (флаг компиляции --default-stream per-thread).
• Память, выделенная не как pinned, делает cudaMemcpyAsync фактически синхронным на уровне реального перекрытия операций — копирование всё равно выполнится, но без выигрыша от асинхронности, и это легко не заметить, поскольку код формально компилируется и работает корректно, просто медленнее.
• Избыточное число потоков (больше, чем GPU реально может параллельно использовать с учётом числа DMA-каналов и SM) не даёт дополнительного выигрыша и усложняет код — нужно соотносить число потоков с реальными возможностями конкретного GPU.
• Забытая синхронизация перед использованием результата на host. Если не вызвать cudaStreamSynchronize() (или cudaDeviceSynchronize()) перед чтением hostData на CPU, есть риск прочитать данные до завершения асинхронного копирования с устройства — поскольку асинхронные вызовы возвращают управление немедленно, не дожидаясь реального завершения операции на GPU.