Как вы читаете профиль Nsight Systems для поиска bottlenecks в vLLM?

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

Nsight Systems — это системный профайлер от NVIDIA, который позволяет видеть временную шкалу выполнения GPU-ядер, CPU-активности, передач данных и вызовов CUDA API. При профилировании vLLM (библиотеки для эффективного инференса LLM) мы ищем характерные паттерны: длинные промежутки между ядрами (CPU overhead|launch overhead|launch overhead), периоды простоя GPU (memory stalls), большие передачи данных через PCIe и задержки в CUDA API. Анализ этих метрик помогает выявить узкие места и оптимизировать производительность инференса.

1. Термины: Nsight Systems, vLLM, bottleneck

Nsight Systems — инструмент для профилирования приложений, использующих GPU. Он записывает события на CPU и GPU, строит временную диаграмму (timeline) и собирает статистику по вызовам CUDA, ядрам, передачам памяти.

vLLM — open-source библиотека для инференса больших языковых моделей, использующая PagedAttention и batching|непрерывное пакетирование (batching|continuous batching) для достижения высокой пропускной способности.

Bottleneck (узкое место) — компонент системы, который ограничивает общую производительность. В контексте GPU-инференса это может быть:

• CPU launch overhead — время, которое CPU тратит на запуск GPU-ядер (kernel launches).
• Memory stalls — периоды, когда GPU ждёт данные из памяти.
• PCIe transfers — копирование данных между CPU и GPU.
• CUDA API latency — задержки при вызове функций CUDA (например, cudaMemcpy, cudaLaunchKernel).

2. Основные метрики в профиле Nsight Systems

При открытии профиля vLLM в Nsight Systems GUI мы видим несколько дорожек (rows). Ключевые:

Основные метрики для поиска bottlenecks:

• GPU Utilization — процент времени, когда GPU занят вычислениями. Низкое значение (< 50%) указывает на CPU launch overhead или memory stalls.
• Kernel Duration — время выполнения каждого ядра. Сравниваем с теоретическим минимумом (roofline model).
• Gap между kernels — время между окончанием одного ядра и началом следующего. Если gaps велики (> 10 мкс), это CPU launch overhead.
• CUDA API Call Latency — время, затраченное на cudaLaunchKernel. Обычно 5–20 мкс, но при большом количестве ядер может накапливаться.
• Memory Throughput — скорость чтения/записи в глобальную память GPU. Если ниже пиковой пропускной способности, возможны memory stalls.

3. Как читать timeline: визуальный анализ

Откроем профиль vLLM в Nsight Systems. Типичная картина для инференса LLM:

1. Prefill-фаза (обработка входного промпта): несколько длинных ядер (матричное умножение, attention). GPU загружен почти 100%.
2. Decode-фаза (генерация токенов): множество коротких ядер (по одному на каждый токен в batch). Здесь часто возникают bottlenecks.

На timeline decode-фазы мы видим:

• Зелёные блоки (CUDA HW) — ядра. Они короткие (единицы микросекунд).
• Между блоками — пустые промежутки (gaps). Если gaps > 5–10 мкс, это CPU launch overhead.
• На дорожке CUDA API — частые вызовы cudaLaunchKernel (оранжевые точки). Если они занимают значительное время, это overhead.
• На дорожке Memory — редкие передачи данных (например, копирование KV cache при переключении batch).

Пример паттерна bottleneck

[CPU]  cudaLaunchKernel (5 мкс)
[GPU]  <<<kernel A>>> (2 мкс)
[CPU]  cudaLaunchKernel (5 мкс)
[GPU]  <<<kernel B>>> (2 мкс)
...

Здесь GPU простаивает 5 мкс между ядрами, а работает всего 2 мкс. Utilization GPU = 2/(2+5) ≈ 28%. Это классический CPU launch overhead bottleneck.

4. Конкретные паттерны bottlenecks в vLLM

4.1 CPU launch overhead (kernel launch bound)

• Причина: vLLM запускает много мелких ядер (attention, MLP, layer norm) для каждого токена. CPU не успевает их подавать.
• Признаки в профиле: частые gaps между ядрами > 5 мкс, много вызовов cudaLaunchKernel на дорожке CUDA API, низкая GPU utilization.
• Решение: увеличить batch size (continuous batching), использовать CUDA graphs (захват последовательности ядер в один граф), объединять ядра (kernel fusion).

4.2 Memory stalls (memory bound)

• Причина: ядра упираются в пропускную способность памяти (например, attention считывает KV cache).
• Признаки: ядра имеют высокую длительность, но низкую вычислительную загрузку (compute utilization). В Nsight Systems можно посмотреть Memory Throughput — если она близка к пиковой, а compute низкая, то memory bound.
• Решение: использовать более эффективные форматы (FP8, INT4), оптимизировать layout памяти (PagedAttention уже помогает), уменьшить размер KV cache.

4.3 PCIe transfers (data transfer bound)

• Причина: копирование входных данных или результатов между CPU и GPU.
• Признаки: большие блоки cudaMemcpy на дорожке Memory, длительностью > 100 мкс.
• Решение: асинхронные копирования, перекрытие с вычислениями (streams), пакетная обработка.

