Shader Flame Graph

August 7, 2026 · View on GitHub

A flame graph of where a shader — or a whole frame — spends its GPU work, broken down to the WGSL statement.

There are two views:

Shader Cost (modeled) — in the Inspect panel, under a shader module. One entry point at a time, showing the cost of a single invocation.
Shader Flame Graph — a Capture tab (the button next to Analyze Shaders). The whole frame: pass → pipeline → shader stage → statement.

Reading the graph

Width is cost, depth is nesting. The y-axis stacks entry point → called functions → loops → branches → statements, so a wide frame near the bottom is something the shader spends real time in, and the frames above it say why.

Click a frame to zoom into it; the breadcrumb above the graph walks back out. In the per-shader view, clicking a frame also jumps the shader editor to that line.

Colors say which kind of work dominates a frame:

Color	Dimension	What it counts
Blue	ALU	Plain arithmetic and logic
Orange	SFU	Transcendentals, `sqrt`, division — special-function-unit ops
Red	Texture	`textureSample` / `textureLoad` / `textureStore`
Purple	Memory	Storage and uniform buffer accesses

A shader that is mostly red is texture-bound; mostly purple is bandwidth-bound. That distinction usually matters more than the raw total.

Where the numbers come from

Three sources feed the graph, and they are not equally trustworthy:

Measured. Pass GPU durations come from timestamp queries, available when the capture was taken with Profile Passes enabled. Fragment invocation counts come from replaying the frame on the GPU (the Measure fragments button). Per-draw GPU times come from replaying each draw with timestamp queries (the Measure draw times button) — see Per-draw timing.

Exact. Vertex and compute invocation counts are read straight out of the captured draw and dispatch arguments — draw(vertexCount, instanceCount), dispatchWorkgroups(x, y, z) times the entry point's @workgroup_size.

Modeled. The per-invocation instruction mix comes from a static analysis of the WGSL AST: an estimated op cost per expression, multiplied by loop trip counts, summed up the call tree. This is what distributes a measured pass duration across the shaders inside it.

So in the frame view with pass timings available, the root and each pass width are real milliseconds; only the split within a pass is modeled. Without timestamps the graph falls back to modeled op units, which compare to each other but not to wall-clock time.

Fragment counts

Fragment shader invocation counts cannot be derived from a capture — only rasterization knows them. The Measure fragments button replays the frame on the DevTools GPU device with a stub fragment shader that counts coverage, the same machinery the overdraw view uses.

Those counts are fragments that reach rasterization: after culling, viewport, scissor and depth-clip, but before the depth test. They are therefore an upper bound on fragment shader invocations, and the gap between the count and reality is exactly the work early-Z is saving you.

Until you run the measurement, fragment stages appear as zero-width frames labelled invocation count unknown, rather than being filled in with a guess.

Per-draw timing

Measure draw times replays each draw — and each compute dispatch — on its own with timestamp queries, so every item's width becomes a measured number. Without it, a pass's measured milliseconds are divided among its draws by the model; with it, the model's only remaining job is splitting each draw's measured time across that draw's shader stages and statements.

Two encodings, picked automatically:

inside-passes — writeTimestamp() between draws in a single pass, keeping the pass structure intact. Needs Chrome's chromium-experimental-timestamp-query-inside-passes.
split-passes — one pass per draw. Portable, and measured to agree with the above to within 1% on desktop hardware, where an empty pass costs nothing measurable. On a tile-based (mobile) GPU each pass forces a tile flush, so treat those numbers as relative rather than absolute.

Three caveats, all surfaced in the UI:

