Как работает asynchronous execution на Hopper (copy engine vs compute)?

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

Архитектура NVIDIA Hopper (H100) вводит два выделенных copy engine для асинхронного копирования данных между хостом и устройством (H2D/D2H), которые могут работать параллельно с compute engine (исполнением ядер). Дополнительно Tensor Memory Accelerator (TMA) обеспечивает асинхронное копирование между глобальной и разделяемой памятью (D2D). Это позволяет перекрывать (overlap) передачу данных с вычислениями, скрывая latency и повышая утилизацию GPU. Для LLM-инференса и агентных RAG-систем такой механизм критичен: можно prefetch’ить следующий батч или контекст, пока обрабатывается текущий.

1. Термины и контекст

Asynchronous execution — выполнение операций (копирование, вычисления) без блокировки основного потока. В CUDA это реализуется через streams и events.

Copy engine — аппаратный блок GPU, специализирующийся на копировании данных между памятью хоста и устройства (H2D/D2H) или между устройствами (P2P). В H100 их два: один для H2D, другой для D2H.

Compute engine — блок, исполняющий CUDA-ядра (kernel). В H100 это набор SM (Streaming Multiprocessors) с поддержкой Tensor Cores.

TMA (Tensor Memory Accelerator) — новый блок в Hopper для асинхронного копирования данных из глобальной памяти в разделяемую (shared memory) и обратно. Поддерживает многомерные массивы, выравнивание и трансформации.

Overlap — техника, при которой передача данных и вычисления выполняются одновременно, скрывая задержки (latency) копирования.

2. Архитектура H100: два copy engine + compute

В H100 реализована асимметричная схема:

Все три engine могут работать параллельно, если они используют разные streams и не конфликтуют по ресурсам памяти.

Пример параллелизма:

• Stream A: копирование следующего батча с хоста на GPU (H2D)
• Stream B: выполнение ядра для текущего батча
• Stream C: копирование результатов предыдущего батча на хост (D2H)

Таким образом, за время одного kernel launch можно полностью скрыть latency передачи данных.

3. Как работает copy engine: DMA, streams, events

Copy engine использует DMA (Direct Memory Access) — прямой доступ к памяти без участия ядер. Операции запускаются через cudaMemcpyAsync() и выполняются в указанном stream.

Ключевые моменты:

• Копирование в одном stream блокирует последующие операции в том же stream, но не в других.
• Для синхронизации используются events: cudaEventRecord() и cudaStreamWaitEvent().
• Copy engine может работать только с памятью, закреплённой (pinned) на хосте (cudaHostAlloc или cudaMallocHost). Без pinned memory копирование идёт через промежуточный буфер и не может быть асинхронным.

Пример кода:

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

// Асинхронное копирование H2D в stream1
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);

// Запуск ядра в stream2 (параллельно копированию)
kernel<<<grid, block, 0, stream2>>>(d_other);

// Синхронизация: ждём завершения копирования
cudaStreamSynchronize(stream1);

4. TMA (Tensor Memory Accelerator) — асинхронное D2D

TMA — это аппаратный блок, который может копировать данные из глобальной памяти в разделяемую (shared memory) асинхронно, без участия потоков (warps). Он поддерживает:

• Многомерные тензоры (1D, 2D, 3D, 4D) с произвольными шагами (strides).
• Трансформации (например, транспонирование, swizzle).
• Выравнивание по границам кэш-линий.
• Асинхронность: копирование запускается одной инструкцией (cp.async.bulk), а завершение проверяется через cp.async.bulk.wait_group.

Сравнение TMA и обычного copy engine:

TMA особенно полезен для операций, требующих загрузки данных в shared memory перед вычислением (например, матричное умножение, attention).

5. Пример: асинхронный pipeline для LLM inference

Рассмотрим батч-обработку запросов в LLM (например, для агентного RAG, где каждый запрос требует поиска и генерации).

Без асинхронности:

1. Копировать входные данные (prompt) на GPU.
2. Выполнить forward pass.
3. Копировать результаты на хост.
4. Повторить для следующего батча.

Всё последовательно → GPU простаивает во время копирования.

С асинхронностью (overlap):

cudaStream_t computeStream, h2dStream, d2hStream;
// Инициализация streams

for (int i = 0; i < num_batches; ++i) {
    // Prefetch: копируем следующий батч в h2dStream
    if (i + 1 < num_batches)
        cudaMemcpyAsync(d_input_next, h_input[i+1], size, cudaMemcpyHostToDevice, h2dStream);

    // Вычисляем текущий батч в computeStream
    launch_llm_kernel(d_input[i], d_output[i], computeStream);

    // Копируем результат предыдущего батча на хост в d2hStream
    if (i > 0)
        cudaMemcpyAsync(h_output[i-1], d_output[i-1], size, cudaMemcpyDeviceToHost, d2hStream);

    // Синхронизация перед следующей итерацией (можно через events)
    cudaStreamSynchronize(computeStream);
}

В результате latency копирования скрывается за вычислениями. Для LLM с большими батчами (например, 4096 токенов) это даёт прирост throughput до 20–30%.

6. Применение в Agentic RAG

Агентные RAG-системы часто выполняют несколько параллельных действий:

• Поиск в векторной БД (CPU-bound).
• Загрузка документов в память GPU (H2D).
• Генерация ответа LLM (GPU compute).

Используя асинхронные copy engine, можно:

• Пока LLM генерирует ответ для одного запроса, prefetch’ить контекст для следующего.
• Если агент вызывает несколько инструментов (например, поиск + суммаризация), загружать их результаты на GPU параллельно.

Пример архитектуры:

Поток 1 (CPU): поиск документов → формирование промпта → pinned buffer
Поток 2 (GPU H2D): асинхронное копирование промпта на GPU
Поток 3 (GPU compute): инференс LLM
Поток 4 (GPU D2H): копирование ответа на хост

Все потоки работают одновременно, утилизация GPU близка к 100%.

7. Ограничения и best practices

• Pinned memory обязательно: без неё cudaMemcpyAsync ведёт себя как синхронный.
• Не все копирования можно перекрыть: если kernel использует ту же память, что копируется, нужна синхронизация.
• TMA требует поддержки в коде: нужно использовать специальные инструкции PTX или библиотеки (CUTLASS, CuTe).
• Число streams ограничено: обычно 2–4 достаточно.
• Профилирование: используйте nsys (NVIDIA Nsight Systems) для визуализации overlap.

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

Задача: Реализовать асинхронный pipeline для batch-инференса небольшой LLM (например, GPT-2) с overlapping копирования и вычислений.

Инструменты: Python + PyTorch + CUDA streams (через torch.cuda.Stream), или C++/CUDA.

Шаги:

1. Создать три CUDA streams: h2d, compute, d2h.
2. Подготовить pinned memory на хосте для входных данных (использовать torch.utils.cpp_extension или cudaHostAlloc).
3. В цикле по батчам:
   • Запустить асинхронное копирование следующего батча в h2d stream.
   • Запустить kernel (forward) в compute stream.
   • Запустить копирование предыдущего результата в d2h stream.
4. Измерить время выполнения с overlap и без.
5. Визуализировать timeline через Nsight Systems.

Ожидаемый результат: Ускорение в 1.5–2 раза по сравнению с последовательным выполнением для 4–8 батчей.

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

10. Навигация

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

Навигация

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