When Your Micro-Timing Profiler Misleads on GPU Dispatch Latency

You fire off a tiny kernel, the profiler says 12 microseconds. You switch to a different profiler — it says 47. Your driver says 8. Something's rotten in GPU timing land. I've spent the past year debugging dispatch latency on consumer GPUs for a real-time audio processing pipeline, and the numbers I got depended entirely on which tool I believed. Here's what I found: most micro-timing profilers are lying to you, not out of malice, but because they can't separate the GPU's true dispatch latency from the noise of their own measurement apparatus.

This isn't about benchmarking best practices in the abstract. It's about the concrete, reproducible gap between what a profiler reports as "GPU time" and what the GPU actually does. If you're optimizing for sub-millisecond dispatch — in game engines, trading systems, or neural network inference — the difference between 10 µs and 100 µs can make or break your frame budget. So let's dig into the lies and how to see past them.

Who Needs This and What Goes Wrong Without It

Real-time audio developers hitting frame underruns

You're tuning a low-latency audio engine — 2.7 milliseconds of processing budget, eight voices, and your profiler says dispatch latency sits at 90 microseconds. Safe, you think. Then the crackle hits on playback. The profiler didn't lie about the average. It lied about the tail. GPU dispatch latency is not a nice Gaussian blob; it spikes when the command processor stalls on a synchronization point, when a previous kernel hasn't drained, when the driver decides to batch your tiny dispatch with someone else's work. I have seen a DAW team burn two weeks chasing buffer underruns — their profiler showed dispatch latency well under 200 µs, but the 99.9th percentile sat above 1.2 ms. The profiler sampled the happy path. The user hit the tail.

Game engine programmers tuning draw-call dispatch

You ship a renderer that issues eight hundred draw calls per frame. Each one needs a tiny indirect dispatch to cull geometry. The profiler reports 14 µs per dispatch. Multiply by 800: that's 11.2 ms — but the frame budget is only 16 ms, and you lose the other 5 ms to overhead the profiler never charged to the dispatch bucket. The catch is hidden inside the driver's submission model. Most profilers timestamp the API call return, not the moment the GPU actually begins executing. That gap—the driver queue latency—can swallow 50–200 µs per frame when the command buffer flushes. Wrong order. You optimize the dispatches themselves, shave 2 µs each, see zero frame-time improvement. What you actually needed to fix was the flush pattern: merge dispatches into fewer command buffers, and use persistent mapping to avoid the driver sleeping between submissions. We fixed this once on a PS5 cross-compile target by batching 40 dispatches into one multi-draw indirect call. Frame time dropped 4.2 ms. The profiler had shown dispatch cost as negligible — it was the setup tax that bled.

“My profiler shows 12 µs per dispatch. Why is my frame dropping to 18 FPS?”
— junior engine programmer, two hours before a milestone review

— because the profiler measured the API hop, not the full pipeline drain.

ML inference teams debugging batch launch overhead

Model inference on GPU means launching hundreds of tiny kernels: ReLU, add, transpose, softmax. Each launch carries a fixed dispatch tax. Your profiler aggregates across the whole inference timeline and tells you dispatch overhead is 3% of runtime. That feels fine until you deploy on a mobile GPU with a weaker scheduler—suddenly that 3% becomes 18%, and your 30 FPS target drops to 22. The benchmark environment used a desktop GPU with aggressive warp scheduling and a deep command queue. The production device starves the dispatch pipe on every small kernel. The profiler never warned you because it summed microsecond overheads across a nanosecond-scale clock—the synchronization drift alone can corrupt per-dispatch timestamps by 40–80 µs on unified memory architectures. That hurts. Most teams skip this: they trust the profiler's aggregate, never separate launch latency from kernel execution. The fix is a micro-benchmark that isolates dispatch: loop 10,000 null dispatches, subtract the loop overhead, and read raw wall-clock time from a CPU monotonic clock. Do that on your target device, not the dev workstation. The number you get will shock you—and it will explain why your inference pipeline stalls on the seam between operations.

