Benchmark Results

A Night in the GPU Mine: Dissecting a vLLM Nsight Trace

I spent a few hours staring at the entrails of a vLLM inference server through Nsight Systems, and it taught me more about how these machines actually operate under the hood.

The Setup, Or: Why I’m Profiling a Server That’s Mostly Asleep

I had a vLLM server running a unsloth-qwen3.6-35b-a3b-nvfp4-fast model (Qwen3 35B, active 3B, NVFP4 quantized, the details matter, we’ll get there), and I ran a multi-phase stress test against it while streaming Prometheus metrics every 2 seconds. The metrics told one that the server that was basically idle, KV cache usage peaking at a laughable 0.023%, zero preemptions, speculative decoding accepting 33.5% of its 45,520 drafted tokens. The CPU was at a pathetic ~1.7% utilization. All of this screams: this model is not being pushed hard enough.

I captured an Nsight Systems trace (nsys) of the same workload (50MB of native report, 149MB of SQLite) and generated the four canonical summary CSVs:

nsys stats --report cuda_gpu_kern_sum --report cuda_api_sum \
           --report osrt_sum --report cuda_gpu_mem_time_sum \
           --format csv --output nsys_report \
           qwen36_trace_%.sqlite

Let me take you through them one by one, because each one is a different window into the same machine, and together they tell a story that no single table can.

The CUDA API Table: Where the CPU’s Time Goes

The first table I always reach for is the API summary, because it tells you what the CPU is doing, and in inference servers, the CPU is usually the bottleneck.

Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
6,7222271872,941,7675102.9,908912.0,1408,40290944,14118642.4,cudaDeviceSynchronize
4,5990097264,411272,14564.8,2896.0,1360,17203136,380234.6,cudaMemcpyAsync
7,4168098480,478,8719871.3,896248.0,3200,216707264,31858475.5,cudaMalloc
4,1503715552,2,751857776.0,751857776.0,710150544,793565008,58982933.1,cudaHostRegister
5,817440160,12474,65531.5,61888.0,8272,702416,32521.1,cudaGraphLaunch_v10000
4,793711744,6727,117989.0,132800.0,832,6431360,139032.1,cudaEventSynchronize
1,734492560,250102,2936.8,736.0,544,44161056,150264.1,cudaStreamSynchronize

The first row is the diagnosis, and it is not good news: cudaDeviceSynchronize consumed 30.6% of all API time, 7.22 seconds total across just 941 calls.

Now, the average is 7.68ms per call. But look at the median: 908µs. And the max: 40ms. That gap between the average and the median is the tell. It means most of those 941 calls are fast (~0.9ms), but a handful are very slow, up to 40ms. The distribution is wildly skewed. Something is causing periodic stalls of 30-40 milliseconds, and the CPU is just sitting there, spinning, waiting for the GPU to finish a batch.

Let me put a number on how bad this is. The GPU kernels themselves only account for ~5-8ms of compute per batch. So when cudaDeviceSynchronize blocks for 38ms, roughly 30ms of that is pure idle: the GPU already finished, but the CPU is waiting anyway because someone called a blocking sync instead of using events. Multiply 30ms of wasted time by 941 calls and you get ~28 seconds of pure wasted wall-clock. Even in this 83-second trace, that’s enormous. This is the single biggest lever in the whole system.

Then there’s cudaMemcpyAsync: 411,272 calls. That’s not a number, that’s a firehose. Four hundred thousand async copy submissions, each taking ~14µs of CPU overhead. This is the second biggest time sink at 6 seconds. And cudaMalloc: only 478 calls, but averaging 8.7ms each, 4.17 seconds total. The server is creating and destroying GPU memory tensors far more often than it should, and each one drags in the driver’s memory allocator.

I want to pause on cudaHostRegister: 2 calls, 1.5 seconds total. That’s 752ms per call. That’s the server pinning host memory for fast DMA. Two calls, each taking three-quarters of a second, you can actually see the server grinding through a huge chunk of host memory when you look at the timeline. That’s the kind of thing that shows up as a mysterious stall if you’re not looking at the right table.

The API table is basically a list of the CPU’s sins. And the biggest sin, by far, is the blocking synchronize.

The GPU Kernel Table: What the Silicon Actually Does

Now flip to the other side: what the GPU actually computes. This is the table that makes the whole thing make sense:

