1. 为什么要在 WebGPU 上折腾 Meshlet Culling第一次看到WebGPU 上做简易版 Nanite这个想法时我脑子里冒出来的第一个念头是浏览器里跑虚拟几何体这不是自找麻烦吗但真把 demo 跑起来之后我发现这件事比想象中靠谱得多。WebGPU 提供的 compute shader 和 storage buffer 能力已经足够支撑一套完整的 meshlet 剔除管线虽然离 UE5 那套 Nanite 的完整实现还有距离但核心思路——把大网格切成小簇、按需剔除、只画可见部分——完全可以在浏览器里复现出来。这篇文章想聊的就是这么一件事用 WebGPU 实现一套简易版的 meshlet culling 系统目标不是复刻 Nanite 的全部特性而是把虚拟几何体最核心的那条链路跑通。适合谁看如果你已经写过 WebGPU 的基础渲染管线对 compute shader 有基本概念又想搞清楚 Nanite 那套东西到底是怎么运作的那这篇内容应该能帮到你。如果你连 WebGPU 的 render pass 都还没写过建议先去补一下基础不然中间很多细节会卡住。先说清楚这套东西能做什么给定一个高面数模型比如几十万面的雕塑在浏览器里以稳定帧率渲染并且随着相机移动动态剔除掉视锥外和背面朝向的 meshlet。它解决的核心问题是draw call 爆炸和顶点处理浪费——传统方式下你要么一次性提交所有三角形让 GPU 硬扛要么在 CPU 端做粗粒度剔除前者浪费顶点着色器算力后者受限于 CPU 带宽。Meshlet 方案把剔除粒度降到簇级别并且把剔除逻辑搬到 GPU 上并行执行这才是它真正的价值所在。2. Meshlet 与 Nanite 的核心思路拆解2.1 Meshlet 到底是什么为什么它能省性能Meshlet 说白了就是把一个大网格切成很多个小块每块包含固定数量的顶点和三角形常见配置是 64 个顶点、124 个三角形。这个数字不是随便定的它跟 GPU 的 wave/warp 大小、缓存行、以及后续要用的 cluster 剔除算法都有关系。为什么是 64 和 124因为 64 个顶点刚好能塞进一个 64 线程的 compute workgroup124 个三角形则是在 8-bit 索引最多 256 个索引约束下能塞进一个 meshlet 的最大三角形数——每个三角形 3 个索引124 × 3 372超过 256 了所以实际会用 8-bit 局部索引加一个顶点重映射表来压缩。这里有个关键点很多人一开始会忽略meshlet 的顶点是局部索引的。也就是说每个 meshlet 内部用自己的 0~63 编号引用顶点而不是直接用全局顶点索引。这样做的好处是索引可以用 8 位存储节省带宽坏处是每个 meshlet 都要维护一张局部索引 → 全局索引的映射表。这张表在剔除阶段用不到但在最后光栅化阶段必须还原成全局索引否则顶点着色器不知道去哪个 buffer 取数据。Nanite 相比普通 meshlet 的进阶之处在于它做了层级化cluster hierarchy和GPU 驱动的 LOD 选择。简单说就是把 meshlet 按空间关系组织成一棵树父节点是子节点的简化版本运行时根据屏幕空间误差决定渲染到哪一层。这套东西实现起来复杂度很高我们这篇只做简易版所以层级化 LOD 先放一放专注把单层 meshlet 的剔除链路跑通。2.2 为什么剔除要放到 GPU 上做传统 CPU 剔除的问题在于粒度太粗。你可以在 CPU 端对每个物体做视锥剔除但一个物体可能有几万个三角形剔除掉整个物体意味着只要它有一个角在视锥内你就得把全部三角形提交给 GPU。Meshlet 把粒度降到 124 个三角形剔除精度提升了几百倍但代价是剔除次数也提升了几百倍——这时候如果还在 CPU 上做光是遍历几万个 meshlet 的包围球就够 CPU 喝一壶了。GPU 剔除的优势在于并行度。几万个 meshlet 的包围球测试在 compute shader 里可以一次性并行处理每个线程负责一个 meshlet测试视锥、背面、遮挡然后把通过的 meshlet 索引写到一个紧凑的 buffer 里。后续的 draw call 只需要读取这个 buffer 的计数用 indirect draw 一次性提交所有可见 meshlet。整个过程 CPU 只负责提交几个 dispatch 和 draw 命令几乎不参与实际计算。这里有个实操细节indirect draw 的计数器必须放在 GPU 可写的 buffer 里而且要注意 buffer 的 usage 要包含INDIRECT和STORAGE。WebGPU 里这个限制比较严格如果你忘了加INDIRECT运行时会直接报 validation error而且错误信息不一定直观我第一次踩这个坑的时候查了半天。2.3 简易版和完整版 Nanite 的差距在哪得先把预期摆正。完整版 Nanite 包含的东西远超 meshlet culling它有 cluster 层级的自动构建、有基于屏幕空间误差的 LOD 选择、有软件光栅化器处理微小三角形、有流式加载支持超大模型、还有和材质系统的深度集成。我们这套简易版只覆盖其中一小部分特性完整 Nanite本文简易版Meshlet 切分自动层级化单层离线切分剔除粒度Cluster 级Cluster 级LOD 选择GPU 驱动多层级无单层光栅化硬件 软件混合纯硬件遮挡剔除两阶段 HZB可选简化版流式加载支持不支持这个对比不是要贬低简易版而是让你清楚知道哪些地方可以简化、哪些地方简化了会出问题。比如 LOD 这块简易版不做层级化意味着你的模型面数不能太夸张否则单层 meshlet 数量会爆炸。我实测下来单层方案在 50 万面以内还能跑得比较舒服再往上就得考虑层级化了。3. 核心细节解析与实操要点3.1 Meshlet 切分离线工具怎么写切分这一步是在 CPU 端离线做的不在运行时。你需要写一个小工具Python 或 C 都行读入模型文件输出三个 buffermeshlet 顶点数据、meshlet 三角形索引、meshlet 包围球信息。切分算法本身不复杂但有几个坑。最朴素的做法是按三角形顺序每 124 个切一刀但这样切出来的 meshlet 空间分布可能很散包围球会很大剔除效果差。更好的做法是基于空间局部性做聚类比如用 k-means 或者简单的贪心算法把空间上靠近的三角形分到同一个 meshlet 里。我自己的做法是先用网格哈希把三角形按空间位置分桶然后每个桶内再按 124 个一组切分。这样切出来的 meshlet 包围球半径通常能比顺序切分小 30% 到 50%剔除效率提升很明显。包围球的计算用经典的 Ritter 算法或者更简单的 AABB 中心加最大距离就行精度要求不高。顶点去重也要注意。同一个顶点可能被多个 meshlet 引用如果每个 meshlet 都存一份完整顶点数据显存会浪费。常见做法是全局顶点 buffer 只存一份每个 meshlet 存局部索引到全局索引的映射。但这样在剔除阶段用不到映射只有最后光栅化才需要所以映射表可以单独放一个 buffer减少剔除阶段的数据读取量。3.2 包围球测试的数学细节视锥剔除的核心是判断包围球是否和视锥的六个平面相交。标准做法是把包围球中心代入每个平面方程如果六个平面的有符号距离都小于负半径说明球完全在视锥外剔除。这里有个容易搞错的地方平面方程的归一化。如果你的平面法线没有归一化那么距离计算出来的值就不是真实距离和半径比较就会出错。我见过不少人在这里翻车表现是剔除结果时对时错很难 debug。建议在构建视锥平面的时候就直接归一化后面用起来省心。背面剔除更简单判断包围球中心到相机的向量和 meshlet 平均法线的点积如果大于某个阈值就认为背面朝向。但这里有个细节单个 meshlet 的平均法线可能不准特别是形状不规则的 meshlet。更稳妥的做法是用 meshlet 的圆锥体cone表示法线范围判断视锥和圆锥是否相交。不过简易版用平均法线加一个宽容阈值也够用实测下来误剔除率可以控制在可接受范围。3.3 Compute Shader 的线程组织剔除 compute shader 的线程组织直接影响性能。最直接的做法是一个线程处理一个 meshletworkgroup size 设成 64 或 128。但如果 meshlet 数量是几万dispatch 的 workgroup 数量就是几百这个量级对 GPU 来说很轻松。不过有个优化点把视锥平面和相机参数放到 uniform buffer 里而不是 storage buffer。uniform buffer 的访问速度更快而且这些数据是所有线程共享的放 uniform 更合适。meshlet 的包围球数据放 storage buffer因为它是只读的、每个线程读不同的数据。输出方面用一个 atomic counter 来记录可见 meshlet 的数量可见 meshlet 的索引写到一个预分配的 storage buffer 里。这里要注意 atomic 操作的性能如果 meshlet 数量特别大atomic 竞争可能成为瓶颈。一个优化技巧是每个 workgroup 先做局部聚合再用一个 atomic 更新全局计数这样能把 atomic 操作次数降低一个数量级。// 简化的剔除 shader 核心逻辑 group(0) binding(0) varuniform camera: CameraUniforms; group(0) binding(1) varstorage, read meshlets: arrayMeshletData; group(0) binding(2) varstorage, read_write visibleIndices: arrayu32; group(0) binding(3) varstorage, read_write counter: atomicu32; compute workgroup_size(64) fn cullMain(builtin(global_invocation_id) gid: vec3u32) { let idx gid.x; if (idx arrayLength(meshlets)) { return; } let m meshlets[idx]; // 视锥测试 var inside true; for (var i 0u; i 6u; i i 1u) { let dist dot(camera.frustumPlanes[i].xyz, m.center) camera.frustumPlanes[i].w; if (dist -m.radius) { inside false; break; } } if (!inside) { return; } // 背面测试 let toCamera normalize(camera.position - m.center); if (dot(toCamera, m.normal) -0.2) { return; } let slot atomicAdd(counter, 1u); visibleIndices[slot] idx; }这段代码是简化版实际用的时候还要处理 workgroup 内聚合、边界情况等。但核心逻辑就是这么直白测试、通过就写索引、计数加一。3.4 Indirect Draw 的参数组织剔除完成后你需要用 indirect draw 提交渲染。WebGPU 的 indirect draw 参数是一个固定格式的 buffer包含 vertexCount、instanceCount、firstVertex、firstInstance 四个 u32。对于 meshlet 渲染通常用 instanceCount 来表示可见 meshlet 数量每个 instance 对应一个 meshlet。这里有个关键点indirect buffer 的内容必须由 GPU 写入因为可见 meshlet 数量是运行时才知道的。你需要一个小的 compute shader 或者用copyBufferToBuffer把 counter 的值写到 indirect buffer 的对应位置。WebGPU 目前不支持直接在 shader 里写 indirect buffer 的特定字段所以通常的做法是先用一个 compute pass 把 counter 拷贝过去再发起 draw。另一个细节是firstInstance 的用法。如果你想让每个 meshlet 实例通过instance_index去查可见索引表那 firstInstance 设 0 就行shader 里用instance_index直接索引 visibleIndices。但要注意 WebGPU 的instance_index是从 0 开始的和某些图形 API 的行为可能不同跨平台时要注意。4. 完整实操流程与关键环节实现4.1 从模型到 Meshlet Buffer 的完整管线整个离线管线分四步加载模型、切分 meshlet、生成包围球、打包 buffer。我用 Python 写了一个脚本依赖 numpy 和 trimesh处理一个 30 万面的模型大概需要十几秒完全可以接受。第一步加载模型没什么好说的trimesh 直接读 obj 或 glb 都行。第二步切分是重点我用的空间哈希方案先算模型的 AABB按 32×32×32 的网格分桶每个三角形根据重心位置落到对应桶里。然后遍历每个桶如果桶内三角形数超过 124就继续细分或者直接按顺序切。这样切出来的 meshlet 空间局部性很好。第三步包围球我用的是简单版取 meshlet 所有顶点的 AABB 中心作为球心然后算所有顶点到球心的最大距离作为半径。这个方案比 Ritter 算法稍差但实现简单实测半径只大 5% 左右对剔除效果影响很小。第四步打包顶点数据用 f32 存位置和法线索引用 u8 存局部索引映射表用 u16 存全局索引。包围球用 vec4 存xyz 是球心w 是半径法线用 vec3 存。这些数据分别打包成独立的 buffer运行时直接上传。4.2 WebGPU 资源创建与绑定组布局资源创建这块有几个容易踩的坑。首先是 buffer 的 usage剔除阶段要读 meshlet 数据所以 meshlet buffer 要有STORAGE可见索引 buffer 要写所以要有STORAGEindirect buffer 要作为 draw 的参数所以要有INDIRECT。少一个都会在运行时炸掉。绑定组布局要提前规划好。我习惯把剔除相关的绑定放在 group 0渲染相关的放在 group 1。剔除的 group 0 包含camera uniform、meshlet storage、visible indices storage、counter storage。渲染的 group 1 包含顶点 buffer、索引 buffer、映射表 buffer、visible indices storage渲染时也要读。这里有个细节同一个 buffer 可以在不同的绑定组里以不同的 access 模式绑定。比如 visible indices 在剔除阶段是read_write在渲染阶段是read。WebGPU 允许这样做但要注意绑定组布局里的 access 模式要和 shader 里的声明一致否则会报错。4.3 剔除 Pass 与渲染 Pass 的衔接剔除 pass 和渲染 pass 之间的衔接是整套流程里最容易出问题的地方。核心问题是同步渲染 pass 必须等剔除 pass 完全写完 visible indices 和 counter 之后才能开始。WebGPU 的 command encoder 会自动处理同一队列内的依赖但如果你用了多个队列或者跨帧复用 buffer就要手动加 barrier。我的做法是每帧重新创建 command encoder剔除 pass 和渲染 pass 在同一个 encoder 里顺序编码这样 WebGPU 会自动插入必要的同步。counter buffer 每帧开始时要清零可以用clearBuffer命令或者用一个小的 compute shader 写 0。我倾向用clearBuffer更简单直接。渲染 pass 里顶点着色器通过instance_index去 visible indices 表里查当前 meshlet 的索引然后用这个索引去 meshlet buffer 里取顶点数据。这里要注意instance_index的范围如果 indirect draw 的 instanceCount 是 0整个 draw 会被跳过不会执行顶点着色器所以不用担心越界。4.4 性能实测与调优记录我在一台配备集成显卡的笔记本上做了实测模型是一个 28 万面的雕塑切分成约 2300 个 meshlet。不剔除的情况下帧率在 45 帧左右开启视锥剔除后帧率稳定在 60 帧垂直同步上限开启视锥加背面剔除后帧率没有明显变化但 GPU 占用率从 85% 降到了 60% 左右。这个结果说明剔除确实有效但瓶颈可能不在剔除本身。我用 timestamp query 测了一下各阶段耗时剔除 pass 大约 0.3ms渲染 pass 大约 12ms。剔除的开销几乎可以忽略主要时间还是花在光栅化上。这也符合预期毕竟简易版没有做 LOD所有可见 meshlet 都是全精度渲染。调优方面我试过把 workgroup size 从 64 改成 128剔除耗时没有明显变化。试过把 meshlet 数据从 storage buffer 改成 uniform buffer结果因为 uniform buffer 大小限制通常 64KB放不下所有 meshlet 而失败。最后发现最有效的优化是减少 meshlet 数量——把切分粒度从 124 个三角形改成 248 个用 u16 索引meshlet 数量减半剔除耗时降低约 40%但剔除精度下降渲染耗时略有上升。这个权衡要根据具体场景来定。5. 常见问题与排查技巧实录5.1 剔除结果不对的排查思路剔除结果不对是最常见的问题表现可能是该显示的没显示、不该显示的显示了、或者画面闪烁。排查的时候按这个顺序来先检查视锥平面。把六个平面的方程打印出来手动代入几个已知点的坐标看看结果是否符合预期。我遇到过一次平面法线没归一化的问题表现是远处物体被错误剔除近处正常查了半天才发现是归一化漏了。再检查包围球。把 meshlet 的包围球可视化出来可以用一个简单的 debug 渲染 pass 画球看看球的位置和大小是否合理。如果球明显偏大或者偏小说明切分或者包围球计算有问题。最后检查 shader 里的测试逻辑。特别是背面测试的阈值如果阈值设得太激进会导致侧面朝向的 meshlet 被误剔除表现是转动相机时物体边缘闪烁。我一般把阈值设在 -0.2 到 -0.3 之间比较稳妥。5.2 Indirect Draw 不生效的几种情况Indirect draw 不生效画面全黑或者只显示一部分可能的原因有counter 没有正确清零。如果 counter 保留了上一帧的值indirect draw 的 instanceCount 会偏大导致渲染越界或者渲染了不该渲染的 meshlet。检查方法是在剔除 pass 之后读回 counter 的值看看是否符合预期。indirect buffer 的 usage 不对。WebGPU 要求 indirect buffer 必须有INDIRECTusage而且不能同时有MAP_READ。如果你需要读回 indirect buffer 的内容做调试要单独创建一个可映射的 buffer用copyBufferToBuffer拷贝过去再映射。visible indices 的写入顺序有问题。如果多个线程同时写同一个 slot会导致索引覆盖。用 atomicAdd 获取 slot 可以避免这个问题但要确保 atomic 操作的正确性。我见过有人用普通的加法而不是 atomicAdd结果在高并发下出现索引丢失。5.3 性能不如预期的优化方向如果剔除开了但性能没提升可能是这几个原因剔除粒度不够细。如果你的 meshlet 切得太大比如 500 个三角形一个剔除精度就低很多不该渲染的 meshlet 还是会被提交。建议控制在 124 到 256 个三角形之间。瓶颈不在顶点处理。如果场景的片元着色很重比如复杂材质、大量 overdraw剔除顶点带来的收益会被片元处理掩盖。这时候要考虑优化材质或者做 early-z。CPU 端开销太大。如果每帧都在 CPU 端重建绑定组或者重新上传 bufferCPU 开销可能成为瓶颈。建议把不变的资源提前创建好每帧只更新必要的 uniform。问题现象可能原因排查方法画面全黑counter 未清零读回 counter 值检查物体边缘闪烁背面剔除阈值过激放宽阈值到 -0.3远处物体消失视锥平面未归一化检查平面方程帧率无提升瓶颈在片元处理用 timestamp query 测各阶段索引错乱非原子写入改用 atomicAdd5.4 跨平台兼容性注意事项WebGPU 虽然标准统一但不同浏览器的实现还是有差异。我实测下来Chrome 和 Edge 的表现基本一致Firefox 在某些情况下对 storage buffer 的大小限制更严格。如果你的 meshlet 数据超过 128MB在 Firefox 上可能会创建失败需要做分块加载。另一个坑是workgroup size 的限制。WebGPU 规范要求 workgroup size 的乘积不能超过 256而且某些设备对单个维度的限制更小。我建议用 64 或 128 这种保守的值兼容性最好。还有一点是shader 的精度问题。WebGPU 的 f32 在大多数设备上是真 32 位但某些移动设备可能会降级。如果你的剔除测试对精度敏感建议在 shader 里显式声明精度或者用更宽容的阈值。6. 后续可以继续折腾的方向这套简易版跑通之后能扩展的方向其实不少。最直接的是加遮挡剔除用上一帧的深度图构建 HZB在 compute shader 里做层级深度测试把被遮挡的 meshlet 也剔除掉。这一步做完室内场景的剔除效率能再提升一大截。再往上是层级化 LOD把 meshlet 组织成树运行时根据屏幕空间误差选择渲染层级。这块复杂度高很多但也是 Nanite 真正的精髓所在。如果只是想学习可以先从两层 LOD 开始手动构建简化版本跑通之后再考虑自动简化。还有一个方向是流式加载把 meshlet 数据分块根据相机位置动态加载和卸载。这个对超大场景很有用但需要配合资源管理和异步加载工程量不小。我自己在实际操作中的体会是meshlet culling 这套东西看起来复杂但拆开之后每一步都不难难的是把整条链路串起来并且调对。最容易出问题的地方往往不是算法本身而是资源绑定、同步、边界条件这些脏活。建议先把最简单的单 meshlet 渲染跑通再逐步加剔除、加 indirect draw每加一步都验证一下比一次性写完再 debug 要高效得多。