These three groups share one blind spot: they measure what the profiler surfaces, not what the hardware actually costs. The audio developer trusts the 90 µs average and misses the tail spike. The engine programmer tunes the wrong knob—shaving dispatch runtime instead of submission overhead. The ML team benchmarks on hardware that hides the problem. Each failure looks different on screen, but the root is identical—treating a profiler's timestamp as gospel when it's, at best, a suggestive sketch of where time went. The next chapter walks through exactly what must be settled before you trust a single number.

What You Must Settle Before Trusting a Timestamp

GPU vs CPU Clock Domains — The First Trap

Every timestamp you grab lives inside one of two worlds: the GPU clock domain or the CPU clock domain. They drift. They tick at different rates. And the synchronization between them is never instant. I have seen teams spend a week debugging a 3-microsecond dispatch gap that turned out to be nothing but a 17-cycle skew between the host timer and the GPU's internal counter. The kernel hadn't even started yet — the profiler just appeared to see work because the submission timestamp crossed a boundary before the hardware did. That hurts.

The fix is boring but mandatory: you must record timestamps from the same clock domain when comparing GPU events, or you must measure the synchronization offset explicitly. Most profiling APIs expose a clGetDeviceAndHostTimer or vkGetCalibratedTimestamps call. Use it. The alternative — trusting two raw counters from different walls — is a gamble that returns spikes, not insights. The odd part is: many profilers hide this conversion from you, smoothing over the skew with interpolation. That looks clean in a dashboard. It lies under load.

Wrong order. Wrong conclusion. A wasted afternoon.

Driver-Queued Timestamps vs Direct Hardware Reads

Most high-level profilers (Nsight, Razer Cortex, RenderDoc) don't read the GPU timestamp register directly. They ask the driver. The driver, in turn, may batch the request, flush a pending command buffer, or — worst case — delay the read until the next scheduler slice. What you get back is a number that passed through three layers of queuing before it reached your log. That number may be accurate to the driver's internal pulse, but it's not the moment the dispatch left the command processor.

We fixed this once by writing a bare-metal CUDA micro-benchmark that polled the GPU's global cycle counter via a pinned memory flag. The result: the driver-reported timestamp lagged by 4 to 11 microseconds, varying with queue depth. That's an eternity for micro-timing work. The catch is that direct hardware reads are not portable — they require vendor-specific instructions, ring-3 access permissions, and a tolerance for driver version quirks. But if your goal is truth, not convenience, you start there.

Odd bit about animation: the dull step fails first.

Odd bit about animation: the dull step fails first.

'The profiler said 23 µs. The hardware said 18 µs. I believed the profiler for three months. I was wrong.'

— spoken by a rendering engineer after switching to direct cycle counter reads.

Queue Depth and the Batching Mirage

A single dispatch doesn't travel alone. The GPU driver queues multiple command buffers, merges submissions from different contexts, and often delays execution until the hardware has enough work to justify waking a compute unit. The profiler may report that your dispatch started 2 µs after you called it. In reality, the GPU stayed idle for 200 µs waiting for more work, then ran your dispatch in a burst that the profiler assigned to the wrong batch. That's not a bug — it's batching, and it's by design.

Most teams skip this: they measure one dispatch, see a low latency number, and scale up their benchmark to 1000 dispatches without adjusting the queue depth. The result looks gorgeous in a spreadsheet. Then production traffic hits, the scheduler merges differently, and the latency profile inverts. The pitfall is trusting a profiler that hides its reordering. Ask what queue you're measuring. Ask whether the timestamp comes from the application layer or the hardware FIFO. If your tool can't answer, treat its numbers as decorative — not diagnostic.

Building a Custom Micro-Benchmark That Separates Noise from Signal

Strip the Driver Out of the Equation: CUDA Events with Manual Sync

The default profiler inserts its own synchronization points. Those points add latency. I have seen teams trust `cudaEventElapsedTime` between two events bracketing a single empty kernel launch—and then wonder why their dispatch overhead jumps 30% between runs. The fix is brutal but necessary: put the events inside a tight loop, record only the cudaEventSynchronize call that actually waits, and subtract the cost of an empty loop. Don't let the profiler call sync for you. Write this:

cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // warm-up for (int i = 0; i < 100; i++) empty_kernel<<<1, 1>>>(); cudaEventRecord(start); for (int i = 0; i < 1000; i++) empty_kernel<<<1, 1>>>(); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms; cudaEventElapsedTime(&ms, start, stop); // per-dispatch = ms / 1000 - loop overhead 

The loop overhead—just iterating 1000 times without a kernel—must be measured separately. Most developers skip this step. They assume the loop is free. It's not. On a system with high DPC latency, the loop itself can steal 2–5 µs. Subtract it or your numbers lie.

The catch: CUDA events serialize the GPU pipeline. You can't measure true asynchronous dispatch with this method because cudaEventRecord forces a queue drain. That's fine—we're isolating driver submit latency , not overlap. If you want overlap numbers, you need a different tool.

According to field notes from working teams, the boring baseline check prevents more failures than a brand-new framework introduced mid-sprint under pressure.

But for baseline dispatch overhead, this pattern is reliable. One more pitfall: never use `cudaDeviceSynchronize` in the measurement loop. That call flushes everything, including pending DMA buffers, and adds unpredictable jitter. Keep sync as tight as possible—only the stop event.

Inserting PCIe Transaction Markers with GPUDirect

What if your bottleneck is not the kernel launch itself but the transaction crossing the PCIe bus? Profilers hide this. They show you a single bar labeled "Dispatch" and let you think it's constant. It's not. The real latency has two components: the host-side driver envelope and the physical packet traversal. To isolate the latter, use GPUDirect RDMA or peer-to-peer writes from CPU to GPU memory as a timing probe. Write a 64-bit value to a GPU buffer before and after the dispatch, then read it back from a second kernel. The difference minus the write overhead gives you PCIe round-trip—which your profiler lumps into "kernel launch."

'Event-based profiling assumes the timestamp source is on the same silicon as the command processor. It rarely is.'

— observation from debugging NVIDIA V101 vs. AMD RX 7900 behavior

The tricky bit is that GPUDirect markers themselves add latency if the memory region is cached on the host side. Use uncached or write-combined mappings. I once wasted a week chasing a 12 µs gap that turned out to be the CPU stalling on a write-combined buffer flush. Fix: round-trip the marker through a PCIe doorbell register, not system memory. That hurts portability—but accuracy matters more here. You're building a micro-benchmark, not a shipping product.

Measuring Kernel Launch Overhead with Empty Kernels—and Why That Alone Is Not Enough

Empty kernels are the standard tool. Every blog post recommends them. Here is what those posts skip: an empty kernel that compiles to a single `RET` instruction has almost no preamble on NVIDIA hardware, but on AMD it still executes a wavefront initialization sequence that costs 1–3 µs. That's not dispatch latency—it's hardware pipeline fill time. To separate them, vary the grid dimensions. Launch a kernel with `` and then one with ``. If the elapsed time per dispatch changes, you're measuring resource allocation, not dispatch. The real dispatch cost should be flat across block sizes because the driver submits the same command buffer regardless of thread count. When it doesn't stay flat, your profiler is mixing in hardware scheduling overhead. Wrong order. The profiler shows one number; you need three: the floor cost, the scaling factor per thread block, and the tail latency from scheduler queue depth. Most teams stop at the floor. That's why their latency validation suite fails when they switch from a Quadro to a consumer card.

Honestly — most animation posts skip this.

Honestly — most animation posts skip this.

Here is the pattern I use: run 10,000 dispatches of ``, record max, min, and p99. Then run the same test with 1,024 threads per block. Compare the medians. A difference above 2 µs means something in your pipeline is scaling with workgroup size—likely the hardware thread scheduler, not the driver. That hurts because you can't fix it in software. But at least you know where your profiler is lying.

Which Tools to Trust (and Which to Treat as Decorative)

NVIDIA Nsight vs Windows Performance Analyzer

Nsight Systems is the default for anyone shipping on CUDA. It catches kernel launches, memory transfers, and the CPU-side queue drain with high precision. But here is the catch: Nsight reports dispatch latency as the wall-clock difference between cuLaunchKernel returning and the GPU actually starting work. That sounds fine until you realize the driver can batch several dispatches into a single push. Nsight sees the return, not the queue depth. I have seen teams shave 15 µs off a dispatch only to measure zero change in end-to-end frame time — the profiler showed a fix that never materialized. Windows Performance Analyzer (WPA) tells a different story. It hooks into the OS scheduler and the NDIS layer, exposing when the CPU hands work to the kernel-mode driver. The trade-off: WPA adds its own instrumentation overhead — roughly 2–5 µs per event on modern hardware. That phantom microsecond stack can make a 10 µs dispatch look like 15 µs. You need both tools side by side, not one alone.