Time (%),Total Time (ns),Instances,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
5,3819843774,22850,167170.4,80288.0,38720,2506720,424616.9,"void cutlass::device_kernel<...GemmUniversal...MainloopSm120TmaWarpSpecialized<(int)9, (int)2,...>"
9,2591928541,54136,47878.1,3776.0,576,335808,60000.8,"void at::native::vectorized_elementwise_kernel<(int)4, at::native::FillFunctor<int>...>"
4,1946677439,7466,260739.0,5792.0,5472,1472704,432073.0,triton_red_fused__to_copy__unsafe_view_abs_add_clamp_clone_cutlass_scaled_mm_div_max_mean_mul_pow_reciprocal_rsqrt_silu_unsqueeze_view_1
0,1680607021,10891,154311.5,150432.0,48064,1331136,69958.3,"void cutlass::device_kernel<...GemmUniversal...GroupProblemShape...MainloopSm120ArrayTmaWarpSpecializedBlockScaled...>"
9,1290966463,6786,190239.7,3920.0,3136,1456352,385863.6,triton_red_fused__to_copy_abs_add_clamp_cutlass_scaled_mm_div_fused_add_rms_norm_max_mul_reciprocal_unsqueeze_3

There’s a beautiful amount of information encoded in that mangled first kernel name, and it’s worth decoding because it tells you exactly what hardware this is running on. The name contains SM120: that’s Blackwell. TmaWarpSpecialized: Tensor Memory Accelerator, warp-specialized. float_e4m3_t: that’s FP8 (E4M3 format). cutlass::gemm::kernel::GemmUniversal: this is CUTLASS, NVIDIA’s template GEMM library, not cuBLAS. So this is a hand-tuned Blackwell GEMM with FP8 inputs, TMA-based data movement, and warp specialization, the state of the art.

But here’s the kicker: 22,850 instances of the big GEMM, averaging 167µs each. Twelve kernels per millisecond. The GPU is being asked to do thousands of tiny matrix-multiply operations in a tight loop. This is the classic signature of attention with per-token GEMMs: the Qwen model has 3B active parameters, so each token generation step fires off a small GEMM per head. And the median is 80µs while the max is 2.5ms, again the skew, again the outliers. These are not uniform operations; they’re a mix of tiny per-token kernels and occasional big-batch kernels.

Now look at the second row: 54,136 instances of vectorized_elementwise_kernel with FillFunctor<int>. Fifty-four thousand fill operations. The server is writing zeros (or some constant) to GPU memory fifty-four thousand times, and it’s eating 13.9% of total GPU time. This is the KV cache. Every new token position needs a freshly-zeroed key/value buffer, and someone is zeroing it one tiny kernel at a time instead of reusing a pre-zeroed pool. This is the kind of inefficiency that’s invisible in Prometheus metrics but screams from a kernel histogram.

Then the Triton fused kernels, triton_red_fused__to_copy__unsafe_view_abs_add_clamp_clone_cutlass_scaled_mm_div_max_mean_mul_pow_reciprocal_rsqrt_silu_unsqueeze_view_1. That name is a whole sentence: it’s the fused RMSNorm + SiLU + scaled-matmul + clamp + absolute-value + add + copy + view + unsqueeze kernel. This is vLLM’s torch.compile-style fusion collapsing an entire MLP sublayer’s worth of operations into a single kernel. The cutlass_scaled_mm in the name tells you it’s the NVFP4 block-scaled path. And the _red_ prefix means it’s the “red” (reduction) variant: this is the one that does the per-token scaling reductions.

So the kernel table’s story is: the model is doing real, sophisticated, fused compute, but it’s buried under a pile of 54,000 fill kernels and 22,850 tiny GEMMs. The silicon is capable of so much more than it’s being asked to do.

The Memory Table: The 68 Gigabytes That Never Left the GPU

Now the most interesting table of the four. The memory transfer summary:

Time (%),Total Time (ns),Count,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Operation
7,948153344,396828,2389.3,960.0,768,9904128,41710.7,[CUDA memcpy Device-to-Device]
6,839393760,23912,35103.5,768.0,160,753024,64323.8,[CUDA memcpy Host-to-Device]
5,8405470,3628,2316.8,832.0,160,237216,13177.6,[CUDA memset]
2,3517152,1596,2203.7,1536.0,544,24960,1971.2,[CUDA memcpy Device-to-Host]

Read that first line again. 396,828 Device-to-Device copies, 948ms total. That’s 397 thousand copies where the data never left the GPU: it went from one GPU memory location to another GPU memory location. And here’s the punchline: when I drilled into the SQLite database by copy size, the dominant chunk size was exactly 524,288 bytes (512KB), repeated 92,520 times, totaling 46 gigabytes of pure intra-GPU shuffling. Plus ~304,000 tiny 4-64 byte copies, and a few ~400 big 16MB-1GB bursts.

What is being copied 512KB at a time, four hundred thousand times, all within GPU memory?