4.4 CUDA API call latency

• Причина: частые вызовы CUDA API из CPU (например, cudaSetDevice, cudaStreamSynchronize).
• Признаки: на дорожке CUDA API много вызовов с заметной длительностью (десятки микросекунд).
• Решение: минимизировать синхронизации, использовать CUDA streams для параллелизма.

5. Инструменты для детального анализа

Nsight Systems предоставляет несколько способов анализа:

• Timeline View — визуальный просмотр событий. Можно увеличивать масштаб, измерять интервалы.
• Statistics View — сводка по вызовам CUDA API, ядрам, передачам памяти. Показывает общее время, количество вызовов, среднюю длительность.
• GPU Utilization — график загрузки GPU во времени.
• Roofline Analysis (через Nsight Compute) — связь вычислительной интенсивности с пропускной способностью памяти.

Для vLLM полезно добавить NVTX markers (пользовательские метки) в код, чтобы размечать фазы: prefill, decode, scheduling. Это упрощает навигацию по профилю.

Пример добавления NVTX в vLLM (псевдокод):

import nvtx

@nvtx.annotate("prefill", color="green")
def prefill_step(batch):
    ...

@nvtx.annotate("decode", color="blue")
def decode_step(batch):
    ...

6. Пример анализа профиля vLLM

Допустим, мы запустили vLLM с моделью LLaMA-7B, batch size = 1, и получили профиль. На timeline decode-фазы видим:

• 100 ядер (attention, MLP) за 1 мс.
• Средний gap между ядрами: 8 мкс.
• Средняя длительность ядра: 3 мкс.
• GPU utilization: 3/(3+8) ≈ 27%.
• На дорожке CUDA API: 100 вызовов cudaLaunchKernel по 6 мкс каждый → 600 мкс overhead.

Вывод: bottleneck — CPU launch overhead. Решение: увеличить batch size до 8, тогда количество ядер на токен останется тем же, но overhead распределится на большее количество токенов, и GPU utilization вырастет.

После увеличения batch size до 8:

• Ядер стало 800 (на 8 токенов), но gaps сократились до 2 мкс (CPU успевает подготовить следующее ядро, пока GPU занято другими токенами).
• GPU utilization: 3/(3+2) = 60%.

7. Метрики для количественной оценки

При чтении профиля полезно вычислить следующие метрики:

• Kernel Launch Overhead Ratio = (total gap time) / (total kernel time + total gap time). Если > 0.3, это проблема.
• GPU Active Time = сумма длительностей всех ядер.
• CPU Active Time = сумма длительностей вызовов CUDA API.
• Memory Transfer Time = сумма длительностей cudaMemcpy.
• Throughput (токенов/сек) — внешняя метрика, коррелирующая с профилем.

В Nsight Systems можно экспортировать данные в CSV и обработать скриптом:

import pandas as pd

df = pd.read_csv('profile.csv')
kernel_events = df[df['Name'].str.contains('kernel')]
total_kernel_time = kernel_events['Duration'].sum()
gaps = ... # вычислить разницу между концом одного ядра и началом следующего
total_gap_time = gaps.sum()
overhead_ratio = total_gap_time / (total_kernel_time + total_gap_time)
print(f"Overhead ratio: {overhead_ratio:.2f}")

8. Связь с архитектурой vLLM

Bottlenecks в vLLM напрямую связаны с его архитектурными решениями:

• PagedAttention — уменьшает memory stalls за счёт эффективного управления KV cache, но может увеличить количество мелких ядер (page walks).
• Continuous batching — увеличивает batch size, снижая CPU launch overhead, но требует более сложного планировщика.
• Tensor parallelism — распределяет модель на несколько GPU, но добавляет коммуникационные overhead (all-reduce). В профиле это видно как дополнительные ядра и передачи между GPU.

При профилировании с несколькими GPU нужно смотреть на дорожки для каждого устройства и искать дисбаланс.

9. Оптимизации на основе профиля

После идентификации bottleneck можно применить:

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

Задача: Профилировать vLLM с моделью LLaMA-2-7B на одном GPU, найти bottleneck и применить CUDA graphs для ускорения.

Инструменты:

• vLLM (установить pip install vllm)
• Nsight Systems (скачать с NVIDIA)
• Python, CUDA toolkit

Шаги:

1. Установить vLLM и запустить инференс с batch size = 1 (например, через vllm.entrypoints.openai.api_server).
2. Запустить профилирование: nsys profile -o vllm_profile python -c "from vllm import LLM; llm = LLM('meta-llama/Llama-2-7b-hf'); llm.generate(['Hello'])".
3. Открыть профиль в Nsight Systems GUI, найти decode-фазу, измерить gaps и kernel durations.
4. Вычислить overhead ratio.
5. Включить CUDA graphs в vLLM (параметр --enforce-eager отключить, по умолчанию graphs включены для decode). Сравнить профиль.
6. Построить график GPU utilization до и после.

Ожидаемый результат: После включения CUDA graphs overhead ratio снизится с >0.3 до <0.1, GPU utilization вырастет, throughput увеличится.

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

Навигация

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