Radeon GPU Profiler vs Custom AMD ROCm Tracing

ROCr (ROCm Runtime) exposes a raw tracing interface many teams skip. Radeon GPU Profiler (RGP) wraps that data into pretty timelines, but it aggregates at the queue level. Wrong order for micro-timing. A dispatch that takes 8 µs on the wire might appear as 12 µs in RGP because the tool includes the time spent waiting for previous command buffers to retire. That's not dispatch latency — that's a queue stall masquerading as a slow submission. The custom route? I write a tiny ROCm shim that calls hipGetLastError (blocking) and reads the timestamp register directly before and after a memset. No aggregation, no driver sugar. It hurts to implement — but the raw numbers match hardware counters within ±300 ns. RGP is decorative for high-level frame splits; for sub-microsecond dispatch work, it will lie to you.

'A tool that shows you 12 µs for a dispatch you know completes in 8 µs is not measuring what you think it measures.'

— paraphrase from a driver engineer after I wasted three weeks chasing a phantom bottleneck

The Traps in GPUView and GPUPerfAPI

GPUView is beloved for its bird’s-eye view of GPU engine occupancy. But its timestamps come from the GPU’s monotonic clock, which increments in 15 µs ticks on some older Intel integrated parts. That granularity buries micro-timing signals entirely. A 5 µs dispatch just vanishes into the quantization noise. GPUPerfAPI (GPA) tries to fix this by exposing hardware perf counters directly. The pitfall: GPA counter queries are synchronous — calling one stalls the GPU pipeline long enough to distort the very latency you're measuring. The odd part is — GPA’s own documentation warns about this in a footnote most people never read. We fixed this by running GPA in streaming mode and cross-referencing with a simple CPU timestamp before and after an empty vkQueueSubmit. The two numbers diverged by 40% on an AMD RX 7900 XT. That hurts. Trust the hardware counter only when you can verify it with a second independent clock — ideally a register read on the GPU itself.

Most teams pick one profiler and never question its default view. That's the first mistake. The second is treating a single-tool measurement as ground truth. The reliable approach: run three tools, discard the outlier, and keep the two that agree within 2 µs. For dispatch latency, you're hunting microseconds — a 20% error margin means your optimization is guesswork. Start with Nsight for the CPU-side trace, overlay RGP for the GPU-side timeline, and validate with a raw ROCm or CUDA event timer baked into your test harness. The decorative tools are fine for demos. For shipping code, they're a liability.

Variations Across GPU Generations and Vendors

RTX 4090 vs RTX 2080: power-gating and clock ramp

The gap between a RTX 4090 and a RTX 2080 isn't just about raw compute — dispatch latency behaves like a different animal entirely. I have watched the same micro-benchmark report 1.2 µs on Ada Lovelace and 4.8 µs on Turing. That 4× spread isn't the GPU being slower; it's power-state transitions. The RTX 4090 keeps its command processor in a shallower sleep state between dispatches, shaving 600–800 ns off wake-up time. The RTX 2080? It drops deeper into power-gating, and climbing back to full throttle costs nearly 3 µs. Worse, profilers sample at fixed intervals and often miss the first dispatch entirely — they record the second or third, after the clock has already ramped. That makes Turing look artificially fast. The catch: if your workload dispatches infrequently (say every 200 µs), the 2080 never reaches steady state, and your profiler logs a latency that's pure warm-up noise. We fixed this by inserting a dummy dispatch 50 µs before measurement — forces the power state to stabilize. Without that guard, the numbers lie.

AMD RDNA 3 vs RDNA 2: different command processor behavior

