@tduncan / Nvidia-Optimization-Stack-VLLM
Code Commits Issues Pulls Cloud Agent CI/CD
main
main
md 346 lines 15.9 KB
Rendered Raw
1 # Profiling vLLM
2
3 !!! warning
4 Profiling is only intended for vLLM developers and maintainers to understand the proportion of time spent in different parts of the codebase. **vLLM end-users should never turn on profiling** as it will significantly slow down the inference.
5
6 !!! tip "Choosing a profiler"
7 - Use **Nsight Systems** for low-overhead, performance-critical profiling.
8 - Use **PyTorch Profiler** for medium-overhead profiling with richer debugging information (e.g., stack traces, memory, shapes). Note that enabling these features adds overhead and is not recommended for benchmarking.
9
10 ## Profile with PyTorch Profiler
11
12 We support tracing vLLM workers using different profilers. You can enable profiling by setting the `--profiler-config` flag when launching the server.
13
14 !!! note
15 The `--profiler-config` flag is available in vLLM v0.13.0 and later. If you are using an earlier version, please upgrade to use this feature.
16
17 To use the `torch.profiler` module, set the `profiler` entry to `'torch'` and `torch_profiler_dir` to the directory where you want to save the traces. Additionally, you can control the profiling content by specifying the following additional arguments in the config:
18
19 - `torch_profiler_record_shapes` to enable recording Tensor Shapes, off by default
20 - `torch_profiler_with_memory` to record memory, off by default
21 - `torch_profiler_with_stack` to enable recording stack information, on by default
22 - `torch_profiler_with_flops` to enable recording FLOPs, off by default
23 - `torch_profiler_use_gzip` to control gzip-compressing profiling files, on by default
24 - `torch_profiler_dump_cuda_time_total` to control dumping and printing the aggregated CUDA self time table, on by default
25 - `torch_profiler_activities` to select worker activities. Defaults are
26 platform-specific: `["CPU"]` on CPU, `["CPU", "CUDA"]` on NVIDIA GPUs, and
27 `["CPU", "XPU"]` on Intel GPUs. Selecting `["CUDA"]` omits CPU annotations
28 and the AsyncLLM CPU trace to reduce profiling overhead and trace size.
29
30 When using `vllm bench serve`, you can enable profiling by passing the `--profile` flag.
31
32 Traces can be visualized using <https://ui.perfetto.dev/>.
33
34 !!! tip
35 You can directly call bench module without installing vLLM using `python -m vllm.entrypoints.cli.main bench`.
36
37 !!! tip
38 Only send a few requests through vLLM when profiling, as the traces can get quite large. Also, no need to untar the traces, they can be viewed directly.
39
40 !!! tip
41 To stop the profiler - it flushes out all the profile trace files to the directory. This takes time, for example for about 100 requests worth of data for a llama 70b, it takes about 10 minutes to flush out on a H100.
42 The engine client waits for this flush to complete without timing out, so simply allow the stop call to run to completion.
43
44 ### Example commands and usage
45
46 #### Offline Inference
47
48 Refer to [examples/features/profiling/simple_profiling_offline.py](../../examples/features/profiling/simple_profiling_offline.py) for an example.
49
50 #### OpenAI Server
51
52 ```bash
53 vllm serve meta-llama/Llama-3.1-8B-Instruct --profiler-config '{"profiler": "torch", "torch_profiler_dir": "./vllm_profile"}'
54 ```
55
56 vllm bench command:
57
58 ```bash
59 vllm bench serve \
60 --backend vllm \
61 --model meta-llama/Llama-3.1-8B-Instruct \
62 --dataset-name sharegpt \
63 --dataset-path sharegpt.json \
64 --profile \
65 --num-prompts 2
66 ```
67
68 Or use http request:
69
70 ```shell
71 # We need first call /start_profile api to start profile.
72 $ curl -X POST http://localhost:8000/start_profile
73
74 # Call model generate.
75 curl -X POST http://localhost:8000/v1/chat/completions \
76 -H "Content-Type: application/json" \
77 -d '{
78 "model": "meta-llama/Llama-3.1-8B-Instruct",
79 "messages": [
80 {
81 "role": "user",
82 "content": "San Francisco is a"
83 }
84 ]
85 }'
86
87 # After need call /stop_profile api to stop profile.
88 $ curl -X POST http://localhost:8000/stop_profile
89 ```
90
91 ## Profile with Triton Proton
92
93 [Proton](https://github.com/triton-lang/triton/tree/main/third_party/proton)
94 is Triton's GPU profiler. It can collect a low-overhead aggregate tree or a
95 Chrome trace and works through the same vLLM profiling controls as the PyTorch
96 and CUDA profilers. Proton currently supports NVIDIA GPUs through CUPTI and
97 supports CUDA graph attribution.
98
99 Start a server with a local output directory and graph attribution:
100
101 ```bash
102 vllm serve meta-llama/Llama-3.1-8B-Instruct \
103 --profiler-config '{
104 "profiler": "proton",
105 "proton_profiler_dir": "./proton_profile",
106 "proton_output_format": "hatchet",
107 "proton_hook": "triton",
108 "proton_graph_attribution": true
109 }'
110 ```
111
112 Then use `/start_profile` and `/stop_profile` as shown above, or pass
113 `--profile` to a vLLM benchmark. Each worker uses a topology- and
114 rank-qualified output name, such as
115 `proton_dp0_pp0_tp0_dcp0_ep0_rank0_pid1234_0123456789abcdef0123456789abcdef_run0.hatchet`,
116 so distributed workers, restarted servers, and repeated profiling runs do not
117 overwrite one another. A `profile_prefix` is included when supplied. Each
118 profile is written by `/stop_profile` and is ready to inspect immediately.
119
120 The Proton-specific options are:
121
122 - `proton_context`: `shadow` (default) or `python`
123 - `proton_data`: `tree` (default) or `trace`
124 - `proton_backend`: `cupti` or automatic
125 - `proton_mode`: an optional backend mode string
126 - `proton_hook`: `triton` to record Triton launch metadata, or unset
127 - `proton_output_format`: `hatchet`, `hatchet_msgpack`, `chrome_trace`, or unset
128 - `proton_graph_attribution`: observe CUDA graph capture for replay attribution;
129 disabled by default and requires `proton_data: "tree"`
130
131 `hatchet` and `hatchet_msgpack` require `proton_data: "tree"`, while
132 `chrome_trace` requires `proton_data: "trace"`.
133
134 Automatic backend selection is recommended. vLLM currently supports Proton's
135 `cupti` backend on NVIDIA GPUs. ROCm support is not yet available. vLLM does not
136 expose Proton's experimental instrumentation backend because current upstream
137 can produce profiles without timing metrics. When `proton_graph_attribution` is
138 enabled, Proton observes vLLM's CUDA graph capture with the configured profiling
139 session active, then deactivates that same session until profiling starts. This
140 lets later profiles attribute replayed kernels without retaining model-startup
141 activity. Backend-specific modes can be selected with `proton_mode`;
142 `pcsampling` synchronizes the CUDA context and therefore requires
143 `--enforce-eager`. When CUDA graphs are enabled (including encoder graphs),
144 Proton requires `proton_graph_attribution: true` to collect replayed kernels.
145 For Chrome traces, disable CUDA graphs with `--enforce-eager`.
146
147 CUDA graph-attributed profiles support repeated `start_profile`/`stop_profile`
148 runs. Each stop flushes and writes one tree-data phase while preserving the
149 graph-aware session. Without `proton_graph_attribution`, each `stop_profile`
150 instead finalizes and writes an independent Proton session.
151
152 CUDA graph attribution requires Triton 3.7 or newer. It uses the phase data API
153 to discard graph-capture activity and separate profiling runs. The `hatchet_msgpack`
154 output format and `periodic_flushing` mode also require Triton 3.7 or newer.
155 `periodic_flushing` cannot be combined with graph attribution because both
156 manage the session's data phases. Ordinary Proton profiling remains available
157 with Triton 3.6.
158
159 Graph attribution retains capture metadata for the worker lifetime. With eager
160 execution or no graphs to capture, profiling uses ordinary independent sessions.
161
162 Inspect tree profiles with:
163
164 ```bash
165 proton-viewer -m time/ns \
166 proton_profile/proton_dp0_pp0_tp0_dcp0_ep0_rank0_pid1234_0123456789abcdef0123456789abcdef_run0.hatchet
167 ```
168
169 Chrome traces (`proton_data: "trace"`) can be opened in
170 <https://ui.perfetto.dev/>. Proton is imported lazily, so selecting another
171 profiler does not require a Proton-capable Triton installation.
172
173 ## Profile with NVIDIA Nsight Systems
174
175 Nsight systems is an advanced tool that exposes more profiling details, such as register and shared memory usage, annotated code regions and low-level CUDA APIs and events.
176
177 [Install nsight-systems](https://docs.nvidia.com/nsight-systems/InstallationGuide/index.html) using your package manager.
178 The following block is an example for Ubuntu.
179
180 ```bash
181 apt update
182 apt install -y --no-install-recommends gnupg
183 echo "deb http://developer.download.nvidia.com/devtools/repos/ubuntu$(source /etc/lsb-release; echo "$DISTRIB_RELEASE" | tr -d .)/$(dpkg --print-architecture) /" | tee /etc/apt/sources.list.d/nvidia-devtools.list
184 apt-key adv --fetch-keys http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/7fa2af80.pub
185 apt update
186 apt install nsight-systems-cli
187 ```
188
189 !!! tip
190 When profiling with `nsys`, it is advisable to set the environment variable `VLLM_WORKER_MULTIPROC_METHOD=spawn`. The default is to use the `fork` method instead of `spawn`. More information on the topic can be found in the [Nsight Systems release notes](https://docs.nvidia.com/nsight-systems/ReleaseNotes/index.html#general-issues).
191
192 The Nsight Systems profiler can be launched with `nsys profile ...`, with a few recommended flags for vLLM: `--trace-fork-before-exec=true --cuda-graph-trace=node`.
193
194 ### Example commands and usage
195
196 #### Offline Inference
197
198 For basic usage, you can just append the profiling command before any existing script you would run for offline inference.
199
200 The following is an example using the `vllm bench latency` script:
201
202 ```bash
203 nsys profile \
204 --trace-fork-before-exec=true \
205 --cuda-graph-trace=node \
206 vllm bench latency \
207 --model meta-llama/Llama-3.1-8B-Instruct \
208 --num-iters-warmup 5 \
209 --num-iters 1 \
210 --batch-size 16 \
211 --input-len 512 \
212 --output-len 8
213 ```
214
215 #### OpenAI Server
216
217 To profile the server, you will want to prepend your `vllm serve` command with `nsys profile` just like for offline inference, but you will need to specify a few other arguments to enable dynamic capture similarly to the Torch Profiler:
218
219 ```bash
220 # server
221 nsys profile \
222 --trace-fork-before-exec=true \
223 --cuda-graph-trace=node \
224 --capture-range=cudaProfilerApi \
225 --capture-range-end repeat \
226 vllm serve meta-llama/Llama-3.1-8B-Instruct --profiler-config.profiler cuda
227
228 # client
229 vllm bench serve \
230 --backend vllm \
231 --model meta-llama/Llama-3.1-8B-Instruct \
232 --dataset-name sharegpt \
233 --dataset-path sharegpt.json \
234 --profile \
235 --num-prompts 2
236 ```
237
238 With `--profile`, vLLM will capture a profile for each run of `vllm bench serve`. Once the server is killed, the profiles will all be saved.
239
240 #### Analysis
241
242 You can view these profiles either as summaries in the CLI, using `nsys stats [profile-file]`, or in the GUI by installing Nsight [locally following the directions here](https://developer.nvidia.com/nsight-systems/get-started).
243
244 ??? console "CLI example"
245
246 ```bash
247 nsys stats report1.nsys-rep
248 ...
249 ** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):
250
251 Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
252 -------- --------------- --------- ----------- ----------- -------- --------- ----------- ----------------------------------------------------------------------------------------------------
253 46.3 10,327,352,338 17,505 589,965.9 144,383.0 27,040 3,126,460 944,263.8 sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_of…
254 14.8 3,305,114,764 5,152 641,520.7 293,408.0 287,296 2,822,716 867,124.9 sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize256x128x64_warpgroupsize2x1x1_execute_segment_k_of…
255 12.1 2,692,284,876 14,280 188,535.4 83,904.0 19,328 2,862,237 497,999.9 sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize64x128x64_warpgroupsize1x1x1_execute_segment_k_off…
256 9.5 2,116,600,578 33,920 62,399.8 21,504.0 15,326 2,532,285 290,954.1 sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize64x64x64_warpgroupsize1x1x1_execute_segment_k_off_…
257 5.0 1,119,749,165 18,912 59,208.4 9,056.0 6,784 2,578,366 271,581.7 void vllm::act_and_mul_kernel<c10::BFloat16, &vllm::silu_kernel<c10::BFloat16>, (bool)1>(T1 *, cons…
258 4.1 916,662,515 21,312 43,011.6 19,776.0 8,928 2,586,205 199,790.1 void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMa…
259 2.6 587,283,113 37,824 15,526.7 3,008.0 2,719 2,517,756 139,091.1 std::enable_if<T2>(int)0&&vllm::_typeConvert<T1>::exists, void>::type vllm::fused_add_rms_norm_kern…
260 1.9 418,362,605 18,912 22,121.5 3,871.0 3,328 2,523,870 175,248.2 void vllm::rotary_embedding_kernel<c10::BFloat16, (bool)1>(const long *, T1 *, T1 *, const T1 *, in…
261 0.7 167,083,069 18,880 8,849.7 2,240.0 1,471 2,499,996 101,436.1 void vllm::reshape_and_cache_flash_kernel<__nv_bfloat16, __nv_bfloat16, (vllm::Fp8KVCacheDataType)0…
262 ...
263 ```
264
265 GUI example:
266
267 <img width="1799" alt="Screenshot 2025-03-05 at 11 48 42 AM" src="https://github.com/user-attachments/assets/c7cff1ae-6d6f-477d-a342-bd13c4fc424c" />
268
269 ## Continuous Profiling
270
271 There is a [GitHub CI workflow](https://github.com/pytorch/pytorch-integration-testing/actions/workflows/vllm-profiling.yml) in the PyTorch infrastructure repository that provides continuous profiling for different models on vLLM. This automated profiling helps track performance characteristics over time and across different model configurations.
272
273 ### How It Works
274
275 The workflow currently runs weekly profiling sessions for selected models, generating detailed performance traces that can be analyzed using different tools to identify performance regressions or optimization opportunities. But, it can be triggered manually as well, using the Github Action tool.
276
277 ### Adding New Models
278
279 To extend the continuous profiling to additional models, you can modify the [profiling-tests.json](https://github.com/pytorch/pytorch-integration-testing/blob/main/vllm-profiling/cuda/profiling-tests.json) configuration file in the PyTorch integration testing repository. Simply add your model specifications to this file to include them in the automated profiling runs.
280
281 ### Viewing Profiling Results
282
283 The profiling traces generated by the continuous profiling workflow are publicly available on the [vLLM Performance Dashboard](https://hud.pytorch.org/benchmark/llms?repoName=vllm-project%2Fvllm). Look for the **Profiling traces** table to access and download the traces for different models and runs.
284
285 ## Profiling vLLM Python Code
286
287 The Python standard library includes
288 [cProfile](https://docs.python.org/3/library/profile.html) for profiling Python
289 code.
290
291 ### Example usage - function call
292
293 If a filename is specified, the profile will be saved to that file. If no
294 filename is specified, profile data can be printed to stdout.
295
296 ```python
297 import cProfile
298
299
300 def expensive_function():
301 # some expensive code
302 pass
303
304
305 profiler = cProfile.Profile()
306 profiler.runcall(expensive_function)
307 profiler.dump_stats("expensive_function.prof")
308 ```
309
310 ### Example usage - context manager style
311
312 ```python
313 import cProfile
314
315
316 def another_function():
317 # more expensive code
318 pass
319
320
321 profiler = cProfile.Profile()
322 profiler.enable()
323 try:
324 another_function()
325 finally:
326 profiler.disable()
327 profiler.dump_stats("another_function.prof")
328 ```
329
330 ### Analyzing Profile Results
331
332 There are multiple tools available that can help analyze the profile results.
333 One example is [snakeviz](https://jiffyclub.github.io/snakeviz/).
334
335 ```bash
336 pip install snakeviz
337 snakeviz expensive_function.prof
338 ```
339
340 ### Analyzing Garbage Collection Costs
341
342 Leverage VLLM_GC_DEBUG environment variable to debug GC costs.
343
344 - VLLM_GC_DEBUG=1: enable GC debugger with gc.collect elapsed times
345 - VLLM_GC_DEBUG='{"top_objects":5}': enable GC debugger to log top 5
346 collected objects for each gc.collect