This is the CUDA graph. When vLLM captures an operation as a CUDA graph (which it does to amortize launch overhead), every cudaMemcpyAsync inside the captured region gets recorded into the graph. And when the graph replays, those copies replay too. The 512KB uniform chunk is the tell: it’s a tiled layout. The framework is staging data through a scratch buffer in 512KB tiles, copying tile-by-tile through shared memory or a staging area, all so that the async copy engine (the copy DMA) can overlap with the compute (the SM120 MMA units). This is the classic Blackwell-style “pingpong” TMA pattern (KernelTmaWarpSpecializedPingpong), where you double-buffer tiles and use TMA loads to prefetch the next tile while the current one computes.

So the 68 gigabytes of D2D copies aren’t waste: they’re the pipeline. They’re the plumbing that lets compute and memory move in parallel. The tragedy is that they’re being replayed 397,000 times because every graph launch replays the whole recorded sequence, even when only a tiny slice of it is needed for the current token.

The OS Runtime Table: The Wait That Hides in the Kernel

The fourth table is the one people skip, and it’s the one that explains the 40ms stalls:

Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
1,1121839560352,58476,19184615.2,10414256.0,1008,92819426416,663048606.1,epoll_wait
9,971480778944,11599,83755563.3,100155696.0,1024,4000870368,264600554.5,pthread_cond_timedwait
4,665111279632,407,1634180048.2,499920.0,1040,33510823888,7069034354.9,futex
1,220011952976,34,6470939793.4,10000302136.0,21136,10000985024,4850792304.7,sem_timedwait
1,219742545312,2858,76886824.8,9600.0,1008,14124115968,525102047.1,poll
9,141953547152,457,310620453.3,31582240.0,8896,42509471648,2603127118.6,sem_wait
4,124212057008,1492,83252048.9,3244624.0,1008,1002528768,231370028.1,epoll_pwait
7,96916214960,23844,4064595.5,2096.0,1008,85922753568,557310243.8,read

There it is. 58,476 calls to epoll_wait, 1,122 seconds of wall-clock. The server is spending 31% of its time in the event loop, waiting for something to arrive on a socket. This is the HTTP serving layer: vLLM’s async engine sitting in epoll_wait, idle, waiting for the next request. And futex: 407 calls but one of them blocked for 33.5 seconds. A futex is the kernel primitive behind mutex/condvar. A 33-second futex block means a thread was parked on a lock for half a minute. That’s the scheduler thread contending with the model executor: the classic vLLM pattern where the Python scheduler and the CUDA executor fight over a shared lock.

And look at the max on pthread_cond_timedwait: 4 seconds. That’s a thread sleeping on a condition variable for four full seconds. In an inference server, that’s the executor waiting for work to be scheduled.

Here’s the beautiful irony that ties the whole trace together. The OS table shows the server is asleep for 31% of its life, waiting on sockets and locks. The API table shows the CPU blocked on cudaDeviceSynchronize for another huge chunk. And the kernel table shows that when it finally does compute, it wastes half its silicon on fill kernels and tiny GEMMs. The machine is simultaneously under-utilized (sleeping on epoll, waiting for requests) and poorly utilized (blocking on syncs, fragmenting compute into 54k fills and 22k tiny GEMMs). Both problems are real, and they compound. You can’t fix one without the other.

Putting the Four Tables Together: The Anatomy of a Stall

Let me walk you through what actually happens, microsecond by microsecond, when this server generates a token. Because the four tables, taken together, are a biography of one token:

The scheduler wakes. epoll_wait returns. The Python scheduler sees the request, walks the attention tree, decides what to do. It grabs a lock, and sometimes waits 33 seconds on a futex because the executor thread holds it. (OS table: futex.)
The executor blocks on the GPU. The CUDA executor launches the batch. Then, instead of using an async event and going back to the scheduler, it calls a blocking cudaDeviceSynchronize and parks. The GPU runs the batch in ~5-8ms of real compute, but the CPU sits there for up to 40ms because the sync over-waits. (API table: cudaDeviceSynchronize.)
The GPU fires the fused kernels. The RMSNorm+SiLU+scaled-matmul Triton kernels run, beautiful, dense, Blackwell-native compute. But interleaved with them are 54,000 FillFunctor zeroing kernels and 22,850 tiny per-token GEMMs, plus 397,000 graph-replayed D2D copies. (Kernel table, memory table.)
The GPU finishes. The CPU is still waiting. The sync returns, the executor hands results back, the scheduler formats the response, and the whole thing goes back to sleep in epoll_wait until the next request. (OS table: epoll_wait.)

Every one of those four steps is a place where time leaks. And the trace shows all four leaks simultaneously, because they’re the same leak viewed from different floors of the building.