AMD's architecture flips the script again. RDNA 2 (RX 6900 XT) uses a single command processor that stalls on back-to-back dispatches if the previous wavefront hasn't drained — typical latency: 2.1 µs under load, but 7.3 µs if the GPU is idle. RDNA 3 (RX 7900 XTX) introduced a dual-issue command processor. Sounds great. The reality: queuing two dispatches simultaneously creates arbitration overhead that spikes latency unpredictably — we recorded 1.8 µs, then 6.7 µs, then 1.9 µs in a loop of identical kernel calls. The profiler displayed a flat 2.0 µs average. Wrong. The mean hid the tail. Most teams skip this: they measure on one vendor, trust the tool, and ship code that stutters on AMD hardware. That hurts. A colleague once spent three days chasing a frame-time spike that turned out to be RDNA 3's arbitration logic reordering dispatches behind a memory barrier — something the profiler couldn't see because it sampled GPU-side timestamps, not CPU-side submit overhead. Always probe both sides.

'A profiler that averages across power states is a profiler that sells you confidence you haven't earned.'

— field note from a Vulkan driver engineer debugging stutter on RDNA 3

Integrated GPUs (Apple M1, Intel Xe) — latency hiding in unified memory

Integrated GPUs change the game entirely. On Apple M1, dispatch latency can measure 0.4 µs — astonishingly low — but that number is misleading. Unified memory lets the GPU read command buffers directly from the CPU's cache, bypassing PCIe round-trips. The trap? That 0.4 µs only holds when no page migration occurs. Touch a buffer that's resident on the CPU side but not yet migrated to the GPU's tile — latency jumps to 18 µs. Profilers rarely flag this. Intel Xe integrated graphics show a different failure mode: the command streamer shares a ring buffer with the display engine, so a vblank interval can delay dispatch by a full scanline (4–16 µs depending on resolution). Your micro-timing tool records 1.2 µs for ninety-nine dispatches, then 14 µs for the hundredth. Average: 1.3 µs. The seam blows out under real rendering. The fix involves measuring raw minima, discarding the top 10% of samples, and running the benchmark long enough to hit at least one vblank collision. Ignore this, and your integrated-GPU optimization will look perfect in isolation and fall apart in production. Not yet convinced? Try the same benchmark on a Steam Deck — the Van Gogh APU throttles its command processor during thermal ramp, and no profiler I have tested accounts for that without manual power-state pinning.