Items are timed in isolation. They don't overlap or share state the way they do in a real pass, so the per-draw times do not sum to the pass's own measured duration — they overstate it. Compare draws against each other, not against the pass total. The root frame is labelled isolated draw time rather than GPU time for exactly this reason.
Depth starts cleared. Each replayed render pass gets fresh attachments, so depth written by earlier passes is absent and draws that the real frame hides behind others will measure more expensive than they are. The compute equivalent: a dispatch's buffers hold their captured contents, not what earlier passes in the frame would have written into them.
Timestamps are quantized (~1µs on the hardware tested, unchanged by disabling Dawn's quantization — it appears to be the hardware tick). A draw costing a microsecond is indistinguishable from noise in one shot, so cheap draws are automatically re-measured with the draw encoded many times and the total divided. That recovers draws down to ~100ns, at the cost of extra submissions.

Ablation: measuring individual statements

Per-draw timing says which draw is expensive. Ablation says which line is, by measuring the shader with progressively more of its body executed and differencing: the gap between the cut that stops before a statement and the one that stops after it is that statement's cost.

instrumentForAblation() in shader_ablation.js rewrites a shader so its entry point can be stopped short at run time:

struct WGPUInspectorAblation { cut : u32 }
@group(1) @binding(0) var<uniform> _wgpuInspectorAblation : WGPUInspectorAblation;

@fragment fn fsMain() -> @location(0) vec4f {
  if (_wgpuInspectorAblation.cut == 0u) { return vec4f(); }
  var acc = 0.0;
  if (_wgpuInspectorAblation.cut == 1u) { return vec4f(); }
  for (var i = 0u; i < n; i = i + 1u) {
    if (_wgpuInspectorAblation.cut == 2u) { return vec4f(); }
    acc = acc + shade(i);
    if (_wgpuInspectorAblation.cut == 3u) { return vec4f(); }   // end of body
  }
  if (_wgpuInspectorAblation.cut == 4u) { return vec4f(); }
  ...
}

Cut 1 and cut 4 bracket the whole loop — the guards inside the body don't match while cut is 4, so every iteration runs. Cuts 2 and 3 bracket one statement of one iteration. See Inside a loop.

The cut point is a uniform, not a compile-time constant. That detail is the whole design, and it is worth explaining.

Why not just compile a truncated shader per cut?

Because the driver deletes the work you are trying to measure. Both variants were measured on real hardware (ablation_harness.html), against a shader whose true cost is 0.49 ms:

Approach	Cost curve across cuts	Signal
Truncate, return a constant	0.00205, 0.00205, 0.00205, 0.00205, 0.00205	100% lost
Truncate, return a sink value	0.00205, 0.00205, 0.10035, 0.10035, 0.49050	preserved
Uniform guard (this design)	0.00307, 0.00291, 0.10035, 0.10026, 0.49152	preserved

With a constant return, nothing consumes the statements above the cut, so dead-code elimination removes them and the curve is perfectly flat — it measures nothing at all. Returning a sink that consumes the accumulated value keeps the work live and does work, but synthesizing that sink means knowing every in-scope variable and its type at every cut point.

The uniform guard needs no sink, because for a large cut the whole body runs and the compiler cannot prove any statement unreachable. It also compiles once instead of once per cut. Its early return uses WGSL's zero-value expression T(), which is valid for any constructible type, so only the return type's name has to be recovered.

Accuracy

On the same harness, against statements with a known 4:1 cost ratio:

Statement	Measured	Expected
`acc = acc + work(u.base * 8u)`	0.0974 ms	1×
`acc = acc + work(u.base * 32u)`	0.3913 ms	4× → ratio 4.02
trivial statements	±0.0002 ms	~0

Instrumentation overhead measured at 0.2% of the shader's cost.

Using it

Select a frame in the frame flame graph and press Measure statements. Clicking a shader-stage frame targets that stage; clicking a draw frame targets its fragment stage, which matters because an unmeasured fragment stage has zero width and is culled before you can click it. Compute stages are selected the same way — a dispatch's invocation count is exact, so its stage frame always has width.

Results appear below the graph, ranked by cost, with clickable line numbers that jump to the shader source. In grouped mode the first draw of the group stands in for the rest, and the header says so.

Vertex ablation stubs the fragment stage

Cutting a vertex shader short makes it return a degenerate position, which rasterizes nothing — so with the real fragment shader attached, the entire fragment cost lands on whichever vertex statement finally produces a valid position. Measured on the test shader, that misattributed 0.4915 ms of fragment work onto one vertex line.

Vertex ablation therefore swaps in a trivial fragment shader whose outputs match the attachment formats. The same shader then measures 0.0012 ms, which is the real vertex cost. This is reported as a note whenever it happens.

A side effect worth knowing: stubbing the fragment stage removes bindings only the original fragment shader used from an auto pipeline layout, so the replayed bind groups are filtered to the bindings the vertex stage actually reads — the same thing the overdraw replay does for the same reason.

Compute ablation

Dispatches ablate the same way, and the compute path is the simpler one: no attachments to rebuild, no second stage to stub, and a compute entry point returns void, so the injected early return is a bare return rather than a synthesized zero value.

Two things are specific to compute:

Buffer state is restored before every cut. A compute shader writes to the buffers it reads, so without a reset each cut would start from whatever the previous cut left behind — a data-dependent loop would take a different trip count at every cut and the differences would be meaningless. The captured bytes are re-uploaded before each measurement. Within a single measurement the repeats still run back to back, which is reported as a note when the dispatch is cheap enough to need them.
Indirect dispatches work, with the workgroup count coming from the captured argument buffer — so it is whatever those bytes held at capture time, not what the frame's earlier passes would have computed into them.

Barriers are safe to cut across. The cut point is a uniform, so the guard's condition is the same for every invocation in the workgroup: either they all return before the barrier or none of them do.

Measured on ablation_measure_harness.html against a compute shader with the same known 4:1 statement ratio as the render case, through both auto and explicit pipeline layouts and an indirect dispatch:

Statement	Measured	Expected
`acc = acc + work(u.base * 8u)`	0.0166 ms	1×
`acc = acc + work(u.base * 32u)`	0.0658 ms	4× → ratio 3.95
trivial statements	±0.0003 ms	~0

Inside a loop

A loop's own row tells you the loop is expensive, not which line in it is. Cut points therefore go inside loop bodies too, and the results panel indents them under their loop.

This works because a cut means different things at the two levels. A cut inside a body stops execution during the first iteration, so differencing two cuts in the same body gives the cost of one statement for one execution. A cut outside the loop lets every iteration run, because the guards inside the body never match — so the loop's own row still measures all of it. Neither level disturbs the other.

Two things follow from that:

The last statement in a body needs a partner. The next cut in source order sits after the loop, and differencing against it would compare one partial iteration against all of them. A synthetic cut is injected at the end of every body to difference against instead.
Bodies that can break, continue or return are left whole, because control could skip that end-of-body marker, let the loop keep running, and turn the difference into nonsense. This is reported when it happens.

Per-execution costs are then scaled to a share of the loop's own measured time, so a body's rows sum to their loop rather than double-counting against it. The raw per-execution figure and the share are in each row's tooltip.

Measured against a loop running 16 times whose first body statement does 4× the work of its second:

	Per execution	Share of the loop	Distributed
`acc = acc + work(u.base * 2u)`	0.024576 ms	80%	0.3932 ms
`acc = acc + work(u.base / 2u)`	0.006144 ms	20%	0.0983 ms
the loop itself			0.4915 ms

The guards inside the body run every iteration, which is the price of the finer breakdown. On the harness shader it did not move the total at all — the uniform load hoists out of the loop and the compare costs nothing next to the body — but a loop with a trivial body will pay proportionally more.

Only loop bodies are descended into. An if body would be sound to cut the same way, but the invocations that don't take the branch run the entire shader in every one of those measurements, so the statement's cost becomes a small difference between two large numbers and drowns in the noise.

Limits

Branches are not broken down. A statement inside an if or switch is attributed to the enclosing statement, for the reason just given.
Negative costs happen. A statement cheaper than the measurement noise floor can difference to below zero. Those are reported as too small to measure rather than clamped to zero, because clamping would quietly turn noise into a confident-looking figure. A nested statement that measures negative gets no share of its loop rather than eating into its siblings'.
Needs a spare bind group. The uniform has to live somewhere; a shader already using all four bind groups can't be instrumented, and says so.
Nesting backs off when it has to. Past 64 cut points the sweep costs more than the detail is worth, and a nested shader the driver rejects (WGSL's uniformity rules around barriers are the likely cause) falls back to the top-level-only instrumentation. Both cases are reported rather than silently changing what you are looking at.

What the model cannot see

The cost model reads source, not machine code. It does not know about:

The driver's optimizer. Dead code, common subexpressions and constant folding all still show up as cost.
Occupancy and register pressure. A shader that spills is far slower than its instruction count suggests.
Divergence. Costs are per-invocation averages. An if is charged as the average of its arms (the expected cost of a uniform branch), so a branch that diverges across a warp is under-counted.
Memory locality. Every buffer access costs the same whether it hits cache or not.

It also has to assume things the source doesn't state:

Unbounded loops. while, loop, and for loops whose bounds depend on buffer data are assumed to run 8 times. Every such assumption is listed under the graph and marks its frames as estimated.
Indexed draws. drawIndexed(indexCount, …) counts indexCount vertex invocations, which is an upper bound — the GPU reuses post-transform vertices for repeated indices.
Indirect draws and dispatches. Their counts live in a GPU buffer and are read at execution time, so they have no invocation count and are shown unweighted, with a note.

Loop bounds are derived exactly when they can be: literal bounds, module-scope const and override values, and the usual i += k / i = i + k increment forms all resolve to a real trip count.

Practical use

The graph is best at answering "which shader, and which part of it, is worth looking at" — not at predicting absolute timings. A reasonable loop:

Capture with Profile Passes enabled so the graph is in milliseconds.
Open Shader Flame Graph and find the widest pass.
Run Measure fragments if a render pass looks suspiciously cheap — an unweighted fragment stage is invisible until you do.
Run Measure draw times to replace the modeled split between draws and dispatches with measurement. Tick Per draw first if you want individual items rather than per-pipeline groups.
Zoom into the widest shader stage and look at the color: red says cut texture work, purple says cut buffer traffic, orange says a transcendental is in a loop.
Press Measure statements on that stage to find the line inside it. This works for vertex, fragment and compute stages.
Cross-check against Analyze Shaders, which flags specific patterns (expensive builtins in loops, loop-invariant expressions) on the same lines.

Note that Per draw re-buckets the frame, which invalidates fragment counts measured in the other mode (a per-draw bucket is a subset of a pipeline one, so old counts can't be re-attributed). The graph clears them and says so; measure again in the mode you want. Draw timings are keyed by the command itself and survive the toggle.