3.3 GPU 计时与优化闭环
前两节都在量 CPU。但一帧到底卡在谁身上，不看 GPU 是说不清的。难点在于：GPU 计时不像 performance.now() 那样随手可得，它依赖 timestamp-query 特性，而不同实现的支持度天差地别。这一节的做法是先探测、再量、再交叉印证，并给出一条不依赖 GPU 的替代测量路径。
这是全书第 3 章第 3 节，也是本章的收官。测法与闭环用 TypeScript / Node 26 直跑；GPU 计时走 headless Chromium（L2），原始数字来自本机实跑（darwin/arm64，Node 26.5.0）。
timestamp-query：能拿到，但要看它可不可信
WebGPU 的 timestamp-query 特性允许在 pass 两端写入时间戳，resolveQuerySet 后回读即可得到 GPU 侧纳秒。但特性存在不等于数值可信：不同后端、不同驱动对时间戳的实现质量差异很大。所以我们的门禁不只问「有没有」，还问「两种测法对不对得上」：
- 时间戳法：
timestampWrites写 pass 首尾两个时间戳，回读求差。 - 墙钟法：
performance.now()包住submit，再await queue.onSubmittedWorkDone()，得到「提交→GPU 完成」的墙钟。
两者测的是同一段 GPU 工作，若数量级一致，说明时间戳可用。本机 headless 用的是 SwiftShader 软件光栅化（adapter {vendor:"google", arch:"swiftshader"}），所以：这里的绝对毫秒数只反映软件光栅化，绝不代表真实 GPU 性能——真实设备的 GPU 计时属 L4，本节需在真实设备 / 浏览器中手动验证，命令行不可复现。本节在命令行里验证的是测量机制本身与两种测法的一致性。
可运行产物：L2 GPU 计时门禁
下面的门禁起一个 loopback http 服务，在 headless Chromium 里跑：对 4 个不同计算量的 compute pass，各重复 5 次取最小/中位，同时给出时间戳与墙钟两条读数；末尾再测一遍「连续提交空 render pass」的 CPU 提交吞吐。
// file: tools/gpu-timing-l2.mjs
// tier: L2
// verify: node tools/gpu-timing-l2.mjs
//
// helio 工程卷 3.3：GPU 计时与优化闭环（L2，headless Chromium）。
// 1) 探测 timestamp-query 特性；2) 可用则在 compute pass 两端写时间戳，resolve 回读，
// 得 GPU 侧 ns；3) 无论时间戳是否可用，都给出可复现的替代量：CPU 提交耗时 + 提交→
// GPU 完成墙钟。每负载重复 5 次取最小/中位，暴露单次采样噪声。
// 注意：headless 用 SwiftShader 软件光栅化，绝对耗时**不代表真实 GPU**（属 L4），
// 这里验证的是「测量机制」本身与两种测法的一致性。
import http from 'node:http';
import { writeFileSync, mkdirSync } from 'node:fs';
import { chromium } from '/opt/homebrew/lib/node_modules/playwright/index.mjs';
const PAGE = `<!doctype html><meta charset=utf-8><canvas id="c" width=64 height=64></canvas>
<script type="module">
const med=(a)=>{const s=[...a].sort((x,y)=>x-y);const m=s.length>>1;return s.length%2?s[m]:(s[m-1]+s[m])/2;};
const mn=(a)=>Math.min(...a);
window.__run = async () => {
const out = { workloads: [], cpu: {} };
const adapter = await navigator.gpu.requestAdapter();
out.adapter = adapter.info ? { vendor: adapter.info.vendor, arch: adapter.info.architecture } : null;
out.hasTS = adapter.features.has('timestamp-query');
const device = await adapter.requestDevice({ requiredFeatures: out.hasTS ? ['timestamp-query'] : [] });
const N = 262144;
const data = new Float32Array(N); for (let i = 0; i < N; i++) data[i] = i % 97;
const buf = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
device.queue.writeBuffer(buf, 0, data);
const mkPipe = (iter) => device.createShaderModule({ code: \`
@group(0) @binding(0) var<storage, read_write> d: array<f32>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) g: vec3<u32>) {
let i = g.x; if (i >= arrayLength(&d)) { return; }
var v = d[i];
for (var k = 0u; k < \${iter}u; k = k + 1u) { v = v * 1.0000001 + 0.5; }
d[i] = v;
}\` });
for (const iter of [1, 64, 512, 4096]) {
const pipe = device.createComputePipeline({ layout: 'auto', compute: { module: mkPipe(iter), entryPoint: 'main' } });
const bg = device.createBindGroup({ layout: pipe.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: buf } }] });
const cpuS = [], wallS = [], tsS = [];
for (let rep = 0; rep < 5; rep++) {
const qs = out.hasTS ? device.createQuerySet({ type: 'timestamp', count: 2 }) : null;
const rb = out.hasTS ? device.createBuffer({ size: 16, usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC }) : null;
const rr = out.hasTS ? device.createBuffer({ size: 16, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ }) : null;
const enc = device.createCommandEncoder();
const pass = enc.beginComputePass(out.hasTS ? { timestampWrites: { querySet: qs, beginningOfPassWriteIndex: 0, endOfPassWriteIndex: 1 } } : {});
pass.setPipeline(pipe); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(Math.ceil(N / 64)); pass.end();
if (out.hasTS) { enc.resolveQuerySet(qs, 0, 2, rb, 0); enc.copyBufferToBuffer(rb, 0, rr, 0, 16); }
const t0 = performance.now();
device.queue.submit([enc.finish()]);
cpuS.push(performance.now() - t0);
await device.queue.onSubmittedWorkDone();
wallS.push(performance.now() - t0);
if (out.hasTS) { await rr.mapAsync(GPUMapMode.READ); const ts = new BigUint64Array(rr.getMappedRange().slice(0)); rr.unmap(); tsS.push(Number(ts[1] - ts[0]) / 1e6); }
}
out.workloads.push({ iter, cpuSubmitMinMs: +mn(cpuS).toFixed(4),
gpuTsMinMs: out.hasTS ? +mn(tsS).toFixed(3) : null, gpuTsMedMs: out.hasTS ? +med(tsS).toFixed(3) : null,
submitToDoneMinMs: +mn(wallS).toFixed(3), submitToDoneMedMs: +med(wallS).toFixed(3) });
}
const fmt = navigator.gpu.getPreferredCanvasFormat();
const tex = device.createTexture({ size: [64, 64], format: fmt, usage: GPUTextureUsage.RENDER_ATTACHMENT });
const FRAMES = 300, t0 = performance.now();
for (let f = 0; f < FRAMES; f++) {
const enc = device.createCommandEncoder();
const p = enc.beginRenderPass({ colorAttachments: [{ view: tex.createView(), clearValue: { r: 0, g: 0, b: 0, a: 1 }, loadOp: 'clear', storeOp: 'store' }] });
p.end(); device.queue.submit([enc.finish()]);
}
await device.queue.onSubmittedWorkDone();
const ms = performance.now() - t0;
out.cpu = { frames: FRAMES, totalMs: +ms.toFixed(1), submitFps: +(FRAMES / ms * 1000).toFixed(0) };
return out;
};
</script>`;
const srv = http.createServer((_q, res) => {
res.setHeader('content-type', 'text/html; charset=utf-8');
res.setHeader('Cross-Origin-Opener-Policy', 'same-origin');
res.setHeader('Cross-Origin-Embedder-Policy', 'require-corp');
res.end(PAGE);
});
await new Promise((r) => srv.listen(0, '127.0.0.1', r));
const browser = await chromium.launch({ headless: true, args: ['--enable-unsafe-webgpu', '--use-webgpu-adapter=swiftshader', '--enable-unsafe-swiftshader'] });
try {
const page = await browser.newPage({ viewport: { width: 64, height: 64 } });
await page.goto(`http://127.0.0.1:${srv.address().port}/`);
const r = await page.evaluate(() => window.__run());
console.log(`timestamp-query 可用: ${r.hasTS} adapter: ${JSON.stringify(r.adapter)}`);
console.log('iter cpuSubmit(min,ms) gpuTs(min,ms) gpuTs(med,ms) submit→done(min,ms) submit→done(med,ms)');
for (const w of r.workloads) console.log(
`${String(w.iter).padStart(4)} ${String(w.cpuSubmitMinMs).padStart(16)} ${String(w.gpuTsMinMs).padStart(13)} ` +
`${String(w.gpuTsMedMs).padStart(13)} ${String(w.submitToDoneMinMs).padStart(19)} ${String(w.submitToDoneMedMs).padStart(19)}`);
console.log(`\n空 render pass 提交吞吐：${r.cpu.frames} 帧 / ${r.cpu.totalMs}ms = ${r.cpu.submitFps} fps（无 vsync，仅测 CPU 提交吞吐）`);
mkdirSync('/tmp/eb-eng/ch3/artifacts', { recursive: true });
writeFileSync('/tmp/eb-eng/ch3/artifacts/gpu-timing-l2.json', JSON.stringify(r, null, 2));
} finally {
await browser.close();
srv.close();
}
验证命令：node tools/gpu-timing-l2.mjs。
可运行产物：优化闭环（B）
拿到数字只是开始，闭环还得有「改完再测」。下面这个基准把CPU 侧命令编码当被测对象：同一批 20 万个渲染项，naive 每项一次 draw，sorted 只排序不实例化，batched 先按键合批、每批一次实例化 draw。三种写法都用真实的 CommandListRecorder 与 radixSortKeys，量的是真开销。
// file: tools/opt-loop.mjs
// tier: B
// verify: node tools/opt-loop.mjs
//
// helio 工程卷 3.3 基准：优化闭环（测 → 改 → 再测）。
// 被测对象是「CPU 侧命令编码」：同一批渲染项，
// naive ：每项一次 draw（setPipeline + setBindGroup + draw）；
// batched：先按键排序合批，每批一次实例化 draw。
// 命令条数与编码耗时都用真实对象（CommandListRecorder / radixSortKeys）量出来。
import { CommandListRecorder } from '../src/render/rhi/command-recorder.ts';
import { makeBatchItems, packBatchKey, buildBatches, summarize } from '../src/render/sort/batch.ts';
import { radixSortKeys } from '../src/render/sort/sort-key.ts';
const N = 200000;
const MESHES = 64;
const REPS = 20;
function encodeNaive(items) {
const rec = new CommandListRecorder();
rec.beginRenderPass('main');
for (const it of items) {
rec.setPipeline(`p${it.pipelineId}`);
rec.setBindGroup(0, `m${it.meshId}`);
rec.draw(36);
}
rec.endRenderPass();
return rec.finish().length;
}
function encodeSorted(items, order) {
const rec = new CommandListRecorder();
rec.beginRenderPass('main');
for (const idx of order) {
const it = items[idx];
rec.setPipeline(`p${it.pipelineId}`);
rec.setBindGroup(0, `m${it.meshId}`);
rec.draw(36);
}
rec.endRenderPass();
return rec.finish().length;
}
function encodeBatched(draws) {
const rec = new CommandListRecorder();
rec.beginRenderPass('main');
for (const d of draws) {
rec.setPipeline(`p${d.pipelineId}`);
rec.setBindGroup(0, `m${d.meshId}`);
rec.draw(d.instanceCount);
}
rec.endRenderPass();
return rec.finish().length;
}
const bench = (fn) => {
fn(); // 热身
const s = [];
for (let r = 0; r < REPS; r++) {
const t = performance.now();
fn();
s.push(performance.now() - t);
}
s.sort((a, b) => a - b);
return s[0];
};
const items = makeBatchItems(N, 21, MESHES, 0);
const keys = Uint32Array.from(items, packBatchKey);
const order = radixSortKeys(keys, N);
const draws = buildBatches(items, order);
const bsum = summarize(draws);
const naiveCmds = encodeNaive(items);
const sortedCmds = encodeSorted(items, order);
const batchedCmds = encodeBatched(draws);
const naiveMs = bench(() => encodeNaive(items));
const sortedMs = bench(() => encodeSorted(items, order));
const batchedMs = bench(() => encodeBatched(draws));
console.log('=== 优化闭环：命令编码（CPU 侧，纯 L1）===');
console.log(`渲染项 ${N}，网格种类 ${MESHES}，重复 ${REPS} 次取最小`);
console.log('方案 drawCalls 命令条数 编码耗时(ms) 每千项(us)');
for (const [name, dc, cmds, ms] of [
['naive', N, naiveCmds, naiveMs],
['sorted', N, sortedCmds, sortedMs],
['batched', bsum.drawCalls, batchedCmds, batchedMs],
]) {
console.log(
`${name.padEnd(8)} ${String(dc).padStart(9)} ${String(cmds).padStart(9)} ${ms.toFixed(2).padStart(11)} ${((ms / N) * 1000).toFixed(2).padStart(9)}`,
);
}
console.log(
`\n闭环结论：只排序不实例化（sorted）编码耗时 ${sortedMs.toFixed(2)}ms，反而比 naive 更慢——` +
`命令条数没变，还多了一层「按 order 下标回查 items」的间接寻址；` +
`真正省命令的是合批：draw 调用 ${N} → ${bsum.drawCalls}（砍 ${(N / bsum.drawCalls).toFixed(0)}x），` +
`编码耗时 ${naiveMs.toFixed(2)}ms → ${batchedMs.toFixed(2)}ms，平均每批 ${bsum.avgInstances.toFixed(1)} 个实例。`,
);
验证命令：node tools/opt-loop.mjs。合批与实例化的渲染侧原理见卷二 helio-render-09-3-instancing-batching
。
真跑出来的两组数字与闭环结论
GPU 计时（headless SwiftShader，每负载 5 次取最小/中位）：
timestamp-query 可用: true adapter: {"vendor":"google","arch":"swiftshader"}
iter cpuSubmit(min,ms) gpuTs(min,ms) gpuTs(med,ms) submit→done(min,ms) submit→done(med,ms)
1 0.01 2.525 3.811 3.705 5.29
64 0.01 3.181 3.63 3.985 4.015
512 0.01 15.387 16.174 16.565 16.885
4096 0.01 113.239 125.844 114.18 126.18
空 render pass 提交吞吐：300 帧 / 6.9ms = 43228 fps（无 vsync，仅测 CPU 提交吞吐）
三点读数。第一，时间戳法（gpuTs）与墙钟法（submit→done）逐行吻合，差值都在 1~2ms 内，说明这台 headless 实现下的 timestamp-query 是可信的——测量机制成立。第二，负载从 iter=1 到 iter=4096，GPU 耗时从 2.5ms 单调涨到 113ms，说明时间戳确实跟着实际计算量走，而不是常量或噪声。第三，cpuSubmit 恒定在 0.01ms：提交本身几乎不花钱，贵的是 GPU 执行——这与前两节 CPU 侧的量级形成对照。再次强调：以上绝对毫秒来自软件光栅化，真实 GPU 数值须在真机复测（L4，命令行不可复现）。
优化闭环（CPU 侧命令编码）：
=== 优化闭环：命令编码（CPU 侧，纯 L1）===
渲染项 200000，网格种类 64，重复 20 次取最小
方案 drawCalls 命令条数 编码耗时(ms) 每千项(us)
naive 200000 600002 10.54 0.05
sorted 200000 600002 40.54 0.20
batched 64 194 0.01 0.00
闭环三步走得很清楚：测——naive 20 万次 draw、60 万条命令、编码 10.54ms；改——按键合批，把 20 万次 draw 并成 64 次实例化 draw、194 条命令；复测——编码降到 0.01ms（已贴近 performance.now() 下限）。中间的 sorted 是个反例：它只排序、不实例化，命令条数一条没少，还多了一层「按 order 下标回查 items」的间接寻址，编码反而涨到 40.54ms——排序本身不省命令，真正省的是合批。这正是闭环的价值：不看复测数字，很容易把「排了序」误当成「优化过」。本章到此收束，下一章转向跨端：同样的引擎，如何在 WebGPU 与 WebGL2 之间做能力探测与降级。
阅读导航：上一节：3.2 GC、内存预算与分配追踪 · 下一节：4.1 WebGPU / WebGL2 能力探测与降级 。
继续阅读
探索更多技术文章
浏览归档，发现更多关于系统设计、工具链和工程实践的内容。