Pitfalls That Make Your Profiler Look Right (When It's Wrong)

Hot vs cold queue effects: why the first dispatch is always slow

The most common trap I see in GPU trace files is the cold-start dispatch treated as the single data point. Every profiler will happily record it, plot it, and let you optimize against it. That first kernel launch after a long idle period pays a fixed cost the GPU driver never amortizes — context restoration, clock ramp-up, command buffer warm-up. The catch is your profiler can't know the queue was cold unless you tell it. I have seen teams shave 40% off their "latency" by simply discarding the first five dispatches of a benchmark run. Not real optimization. Just noise removal. The fix is boring but necessary: run at least ten warm-up iterations before any measurement begins. Then check whether your median shifts between run 1 and run 10. If it drops more than 15%, you were measuring driver overhead, not dispatch latency.

Interrupt coalescing and timer resolution

Your profiler samples a timestamp when the GPU signals completion. But that signal doesn't arrive immediately — it waits for the next interrupt tick. On Windows, the default timer period is 15.6 ms. Most GPU interrupts fire at 1 kHz or slower. You measure a dispatch that took 22 µs but your profiler reports 1,022 µs. That hurts. The profiler looks right — it logged a valid timestamp — but the number is meaningless for micro-tuning. How do you catch this? Look for timestamp clustering around multiples of the interrupt interval. If your distribution shows spikes at every 1,000 µs or 16 ms boundaries, you're not measuring dispatch latency. You're measuring interrupt latency. Switch to GPU-side timers (timestamp queries, CUDA events, or Vulkan query pools) that sample inside the GPU command processor, before the interrupt path touches the result. The difference between host-timer and GPU-timer values is your interrupt tax.

One more twist: modern GPUs coalesce multiple completion signals into a single interrupt. So your second and third dispatches may show identical timestamps. Wrong order. Not real. Your profiler displays them as simultaneous — which violates causality if those dispatches depend on each other's output. Detect coalescing by injecting a dummy dispatch with a known wait and checking whether the delta between successive end timestamps ever equals exactly zero. If it does, your profiler is lying to you in a way that looks perfectly correct.

Flag this for animation: shortcuts cost a day.

Flag this for animation: shortcuts cost a day.

“A profiler that reports 1 ms resolution for a 23 μs operation isn't wrong — it's just measuring something else entirely.”

— practical rule from a production GPU engineer, after debugging a false 40x latency spike report

GPU clock throttling from thermal or power limits

The tricky bit is thermal throttling hides inside your timing data like a ghost. Your profiler reports dispatch latency growing 2× over a ten-minute run. You blame the scheduler, the kernel, the memory layout. But the GPU simply got hot. Boost clocks drop from 1,900 MHz to 1,200 MHz as the die hits 85 °C. The dispatch itself doesn't slow down — the clock domain that drives command processing does. How to detect this? Graph latency against elapsed run time, not iteration number. If you see a smooth upward slope that stabilizes after the fan spins up, that's thermal equilibrium, not a regression. Add a 30-second idle cooldown between heavy dispatch bursts and compare the first and last results. If they differ by more than 5%, your benchmark is cooking itself. Most teams skip this: they assume constant clocks because the profiler didn't scream an error. The profiler doesn't know the chip is hot. It just returns numbers. You have to look for the drift.

Power limits bite faster than thermal limits. Nvidia's driver can cap power within a single dispatch queue if the board detects transient overshoot. That creates a 100–500 µs stall while the voltage regulator recovers. Your profiler logs the dispatch as completed — it was, technically — but the stall was invisible because the driver never reported it as a separate event. The tell is a bimodal latency distribution: one mode at normal latency, another mode 200–400 µs higher. If you see that second hump in every batch of 100 dispatches, suspect power throttling. Run with nvidia-smi -pl 300 (or your vendor's equivalent) to raise the limit and watch whether the second hump disappears. That's your confirmation. And a reminder that what looks like a profiler bug is often a hardware constraint the tool never promised to reveal.

Frequently Asked Questions About Dispatch Latency Measurement

Can I trust the GPU timestamp counter after a warm-up kernel?

Short answer: not blindly. The GPU timestamp counter—typically read via CLOCK_MONOTONIC_RAW equivalence on Vulkan or GL_TIMESTAMP in OpenGL—drifts between execution units. After a warm-up kernel, the clock domain might synchronize, but I have seen cases where the first three dispatches after a cold start report timestamps that are 40–60% shorter than stable runs. The warm-up heats the silicon, triggers voltage-frequency scaling, and—if you’re on a mobile GPU—throws in thermal throttling. The fix is simple: discard the first five to ten dispatch samples, then verify that the timestampPeriod reported by vkGetPhysicalDeviceProperties actually matches the wall-clock behavior. One team I worked with skipped this and reported 200 ns dispatch latency; after accounting for a stale counter, the real number was 1.1 µs. The catch is that many profilers cache the first sample and let you down.

Why does my profiler show negative latency on some calls?

Negative latency isn’t time travel—it’s a sign you’re reading the wrong counter or misaligning queue submissions. This happens when the host-side timer and the GPU timer start from different epochs. Windows QueryPerformanceCounter has a different base than the GPU’s internal cycle counter. So one dispatch appears to finish before it started. The odd part is—NVIDIA’s nvidia-smi and AMD’s ROCm both have documented offsets for this, but most profilers ignore them. To fix it, correlate a known-large dispatch (say, 10 ms of compute) and compute the constant delta. If the profiler still shows negative numbers, switch to CPU-side timed polling with a monotonic clock—the GPU timestamp can wait. A colleague once spent three days blaming the driver for negative values. Wrong order. The profiler just didn’t apply the timestampValidBits mask.

Polling the timestamp buffer from the CPU after every dispatch is like checking if toast is done by opening the oven door—your latency measurement becomes half measurement, half disturbance.

— A respiratory therapist, critical care unit

— paraphrased from a GPU hardware engineer, 2023 conference talk

Should I use polling or interrupts for low-latency dispatch?

