8.3 GPU 驱动剔除与 LOD 协同（compute + indirect）
前两节的剔除都在 CPU 上：CPU 遍历 BVH、算出可见列表、再逐批 draw。这一节把剔除与 LOD 选择整个搬到 GPU——一个 compute 通道读场景、自己判断谁可见、按距离选档，用原子计数器把存活对象追加进对应 LOD 的实例缓冲；第二个通道写出三组 indirect 绘制参数；CPU 只在最后发三个 drawIndirect，全程不碰实例列表。
剔除与选档写进 compute，绘制只用 indirect
数据流：宿主上传对象中心与半径、相机矩阵与六个平面、两个 LOD 距离阈值。cullMain 每个线程处理一个对象——用六个平面对包围球做剔除（球心到平面距离 < -r 即完全在外），存活的按 distance 选 0/1/2 档，atomicAdd 拿到本档的写入槽位，把中心、半径、档位写进该档的实例缓冲。finalize 单线程读计数器，把三档的 vertexCount（10/6/5）与 instanceCount 写进 indirect 缓冲。渲染时三档共用一份几何，靠 firstVertex / vertexCount 选档。着色器 shaders/gpu-driven-cull-lod.wgsl：
// file: shaders/gpu-driven-cull-lod.wgsl
// tier: L2
// verify: node tools/run-l2.mjs src/render/culling/gpu-driven-cull-lod.ts
//
// 8.3 的 L2 验收：把「剔除 + LOD 选择」全部搬到 GPU。
// 一个 compute 通道逐对象做包围球-视锥剔除并按距离选 LOD，用原子计数器把
// 存活对象追加进对应 LOD 的实例缓冲；第二个通道据计数器写出 3 组 indirect 绘制参数。
// 渲染阶段只用 drawIndirect 画三档 LOD，CPU 全程不碰实例列表。
// 每档 LOD 一个独立实例缓冲（偏移 0），不依赖 firstInstance 的实例基址。
// MAX_OBJ 必须与宿主一致（4096）。
const MAX_OBJ: u32 = 4096u;
struct Camera {
vp: mat4x4<f32>,
planes: array<vec4<f32>, 6>,
camPos: vec4<f32>, // xyz = 相机世界坐标
lodDist: vec4<f32>, // x = LOD0 距离上限，y = LOD1 距离上限
};
struct Obj { center: vec4<f32> }; // xyz = 中心，w = 半径
struct IndirectArgs {
vertexCount: u32,
instanceCount: u32,
firstVertex: u32,
firstInstance: u32,
};
@group(0) @binding(0) var<uniform> cam: Camera;
@group(0) @binding(1) var<storage, read> objs: array<Obj>;
@group(0) @binding(2) var<storage, read_write> lodOf: array<u32>;
@group(0) @binding(3) var<storage, read_write> counters: array<atomic<u32>, 3>;
@group(0) @binding(4) var<storage, read_write> indirect: array<IndirectArgs, 3>;
@group(0) @binding(5) var<storage, read_write> inst0: array<f32>;
@group(0) @binding(6) var<storage, read_write> inst1: array<f32>;
@group(0) @binding(7) var<storage, read_write> inst2: array<f32>;
@compute @workgroup_size(64)
fn cullMain(@builtin(global_invocation_id) gid: vec3<u32>) {
let i = gid.x;
if (i >= arrayLength(&objs)) { return; }
let c = objs[i].center;
var vis = true;
for (var p = 0u; p < 6u; p = p + 1u) {
let pl = cam.planes[p];
if (dot(pl.xyz, c.xyz) + pl.w < -c.w) { vis = false; }
}
if (!vis) { lodOf[i] = 0xFFFFFFFFu; return; }
let d = distance(cam.camPos.xyz, c.xyz);
var lod = 0u;
if (d > cam.lodDist.y) { lod = 2u; } else if (d > cam.lodDist.x) { lod = 1u; }
lodOf[i] = lod;
let slot = atomicAdd(&counters[lod], 1u);
let base = slot * 8u; // 每实例 8 个 float：中心 xyz + 半径 + lod + 3 个填充
if (lod == 0u) {
inst0[base + 0u] = c.x; inst0[base + 1u] = c.y; inst0[base + 2u] = c.z; inst0[base + 3u] = c.w;
inst0[base + 4u] = 0.0; inst0[base + 5u] = 0.0; inst0[base + 6u] = 0.0; inst0[base + 7u] = 0.0;
} else if (lod == 1u) {
inst1[base + 0u] = c.x; inst1[base + 1u] = c.y; inst1[base + 2u] = c.z; inst1[base + 3u] = c.w;
inst1[base + 4u] = 1.0; inst1[base + 5u] = 0.0; inst1[base + 6u] = 0.0; inst1[base + 7u] = 0.0;
} else {
inst2[base + 0u] = c.x; inst2[base + 1u] = c.y; inst2[base + 2u] = c.z; inst2[base + 3u] = c.w;
inst2[base + 4u] = 2.0; inst2[base + 5u] = 0.0; inst2[base + 6u] = 0.0; inst2[base + 7u] = 0.0;
}
}
// 单线程收尾：据计数器把三档 LOD 的绘制参数写进 indirect 缓冲。
@compute @workgroup_size(1)
fn finalize() {
for (var l = 0u; l < 3u; l = l + 1u) {
var a: IndirectArgs;
a.instanceCount = atomicLoad(&counters[l]);
a.firstInstance = 0u;
if (l == 0u) { a.firstVertex = 0u; a.vertexCount = 10u; }
else if (l == 1u) { a.firstVertex = 10u; a.vertexCount = 6u; }
else { a.firstVertex = 16u; a.vertexCount = 5u; }
indirect[l] = a;
}
}
// ---- 渲染：三档 LOD 共用一份几何，靠 firstVertex/vertexCount 选档 ----
struct VSOut {
@builtin(position) pos: vec4<f32>,
@location(0) color: vec3<f32>,
};
fn lodColor(lod: u32) -> vec3<f32> {
if (lod == 0u) { return vec3<f32>(0.9, 0.2, 0.2); }
if (lod == 1u) { return vec3<f32>(0.2, 0.9, 0.2); }
return vec3<f32>(0.2, 0.2, 0.9);
}
@vertex
fn vs(
@location(0) corner: vec2<f32>,
@location(1) center: vec3<f32>,
@location(2) radius: f32,
@location(3) lodF: f32, // 实例缓冲里存的是 float，故按 float 读再转 u32
) -> VSOut {
var o: VSOut;
let p = center + vec3<f32>(corner * radius, 0.0);
o.pos = cam.vp * vec4<f32>(p, 1.0);
o.color = lodColor(u32(lodF));
return o;
}
@fragment
fn fs(in: VSOut) -> @location(0) vec4<f32> {
return vec4<f32>(in.color, 1.0);
}
可运行产物：4096 对象、三档 LOD
src/render/culling/gpu-driven-cull-lod.ts（L2）在 CPU 侧用同一套判定算出基准，再与 GPU 的 lodOf 掩码逐对象比对，并核对 indirect 参数与像素：
// file: src/render/culling/gpu-driven-cull-lod.ts
// tier: L2
// verify: node tools/run-l2.mjs src/render/culling/gpu-driven-cull-lod.ts
//
// 8.3 的 L2 验收：GPU 驱动剔除 + LOD 协同。
// compute 通道逐对象做包围球-视锥剔除、按距离选 LOD，用原子计数器把存活对象
// 追加进对应 LOD 的实例缓冲；收尾通道写出 3 组 indirect 参数；渲染只用 drawIndirect。
// 回读 lodOf 掩码与 CPU 逐对象结果比对，并核对 indirect 参数与像素。
// 需 headless Chromium + 本地 http 服务，命令行不可复现。
import { extractFrustumPlanes, lookAt, multiply, perspective, sphereInFrustum, type Frustum, type Vec3 } from './frustum.ts';
const N = 4096, W = 512, H = 512, MAX_OBJ = 4096;
const LOD_DIST = [8, 20] as const; // LOD0 < 8 <= LOD1 < 20 <= LOD2
const CULLED = 0xffffffff;
const CLEAR: readonly [number, number, number] = [0.05, 0.05, 0.1];
export interface GpuDrivenResult {
readonly objects: number;
readonly cpuVisible: number;
readonly gpuVisible: number;
readonly lodMismatch: number;
readonly perLod: readonly [number, number, number];
readonly indirectCounts: readonly [number, number, number];
readonly vertexCounts: readonly [number, number, number];
readonly lodPixels: readonly [number, number, number];
readonly wgslErrors: number;
}
/** 三档 LOD 的三角扇几何：8/4/3 个三角形，顶点数 10/6/5，拼进同一顶点缓冲。 */
function discGeometry(): Float32Array {
const out: number[] = [];
for (const t of [8, 4, 3]) {
out.push(0, 0);
for (let k = 0; k <= t; k++) out.push(Math.cos((k * 2 * Math.PI) / t), Math.sin((k * 2 * Math.PI) / t));
}
return new Float32Array(out);
}
export async function runL2(): Promise<{ checks: { name: string; ok: boolean; detail: string }[]; data: GpuDrivenResult }> {
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) throw new Error('拿不到 WebGPU adapter');
const device = await adapter.requestDevice();
const code = await (await fetch('/shaders/gpu-driven-cull-lod.wgsl')).text();
const module = device.createShaderModule({ code });
const info = await module.getCompilationInfo();
const wgslErrors = info.messages.filter((m) => m.type === 'error').length;
// ---- 场景与 CPU 基准 ----
let s = 5 >>> 0;
const next = (): number => ((s = (Math.imul(s, 1664525) + 1013904223) >>> 0) / 4294967296);
const objs = new Float32Array(N * 4);
for (let i = 0; i < N; i++) objs.set([(next() - 0.5) * 40, (next() - 0.5) * 40, (next() - 0.5) * 40, 0.6], i * 4);
const eye: Vec3 = [0, 0, 0];
const vp = multiply(perspective((60 * Math.PI) / 180, 1, 0.1, 40), lookAt(eye, [0, 0, -1], [0, 1, 0]));
const f: Frustum = extractFrustumPlanes(vp);
const expected = new Uint32Array(N);
const expPerLod = [0, 0, 0];
for (let i = 0; i < N; i++) {
const c: Vec3 = [objs[i * 4]!, objs[i * 4 + 1]!, objs[i * 4 + 2]!];
if (!sphereInFrustum({ center: c, radius: objs[i * 4 + 3]! }, f)) { expected[i] = CULLED; continue; }
const d = Math.hypot(c[0], c[1], c[2]);
const lod = d > LOD_DIST[1] ? 2 : d > LOD_DIST[0] ? 1 : 0;
expected[i] = lod; expPerLod[lod]!++;
}
// ---- 缓冲 ----
const mk = (bytes: number, usage: GPUBufferUsageFlags, data?: Float32Array | Uint32Array): GPUBuffer => {
const b = device.createBuffer({ size: Math.max(4, bytes), usage: usage | GPUBufferUsage.COPY_DST });
if (data) device.queue.writeBuffer(b, 0, data);
return b;
};
const R = GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC;
const objBuf = mk(N * 16, GPUBufferUsage.STORAGE, objs);
const lodOfBuf = mk(N * 4, R);
const inst = [0, 1, 2].map(() => mk(MAX_OBJ * 32, GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX));
const cntBuf = mk(16, R, new Uint32Array(4));
const indBuf = mk(3 * 16, GPUBufferUsage.STORAGE | GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_SRC, new Uint32Array(12));
const geo = discGeometry();
const geoBuf = mk(geo.byteLength, GPUBufferUsage.VERTEX, geo);
const cam = new Float32Array(48);
cam.set(vp, 0);
f.forEach((p, i) => cam.set(p, 16 + i * 4));
cam.set([...eye, 0], 40);
cam.set([LOD_DIST[0], LOD_DIST[1], 0, 0], 44);
const camBuf = mk(192, GPUBufferUsage.UNIFORM, cam);
const comp = (binding: number, type: GPUBufferBindingType) => ({ binding, visibility: GPUShaderStage.COMPUTE, buffer: { type } });
const bindLayout = device.createBindGroupLayout({ entries: [
{ binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform', minBindingSize: 192 } },
comp(1, 'read-only-storage'), comp(2, 'storage'), comp(3, 'storage'), comp(4, 'storage'),
comp(5, 'storage'), comp(6, 'storage'), comp(7, 'storage'),
] });
const pl = device.createPipelineLayout({ bindGroupLayouts: [bindLayout] });
const bind = device.createBindGroup({ layout: bindLayout, entries: [
{ binding: 0, resource: { buffer: camBuf } },
{ binding: 1, resource: { buffer: objBuf } },
{ binding: 2, resource: { buffer: lodOfBuf } },
{ binding: 3, resource: { buffer: cntBuf } },
{ binding: 4, resource: { buffer: indBuf } },
{ binding: 5, resource: { buffer: inst[0]! } },
{ binding: 6, resource: { buffer: inst[1]! } },
{ binding: 7, resource: { buffer: inst[2]! } },
] });
const pipe = (entryPoint: string) => device.createComputePipeline({ layout: pl, compute: { module, entryPoint } });
const cullPipe = pipe('cullMain'), finPipe = pipe('finalize');
const rpipe = device.createRenderPipeline({
layout: 'auto',
vertex: {
module, entryPoint: 'vs',
buffers: [
{ arrayStride: 8, attributes: [{ shaderLocation: 0, offset: 0, format: 'float32x2' }] },
{ arrayStride: 32, stepMode: 'instance', attributes: [
{ shaderLocation: 1, offset: 0, format: 'float32x3' },
{ shaderLocation: 2, offset: 12, format: 'float32' },
{ shaderLocation: 3, offset: 16, format: 'float32' },
] },
],
},
fragment: { module, entryPoint: 'fs', targets: [{ format: 'rgba8unorm' }] },
primitive: { topology: 'triangle-list', cullMode: 'none' },
});
const rbind = device.createBindGroup({ layout: rpipe.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: camBuf } }] });
const tex = device.createTexture({ size: [W, H], format: 'rgba8unorm', usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC });
const enc = device.createCommandEncoder();
const cp = enc.beginComputePass();
cp.setPipeline(cullPipe); cp.setBindGroup(0, bind); cp.dispatchWorkgroups(Math.ceil(N / 64)); cp.end();
const fp = enc.beginComputePass();
fp.setPipeline(finPipe); fp.setBindGroup(0, bind); fp.dispatchWorkgroups(1); fp.end();
const rp = enc.beginRenderPass({
colorAttachments: [{ view: tex.createView(), clearValue: { r: CLEAR[0], g: CLEAR[1], b: CLEAR[2], a: 1 }, loadOp: 'clear', storeOp: 'store' }],
});
rp.setPipeline(rpipe); rp.setBindGroup(0, rbind); rp.setVertexBuffer(0, geoBuf);
for (let l = 0; l < 3; l++) {
rp.setVertexBuffer(1, inst[l]!);
rp.drawIndirect(indBuf, l * 16);
}
rp.end();
const read = (src: GPUBuffer, bytes: number): GPUBuffer => {
const rb = device.createBuffer({ size: bytes, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
enc.copyBufferToBuffer(src, 0, rb, 0, bytes);
return rb;
};
const lodRb = read(lodOfBuf, N * 4), cntRb = read(cntBuf, 12), indRb = read(indBuf, 48);
const bpr = W * 4;
const pxRb = device.createBuffer({ size: bpr * H, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
enc.copyTextureToBuffer({ texture: tex }, { buffer: pxRb, bytesPerRow: bpr, rowsPerImage: H }, [W, H]);
device.queue.submit([enc.finish()]);
await Promise.all([lodRb.mapAsync(GPUMapMode.READ), cntRb.mapAsync(GPUMapMode.READ), indRb.mapAsync(GPUMapMode.READ), pxRb.mapAsync(GPUMapMode.READ)]);
const lodOf = new Uint32Array(lodRb.getMappedRange().slice(0));
const cnt = new Uint32Array(cntRb.getMappedRange().slice(0));
const ind = new Uint32Array(indRb.getMappedRange().slice(0));
const px = new Uint8Array(pxRb.getMappedRange().slice(0));
lodRb.unmap(); cntRb.unmap(); indRb.unmap(); pxRb.unmap();
let lodMismatch = 0, gpuVisible = 0;
for (let i = 0; i < N; i++) {
if (lodOf[i] !== expected[i]) lodMismatch++;
if (lodOf[i] !== CULLED) gpuVisible++;
}
const perLod: [number, number, number] = [cnt[0]!, cnt[1]!, cnt[2]!];
const indirectCounts: [number, number, number] = [ind[1]!, ind[5]!, ind[9]!];
const vertexCounts: [number, number, number] = [ind[0]!, ind[4]!, ind[8]!];
const lodColors = [[0.9, 0.2, 0.2], [0.2, 0.9, 0.2], [0.2, 0.2, 0.9]];
const lodPixels = lodColors.map((c) => {
const want = c.map((v) => Math.round(v * 255));
let n = 0;
for (let i = 0; i < px.length; i += 4) {
if (Math.abs(px[i]! - want[0]!) <= 2 && Math.abs(px[i + 1]! - want[1]!) <= 2 && Math.abs(px[i + 2]! - want[2]!) <= 2) n++;
}
return n;
}) as [number, number, number];
device.destroy();
const data: GpuDrivenResult = {
objects: N, cpuVisible: expPerLod.reduce((a, b) => a + b, 0), gpuVisible, lodMismatch,
perLod, indirectCounts, vertexCounts, lodPixels, wgslErrors,
};
const checks = [
{ name: 'wgsl.0error', ok: wgslErrors === 0, detail: `errors=${wgslErrors}` },
{ name: 'lod.match.cpu', ok: lodMismatch === 0, detail: `mismatch=${lodMismatch}` },
{ name: 'indirect.sum==visible', ok: indirectCounts[0]! + indirectCounts[1]! + indirectCounts[2]! === data.cpuVisible, detail: indirectCounts.join('+') },
{ name: 'indirect.vertexCounts', ok: vertexCounts[0] === 10 && vertexCounts[1] === 6 && vertexCounts[2] === 5, detail: vertexCounts.join(',') },
{ name: 'all.lods.drawn', ok: lodPixels.every((n) => n > 0), detail: lodPixels.join(',') },
];
return { checks, data };
}
node tools/run-l2.mjs src/render/culling/gpu-driven-cull-lod.ts（本机实测）：
PASS wgsl.0error errors=0
PASS lod.match.cpu mismatch=0
PASS indirect.sum==visible 8+204+62
PASS indirect.vertexCounts 10,6,5
PASS all.lods.drawn 9199,87713,4394
4096 个对象撒在 40³ 的立方体里，相机在原点：CPU 与 GPU 都判出 274 个可见（lodMismatch=0，逐对象一致）。三档 LOD 分别命中 8 / 204 / 62 个对象，indirect 参数里读回的 instanceCount 正好是 8+204+62=274，vertexCount 是 10 / 6 / 5；三档颜色的像素数分别为 9199 / 87713 / 4394，都非零，说明三档几何都真的画出来了。
为什么不能靠 firstInstance 选实例区
一个坑值得记下：最初想让三档共用一个实例缓冲、用 firstInstance 指定各档起点，结果在 SwiftShader 上非零偏移读出的是错位数据。改成每档一个独立实例缓冲、绑定偏移恒为 0，才稳定——这也是本节 inst0/inst1/inst2 三个绑定的由来。另一个坑：实例缓冲里 LOD 存的是 float，顶点属性若声明成 uint32，会把 1.0 的位模式读成 1065353216，必须声明 float32 再在着色器里 u32(lodF)。此外，indirect 缓冲漏了 COPY_SRC 会让整个命令缓冲被静默拒绝、回读全零，务必用 pushErrorScope('validation') 兜底。
小结
第 8 章到此收束：从 CPU 视锥剔除（8.1）、BVH 层次剔除（8.2），到把剔除与 LOD 全搬上 GPU（8.3）。三者共用同一套平面与包围体判定，只是执行位置不同。GPU 驱动后，CPU 每帧只发三个 drawIndirect——对象数再多，CPU 侧的提交开销也不涨。下一章处理另一个规模问题：把绘制调用按状态排序，进一步压低提交成本。
阅读导航：上一节：8.2 空间划分与 BVH · 下一节：9.1 排序键与前后绘制 。
继续阅读
探索更多技术文章
浏览归档，发现更多关于系统设计、工具链和工程实践的内容。