Как вы профилируете GPU utilization для LLM serving (nsys, ncu, nvprof)?

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

Профилирование GPU utilization для LLM serving — это процесс измерения и анализа использования вычислительных ресурсов GPU (ядер CUDA, памяти, пропускной способности) при инференсе LLM. Основные инструменты: NVIDIA Nsight Systems (nsys) для системного профилирования, NVIDIA Nsight Compute (ncu) для анализа отдельных ядер, и устаревший nvprof. Ключевые метрики: SM occupancy, warp stall reasons, memory bandwidth utilization, которые помогают выявить узкие места (bottleneck) и оптимизировать serving (например, увеличить batch size, включить FlashAttention или parallelism|tensor parallelism).

1. Введение: зачем профилировать GPU при LLM serving?

LLM serving — это инференс больших языковых моделей (например, LLaMA, GPT) на GPU. Без профилирования невозможно понять, эффективно ли используются ресурсы GPU. Типичные проблемы:

• Низкая occupancy — мало активных warp'ов на SM, GPU простаивает.
• Memory-bound — узким местом является пропускная способность памяти (HBM), а не вычисления.
• Warp stall — warp'ы ждут данные из памяти или синхронизации.

Профилирование даёт численные метрики, на основе которых принимаются решения: увеличить batch size, включить FlashAttention, использовать tensor parallelism или quantization.

2. Инструменты профилирования: обзор

Рекомендация используйте nsys для общего понимания, затем ncu для углублённого анализа горячих ядер.

3. NVIDIA Nsight Systems (nsys) — системное профилирование

nsys записывает временную шкалу (timeline) выполнения: вызовы CUDA API, запуск ядер, копирование памяти, активность CPU. Позволяет увидеть:

• CPU-GPU overlap — перекрывается ли подготовка данных на CPU с вычислениями на GPU.
• Idle time — периоды, когда GPU простаивает (например, из-за ожидания данных).
• Memory transfers — объём и частота копирований между CPU и GPU.

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

nsys profile -o my_profile -t cuda,nvtx,osrt python serve.py

• -t cuda,nvtx,osrt — трассировка CUDA, пользовательских меток (NVTX) и системных вызовов.
• Результат — файл .nsys-rep, открываемый в Nsight Systems GUI.

Ключевые метрики из nsys

• GPU utilization — процент времени, когда GPU занят.
• Memory bandwidth utilization — доля используемой пропускной способности HBM.
• Kernel launch overhead — время на запуск ядер.

4. NVIDIA Nsight Compute (ncu) — kernel-level анализ

ncu профилирует одно или несколько CUDA-ядер (kernels) и выдаёт детальную статистику: occupancy, warp stall reasons, использование регистров, shared memory, инструкции.

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

ncu --set full -o kernel_profile python serve.py

• --set full — собирает все доступные метрики.
• Можно указать конкретное ядро по имени: --kernel-name "attention".

Основные метрики ncu

• SM occupancy — отношение активных warp'ов к максимально возможному на SM.
• Warp stall reasons — причины, по которым warp'ы простаивают (см. таблицу ниже).
• Memory bandwidth utilization — процент использования пропускной способности HBM.
• Compute utilization — процент использования вычислительных блоков (FP32/FP16/Tensor Cores).

Таблица: Warp stall reasons (основные)

5. nvprof (устаревший)

nvprof — предшественник nsys и ncu. Поддерживает как системную трассировку (nvprof --profile-from-start off), так и kernel-метрики (nvprof --metrics all). Однако NVIDIA рекомендует переходить на nsys/ncu. Для совместимости со старыми проектами можно использовать, но новые разработки должны опираться на современные инструменты.

6. Ключевые метрики: подробный разбор

6.1 SM Occupancy

Occupancy = (количество активных warp'ов на SM) / (максимальное количество warp'ов на SM). Высокая occupancy (обычно > 80%) означает, что GPU может скрывать задержки памяти за счёт переключения между warp'ами. Низкая occupancy (< 50%) — признак того, что ядро ограничено регистрами или shared memory.

Формула (упрощённо):

Occupancy = (warp_limit - warp_stall) / warp_limit

где warp_limit — максимальное число warp'ов на SM (зависит от архитектуры: 64 для A100, 48 для V100).

Как интерпретировать