The Technical Concepts, Made Concrete

Since this whole post is one long taxonomy of what those CSV columns actually mean, let me make sure the vocabulary is nailed down:

cudaDeviceSynchronize vs events. The cardinal sin. A blocking sync makes the CPU wait for all prior work on the device, including the hundreds of thousands of replay copies and fill kernels queued behind it. An event (cudaEventRecord + cudaEventSynchronize) lets you wait for one specific operation. The fix: record an event after the last kernel and sync on that event, so the CPU only waits for the actual work.

CUDA graphs. The cudaGraphLaunch calls (12,474 of them) are the framework capturing a sequence of kernels into a DAG and replaying it. The benefit is amortized launch overhead. The cost, visible right here: every replay replays the whole recorded choreography, including 397,000 D2D copies and 54,000 fills, even when the current token only needs a slice. The D2D copy count is 397,000: that’s the graph replaying its pipeline over and over.

TMA and warp specialization. The MainloopSm120TmaWarpSpecialized kernel name is Blackwell’s async-copy + double-buffering pattern. One warp issues TMA loads to prefetch the next tile while the MMA warps compute the current tile. This is good engineering: it’s the only reason the GEMMs aren’t even slower. But it only pays off if the graph isn’t replaying 512KB tiles four hundred thousand times.

Block-scaled quantization (NVFP4). The cutlass_scaled_mm path means weights are stored in 4-bit blocks with per-block scales, and the _red_ (reduction) Triton variant computes the per-token scaling reductions on the fly. So the model can be “35B with 3B active”: the MoE experts are quantized to NVFP4, which is why the memory table shows so much staging traffic, the scaling factors get shuffled around as part of the D2D copies.

The epoll/futex anatomy. epoll_wait is the async engine idle between requests: that’s the throughput problem, the server is under-loaded. futex is the scheduler/executor lock contention: that’s the latency problem, a single 33-second park. They’re opposite diseases: too idle, and too contended.

The Verdict

The trace is one of the clearest I’ve ever read. It’s not a hardware problem. It’s an orchestration problem.

It’s under-loaded. epoll_wait is 31% of wall-clock: the server is asleep, waiting for traffic. The Prometheus metrics confirmed it: 0.023% KV cache usage, the model is being asked to do almost nothing.
It wastes its compute when it does work. 54,136 fill kernels and 22,850 tiny GEMMs and 397,000 D2D copies: that’s the CUDA graph replaying a fat captured region for every single token. The fix is to capture a slice per token, not the whole pipeline, and to pre-zero the KV cache pool once instead of re-filling it 54,000 times.
It blocks instead of streaming. cudaDeviceSynchronize over-waits up to 40ms because the executor parks on a full-device sync instead of an event. Replace blocking syncs with event syncs and you reclaim most of that 40ms.
It contends instead of coordinating. The 33-second futex park is the Python scheduler and the CUDA executor fighting over one lock. Give the scheduler a real event loop and a lock-free handoff, and that contention disappears.

So the machine is simultaneously too idle and too contended: a server asleep on epoll, blocking on syncs, and replaying a fat graph. Blackwell’s TMA pingpong and the fused Triton kernels are doing real work, but they’re buried. Give the scheduler a real event loop, replace blocking syncs with event syncs, pre-zero the KV cache once, and stop replaying the whole graph per token: those four things turn this trace from a portrait of wasted cycles into a portrait of a server actually using its silicon. None of it shows up in Prometheus. All of it shows up in a 50MB nsys trace, if you’re willing to read four CSVs like they’re a blood panel and follow the numbers where they lead.

The Raw Numbers, For The Skeptics

Table	Line	Number	What it means
API	`cudaDeviceSynchronize`	941 calls, 7.22s, 30.6%	CPU blocked waiting on the GPU
API	`cudaMemcpyAsync`	411,272 calls	Firehose of async copies
Kernel	Big GEMM	22,850 instances, 167µs avg	Per-token attention matmuls
Kernel	`FillFunctor`	54,136 instances, 13.9%	KV cache zeroing, replayed per token
Kernel	Triton fused	~6,786 instances	The actual MLP compute
Mem	D2D copies	396,828, 948ms	Graph-replayed staging, 46GB of 512KB tiles
Mem	512KB tile	92,520 copies	The TMA pingpong staging pattern
OS	`epoll_wait`	58,476 calls, 1,122s, 31%	Server idle, waiting for sockets
OS	`futex`	407 calls, one 33.5s block	Scheduler/executor lock contention
OS	`pthread_cond_timedwait`	max 4s	Executor waiting for work

Thanks for reading!

Colin Zhou