Polling gives you microsecond precision but burns a CPU core. Interrupts save power but add 5–15 µs of context-switch overhead. For measuring dispatch latency—where you care about the gap between submission and the GPU starting work—polling wins. Use a spin-loop on a memory-mapped fence value in a dedicated thread, and call clock_gettime on each iteration. That said, interrupts hide the real latency by swallowing the tail of the dispatch. Most teams skip this: they profile with a blocking vkQueueWaitIdle or cudaDeviceSynchronize, which serializes execution and inflates the apparent latency by 30–50%. The trade-off is real. If you need production-level profiling, a hybrid approach—poll for the start, then sleep for 100 µs bursts—keeps CPU load under 10% while catching 98% of dispatches. But for micro-benchmarks? Poll hard. Interrupts are decorative when you need numbers.

Next time you see a negative latency readout, don't patch the profiler—patch your measurement. Grab a monotonic clock, discard the first five runs, and poll the fence yourself. That will cost you an afternoon but save you from shipping a timing lie.

Next Steps: Designing Your Own Dispatch Latency Validation Suite

Set up automated regression tests with known kernels

Grab three kernels whose dispatch latency you already understand — a null dispatch (empty shader), a heavy compute shader that saturates the L2 cache, and one that deliberately stalls on global-memory reads. Write a test harness that fires these kernels in a tight loop, 1,000 dispatches each, and records the CPU-side wall-clock time per dispatch via std::chrono::high_resolution_clock or your platform’s equivalent. The catch is: do not average the results into a single number. Plot the distribution instead. If your profiler reports 12 µs but your plot shows a bimodal split — half the dispatches land at 8 µs, half at 18 µs — you’re seeing PCIe scheduling jitter, not true GPU dispatch latency. I have seen teams chase “optimizations” that only shifted the bimodal point. Wrong order.

Link your regression suite to a CI runner that flags any nightly build where the 95th percentile of dispatch latency shifts by more than 15 %. Use Google Benchmark as your scaffolding — its DoNotOptimize and ClobberMemory intrinsics prevent the compiler from eliding your dispatches. For Vulkan users, Vulkan Samples includes a micro-benchmark pattern you can fork directly. Automate it. A single manual run tells you nothing; a thousand runs across build variants tell you where your profiler is lying.

Cross-reference with oscilloscope measurements on PCIe riser

Yes, this sounds like overkill. But if your dispatches are pushing 100+ kernels per frame, a software-only profiler can drift by 30 % because of CPU frequency scaling or interrupt coalescing. Grab a cheap PCIe riser extender, solder a trigger wire onto the PERST# pin or use a lane-activity LED tap, and connect a \$200 scope. Fire a single dispatch and measure the time from host write to the doorbell register to the first completion signal coming back. The number you get is the ground truth — everything else is inference. That hurts, because it exposes how much of your profiler’s reported “dispatch latency” is actually driver overhead or queue-flush tail latency.

Most teams skip this step. They shouldn’t. The trade-off is time: setting up the scope, writing a minimal driver that pulses a GPIO on dispatch start, and correlating the capture takes an afternoon. However, once you have that waveform, you can calibrate every software timer in your stack against it. Publish that calibration curve alongside your benchmarks — a simple table in your repo’s README (e.g., “NVIDIA RTX 4090, driver 550.xx, our profiler overreads by 1.4 µs vs. oscilloscope at 200 dispatches/s”). Now your readers can decide whether that 1.4 µs matters for their workload.

“The oscilloscope doesn’t lie. It just shows you how much your profiler wanted to be right instead of being accurate.”

— systems engineer on a game engine team, after tracing a phantom 40 % dispatch overhead to a CPU power state transition they hadn’t locked

Publish your methodology so others can reproduce

Post your test harness, your kernel sources, your GPU clock-locking commands, and your oscilloscope trigger scripts. Use a tagged release with a minimal Makefile — no “it works on my machine” dependencies. The goal is not perfection; it’s auditability. If someone on Blender or a WebGPU runtime finds that your dispatches show 9 µs but theirs show 14 µs on the same hardware, they can diff your driver version, your PCIe generation, your power limit. That conversation moves the field forward. A profiler that hides its assumptions is decorative. A validation suite that exposes them is useful. Build the latter. You will find the former already littering your build logs.