• Высокая occupancy — хорошо для скрытия latency, но не гарантирует высокую производительность (может быть memory-bound).
• Низкая occupancy — нужно уменьшить использование регистров или shared memory (например, через __launch_bounds__).

6.2 Warp Stall Reasons

Показывает, почему warp'ы не выполняют инструкции. Если доминирует waiting — ядро memory-bound. Если math_pipe — compute-bound. Если not_selected — occupancy высокая, но warp'ы конкурируют за ресурсы.

6.3 Memory Bandwidth Utilization

Доля от пиковой пропускной способности HBM (например, 2 TB/s для A100). Если utilisation < 50% при высокой occupancy — возможно, плохой паттерн доступа (некоалесцентный). Если utilisation > 90% — ядро memory-bound.

7. Интерпретация и поиск узких мест

Типичный сценарий для LLM serving (например, decoder-only модель):

1. Запуск nsys видим, что GPU utilisation около 60%, есть длинные промежутки между запусками ядер — возможно, bottleneck на CPU (подготовка данных, scheduling).
2. Запуск ncu на ядре attention occupancy 40%, warp stall reasons — waiting 70% → memory-bound.
3. Решение включить FlashAttention (уменьшает чтение/запись HBM) или увеличить batch size (повышает occupancy).

Таблица: Типичные проблемы и решения

8. Пример профилирования LLM serving с vLLM

Сценарий Serving модели LLaMA-2-7B на одной A100 с vLLM. Нагрузка — 100 запросов.

Шаги:

1. Запускаем vLLM с --max-num-seqs 256.
2. В отдельном терминале запускаем nsys profile -o vllm_profile -t cuda,nvtx python -m vllm.entrypoints.openai.api_server --model meta-llama/Llama-2-7b-hf.
3. Отправляем запросы (например, с помощью wrk).
4. После завершения открываем .nsys-rep в Nsight Systems GUI.
5. Выбираем горячее ядро (например, flash_attn_v2_fwd) и запускаем ncu --kernel-name flash_attn_v2_fwd --set full -o flash_attn_profile python ....

Ожидаемые результаты

• nsys: GPU utilisation ~70%, есть небольшие простои между запросами.
• ncu: occupancy 60%, warp stall reasons: waiting 50%, not_selected 30%, math_pipe 20% → memory-bound с умеренной occupancy.
• Рекомендация: увеличить batch size до 512 (если позволяет память) или включить tensor parallelism (2 GPU).

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

После профилирования можно применить следующие техники:

• FlashAttention — уменьшает количество операций чтения/записи HBM, снижает memory-bound.
• Tensor parallelism — распределяет веса между несколькими GPU, уменьшает нагрузку на HBM одной GPU.
• Quantization (FP16, INT8, INT4) — уменьшает размер данных и ускоряет вычисления.
• Fused kernels — объединяет несколько операций в одно ядро (например, fused_attention).
• Dynamic batching — увеличивает occupancy за счёт объединения запросов.

Пример кода для включения FlashAttention в vLLM:

from vllm import LLM
llm = LLM(model="meta-llama/Llama-2-7b-hf", enable_flash_attention=True)

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

Задача Разверните небольшую LLM (например, LLaMA-2-7B) с vLLM на одной GPU, запустите нагрузочное тестирование с помощью wrk или locust, одновременно профилируя GPU с помощью nsys и ncu. Соберите метрики occupancy, memory bandwidth, warp stall reasons. Попробуйте изменить batch size и включить FlashAttention, сравните профили.

Инструменты vLLM, NVIDIA Nsight Systems, NVIDIA Nsight Compute, wrk (или locust), Python.

Шаги:

1. Установите vLLM и зависимости.
2. Запустите сервер vLLM с разными конфигурациями (batch size, FlashAttention on/off).
3. Запустите nsys profile во время нагрузки.
4. Извлеките горячие ядра и профилируйте их ncu.
5. Постройте графики occupancy, bandwidth, warp stall reasons.
6. Напишите отчёт с рекомендациями.

Ожидаемый результат Отчёт в виде Jupyter Notebook или PDF, содержащий:

• Скриншоты из Nsight Systems timeline.
• Таблицы метрик из ncu.
• Сравнение конфигураций (например, occupancy выросла с 40% до 70% при включении FlashAttention).
• Выводы: какая конфигурация оптимальна для данного сценария.

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

12. Навигация

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

Навигация

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