企业官网建设流程全解析

如果你把 GPU 只当作一块“画画的”硬件，那你大概率还没体会到它一半的价值。这句话是我在把粒子系统从 CPU 循环改成 GPU Compute 之后最真实的感慨。Axmol 这次把 RHI（Rendering Hardware Interface）从纯图形渲染抽象升级成覆盖通用计算的能力层，本质上就是在补这个“另一半”。

Axmol 是从 Cocos2d-x 社区里长出来的 C++ 引擎，底子是 Cocos2d-x 4.0 那套多后端渲染架构，但这两年它做了不少“手术”。其中最关键的一刀，就是让 RHI 不再只关心“怎么把一个三角形画出来”，而是开始操心“怎么把一段数据交给 GPU 去算”。熟悉 GPU 编程的人都知道，compute shader 和 fragment shader 完全是两套心智模型：前者没有顶点、没有光栅化、没有 framebuffer，只有一群线程在读写 buffer。这篇文章我会从升级动机讲到接口设计，再到后端差异、实际集成和踩坑调优，适合想在 Axmol 里做 GPU 粒子、后处理、流体或任何需要通用计算的读者。

1. 老的 RHI 边界：画三角形的抽象为什么管不了计算

1.1 渲染抽象的基本盘

Cocos2d-x 4.0 时代引入的 backend 模块，是所有后续 RHI 工作的起点。它把渲染命令抽象成几个核心对象：CommandBuffer负责记录命令，RenderPipeline封装 shader 与混合状态，Buffer和Texture承载资源，RenderPass描述颜色和深度附件的目标。这样的设计在 2D 游戏里非常顺手：一个 Sprite 就是一个 draw call，一个 UI 界面就是一组按顺序提交的渲染命令。引擎层只需要组装参数，后端自动翻译成 OpenGL ES 或 Metal 调用。

但这套抽象有个隐含假设：每一个命令最终都会走到“光栅化”这一步。无论你是画三角形、画线段还是直接 blit 一张纹理，GPU 执行管线总是从顶点着色器开始，经过几何处理、光栅化、片元着色，最后写入 framebuffer。所有缓冲区的语义都是围绕这个流程设计的：顶点缓冲、索引缓冲、统一变量缓冲、纹理采样器。

在这样一个模型里，没有“GPU 自己算一段数据”的位置。你没法说“请 GPU 把这 10 万个粒子的位置加上速度乘以 delta time，然后写回同一块内存”。因为渲染管线的每一步都被定义了用途，而计算没有定义。

1.2 旧时代的“借壳渲染”方案

没有 compute 的时候，大家也不是不能做 GPU 计算。最经典的做法是把数据写在纹理里，用 fragment shader 做一轮运算，再渲染到另一张纹理上。这种做法在图形学里叫 render-to-texture 或 ping-pong 纹理，早期 GPGPU 全靠它。在 Axmol 升级前，如果你想在引擎里做一个 GPU 粒子系统，大体上就是这条路：

• 把粒子的位置、速度编码成 RGBA 纹理；
• 在 fragment shader 里采样上一帧纹理，执行物理更新；
• 把结果渲染到一张新纹理上；
• 采样新纹理，把粒子画成点或实例化网格。

这条路能跑，但实际工程里难受得很。首先是资源切换：每一次计算 pass 其实都是一个渲染 pass，你得准备 framebuffer、RenderPipeline、贴图附件，中间结果的保存和丢弃都要显式管理。其次是精度问题：把浮点数据塞进 8-bit 或 16-bit 纹理，要么精度不够，要么得自己搞定点数编码，debug 起来很痛苦。最后是语义别扭：数据一进 framebuffer，你就失去了对它“作为普通内存”的掌控，想读回、想原子操作，全都无从下手。

我在 GLES 设备上做后处理降采样的时候，为了把 4 个中间结果合并到一张纹理并且不触发额外的 framebuffer 转换，花了两天才把 preserve 和 discard 的状态理清楚。那还是渲染场景，不是通用计算。通用计算的复杂度只会更高。

1.3 升级的本质：执行模型从“绘制”变成“调度”

所以这次 RHI 升级，不能简单理解成“加几个 compute 函数”。它的本质是把 RHI 的执行模型从“绘制”扩展成“调度”。所谓调度，就是引擎告诉 GPU：这里有 N 个线程，每组 M 个，请执行这段 kernel，读写这些资源。至于结果是要画出来，还是要作为下一次计算的输入，都由资源本身的状态决定，而不是由 framebuffer 决定。

要支撑这个模型，四件事是绕不开的：

1. 独立的 compute pipeline，不再绑定顶点布局、混合状态、深度状态这些渲染专属属性；
2. 独立的 command encoder，能够在 CommandBuffer 中插入计算 pass；
3. 可写可读的存储缓冲区（storage buffer），让 shader 能直接修改数据；
4. 明确的内存屏障和同步语义，保证计算与渲染、计算与计算之间的先后关系不会乱。

这套划分思路在 WebGPU、Metal、Vulkan 里都有成熟先例，Axmol 真正要做的，是把它融进一个以 2D 游戏为主要使用场景的 C++ 引擎里。这就涉及很多具体的取舍，下面我会一点一点拆开讲。

2. 这次升级补了什么：资源、调度与同步三件套

2.1 资源形态：从 VertexBuffer 到 StorageBuffer

老的 RHI 里，Buffer的用途是固定的：要么是顶点数据，要么是索引数据，要么是 uniform。你创建 buffer 的时候就要告诉后端“我是干什么用的”。这个设计对驱动优化友好，但对 compute 不友好，因为 compute shader 需要的是既能读又能写的存储缓冲区。

以我拿到的简化接口为例，升级后的 buffer 创建会多出 storage 相关的 usage 标志：

auto* particleBuffer = rhi->createBuffer({ .size = maxParticles * sizeof(ParticleData), .usage = RHIBufferUsage::Storage | RHIBufferUsage::Vertex, .storageMode = RHIBufferStorageMode::DevicePrivate, });

这里有两个关键点。第一，usage增加了Storage选项，它指示后端这块 buffer 不仅可能被顶点阶段读取，还可能被 compute 阶段读写。第二，storageMode变成DevicePrivate，意味着数据默认只存在于 GPU 显存中，CPU 不会直接访问。这两个设计决定直接决定了后端的实现方式：Metal 上它会映射为MTLResourceStorageModePrivate的 buffer，GLES 上需要生成一个 SSBO，Vulkan 上则是带 storage buffer usage flag 的 VkBuffer。

之所以强调“读和写都要支持”，是因为 compute 的核心价值就是迭代。粒子系统每帧都要读上一帧的位置，算完再写回去。如果没有可写存储，就退回到纹理 ping-pong 的老路。这也是区分“真 compute 支持”和“花架子 compute 支持”的最直接标准。

2.2 调度抽象：dispatch 与线程组

后续的调度接口沿用了主流图形 API 的线程组概念。一次计算调度需要描述两个层面的信息：总共启动多少个线程组，以及每个线程组里有多少个线程。在 Metal 里是threadgroupsPerGrid和threadsPerThreadgroup，在 GLES/GLSL 里是gl_WorkGroupSize和glDispatchCompute的 group 数量，在 Direct3D 里则是numthreads和Dispatch的 group 数。

RHI 层的做法是把它们统一成一个 dispatch 调用：

commandBuffer->dispatch( groupX, groupY, groupZ, // 线程组数量 threadsX, threadsY, threadsZ // 每组线程数 );

实际数据规模通常远大于线程组数量，所以 compute shader 里一般通过gl_GlobalInvocationID或等价的全局线程 ID 来索引数据。RHI 不需要把“全局线程 ID”再抽象一遍，它只需要保证：shader 里拿到的全局 ID 能覆盖到整块数据，并且不同线程不会写到同一个位置上（除非你故意做原子操作）。

这里有个容易被忽视的设计细节：当数据量不是线程组整数倍时，最后一个 group 里的部分线程会超出数据范围。所以所有合格的 kernel 第一行几乎都是越界检查。这个责任在引擎层无法替你承担，RHI 只能提供一个约定：dispatch 的 group 数量由调用者自己向上取整。

2.3 同步语义：barrier 不能省也不能太粗暴

compute 引入后，最难处理的是同步。渲染时代，同一个 framebuffer 里的 draw call 天然按顺序执行，你不用担心前一个 draw 写深度、后一个 draw 读深度会乱掉——管线会帮你保证。但 compute 和渲染混在一起时，硬件不会自动帮你排顺序。

举个例子。你在 compute pass 里更新了一整块粒子 buffer，然后立刻在 render pass 里把它作为实例化顶点数据来画。如果中间没有屏障，某些 GPU 上顶点着色器可能读到的还是更新前的数据。因为 compute 的写入还没落盘，或者还在 L2 cache 里没刷出去。

RHI 的同步抽象必须暴露得很克制。它没有把 Vulkan 那套 dstStageMask、srcAccessMask 的完整矩阵搬进来，而是提供了一层更加语义化的 barrier：

commandBuffer->memoryBarrier( RHIBarrierScope::ComputeWrite, RHIBarrierScope::VertexRead );

内部的翻译因后端而异：在 GLES 上是glMemoryBarrier(GL_VERTEX_SHADER_BIT)，在 Metal 上是通过MTLFence或memoryBarrier接口建立依赖，在 Vulkan 上则展开为对应的vkCmdPipelineBarrier调用。抽象的目标不是覆盖全部场景，而是覆盖引擎内部最常用的“计算写、渲染读”“计算写、计算读”两种场景。

我个人的建议是，在 compute pass 结束、渲染 pass 开始之前，默认插入一次全屏障。因为 compute pass 和 render pass 之间的资源依赖关系引擎是能静态分析出来的，提前插屏障是安全的，而且对 2D 游戏的帧耗时影响基本可以忽略。真正需要精细优化的是多个 compute pass 之间，那种场景才值得做细粒度同步。

2.4 能力探测与回退策略

有了上面这些抽象，还不能直接铺开用。因为电脑上跑得好好的，不代表手机上也跑得起来。RHI 必须提供能力查询，否则用户代码会在老设备上直接崩溃。

升级后的 RHI 至少暴露了这样几个查询点：

if (rhi->isComputeSupported()) { // 使用新的 compute 路径 } else { // 回退到 CPU 更新或旧的 fragment 方案 }

除了“是否支持 compute”这种布尔查询，还应该有更细的能力枚举，比如“最大线程组大小”“是否支持 storage texture”“是否支持间接调度”。这些能力在不同平台差异很大，提前查清楚可以避免很多运行时 crash。

3. 后端适配的取舍：Metal、GLES 和 Web 端不是一家人

3.1 三个后端三种脾气

• Metal是对 compute 支持得最“原生”的底层 API。它专门设计了MTLComputeCommandEncoder，线程组内存（threadgroup memory）、同步、原子操作都是一等公民。iOS 从 A7 芯片开始就支持 Metal compute，所以在 iPhone 上做 GPU 粒子几乎是零门槛。

• OpenGL ES要分版本看。GLES 3.1 才引入 compute shader，并且配套的 SSBO 支持比较看驱动。Android 中低端设备的 GLES 驱动对 compute 的支持，往往是“能编译但不快”的状态。而且 GLES 的 compute 没有独立的 command encoder 概念，所有命令都在同一个 command buffer 里，同步全靠glMemoryBarrier，使用体验比 Metal 糙很多。

• Web 端是最麻烦的。WebGL2 到目前为止依然不支持 compute shader（严格说 WebGL2 规范里就没有 compute 入口），Emscripten 环境里只有 WebGPU 能提供真正的 compute pass。所以在 Web 上，升级后的 RHI 必须做好“没有 compute”的准备，继续走旧方案，或者干脆 CPU 计算。

3.2 统一抽象选了什么方案

面对三种不同脾气，RHI 的统一方案是“取最小公倍数，但保留能力出口”。计算管线对象、线程组调度、存储缓冲区、内存屏障这四件事是公因数，必须统一。而像 Metal 的 threadgroup memory、Vulkan 的 descriptor set、GLES 的 SSBO 绑定位置这类各自私有的能力，RHI 不强行统一，而是让上层引擎再封装一层。

一个典型的取舍案例是读回（readback）：CPU 想读 GPU 计算的结果，在 Metal 上是[MTLBuffer contents]，在 GLES 上是glMapBufferRange，在 Vulkan 上要绕一圈 staging buffer。RHI 最终只提供一个异步读回的接口。同步读取很容易阻塞管线，异步则是轮询或回调。宁愿接口难看一点，也要保证主流路径不会把帧率带崩。

另一个取舍是间接调度（indirect dispatch）。编程模型里一个很有意思的特性是：GPU 自己算出来需要跑多少个线程组，然后 GPU 自己发起下一次 dispatch。这个能力在 Metal 和 Vulkan 里都有，但 GLES 3.1 的间接调度支持参差不齐。RHI 早期版本干脆没有公开 indirect dispatch 接口，只支持 CPU 侧 dispatch。这样虽然损失了一部分性能优化空间，但保证了所有后端的语义一致性。等能力探测体系完善后，再逐步放开会稳得多。

3.3 平台能力矩阵

以我现在手上的 RHI 版本为例，平台支持情况大致如下：

这个矩阵对上层决策很有用。引擎层拿到isComputeSupported()之后，最好还能拿到“支持强度”信息，比如最高线程组大小、是否支持 storage texture。这样上层系统可以做多级回退：最强的走完整 compute，中等的只做局部计算，最弱的直接 CPU。

4. 实战：在 Axmol 里跑一个十万粒子系统

4.1 粒子数据的 Buffer 设计

理论讲再多，不如一个能跑的例子。我挑粒子系统来演示，是因为它最能体现 compute 的价值：数据量大、每帧都要更新、更新逻辑简单、渲染端用实例化就能画完。而且粒子系统的迭代状态天然存在，特别适合展示“读改写同一块 buffer”的用法。

粒子数据我用一个结构体表示，位置和速度各三个 float：

struct ParticleData { float position[3]; float velocity[3]; };

场景里预分配 10 万个粒子的 buffer，usage 同时带Storage和Vertex。前者给 compute shader 读写，后者给渲染管线当实例数据用。再加一个 CPU 侧的 staging buffer 用于初始化时写入初始位置和速度，初始化完成后 staging buffer 可以释放。

这里有个细节：Storage和Vertex同时标记，不是所有后端都默认支持，但 Metal 和现代 GLES 都可以。如果某个后端不支持同一块 buffer 同时用于两种用途，RHI 层会自动创建一份内部副本并在 barrier 时同步。这是抽象层的价值所在，引擎上层代码不用关心。

4.2 Compute Shader 的编写

粒子更新的 kernel 我用 GLSL 风格来写，RHI 内部会把它包装成各后端对应的 shader 入口。GLES 版本长这样：

#version 310 es layout(local_size_x = 256) in; layout(std430, binding = 0) buffer ParticleBuffer { ParticleData particles[]; }; uniform float u_dt; uniform vec3 u_gravity; uniform float u_damping; void main() { uint index = gl_GlobalInvocationID.x; if (index >= u_particleCount) { return; } ParticleData p = particles[index]; p.velocity += u_gravity * u_dt; p.velocity *= u_damping; p.position += p.velocity * u_dt; particles[index] = p; }

注意local_size_x = 256这个值的来历：它不是一个拍脑袋的数。NVIDIA 和 AMD 的 GPU 一个 wavefront/warp 通常是 32 或 64 线程，256 是 64 的整数倍，也基本是各平台 “一个线程组能塞下且不爆寄存器” 的安全值。移动端 GPU 的 wavefront 大小不等，但 256 对大多数 Mali 和 Adreno 设备也都是友好区间。后面调优部分我会再讲怎么根据具体 GPU 微调。

越界检查if (index >= u_particleCount)是必须的。因为 10 万不是 256 的整数倍，向上取整后最后一个线程组会多出一些线程，它们必须直接退出，否则就会越界写坏内存。这类 bug 在 CPU 上根本不会出现，但 GPU 并行世界里它是头号 crash 来源。

4.3 提交 Compute Pass 并衔接渲染

引擎层的调用流程是这样的：在CommandBuffer里先开始一个 compute pass，执行粒子更新，结束；再开始 render pass，把更新后的粒子 buffer 绑到顶点输入上，实例化绘制。

// 计算阶段 commandBuffer->beginComputePass(computeDescriptor); commandBuffer->setComputePipeline(particleComputePipeline); commandBuffer->bindBuffer(particleBuffer, 0); commandBuffer->bindUniform(u_dt, 1.0f / 60.0f); commandBuffer->bindUniform(u_gravity, vec3(0.0f, -9.8f, 0.0f)); commandBuffer->dispatch(ceil(maxParticles / 256.0f), 1, 1); commandBuffer->endComputePass(); // 渲染阶段 commandBuffer->beginRenderPass(renderPassDescriptor); commandBuffer->setRenderPipeline(particleRenderPipeline); commandBuffer->bindVertexBuffer(particleBuffer); commandBuffer->drawInstanced(maxParticles, 1); commandBuffer->endRenderPass();

注意计算 pass 和渲染 pass 之间的衔接：RHI 会在endComputePass时自动插入一个从计算写到顶点读的内存屏障。这是引擎层帮你做好的事情，但上层代码要意识到它的存在——如果你自己手动管理 dispatch 和 draw call 的顺序，而引擎没有自动插屏障，你就得手动调用 memory barrier 接口。

粒子的渲染我直接用了点精灵（point sprite）加实例化。10 万个点的逐顶点数据只有位置，每帧从 compute 更新后的 buffer 读取。这种组合在移动端非常常见：计算 pass 更新数据，渲染 pass 消费数据，整个管线里没有一次 CPU 回读，GPU 自产自销。

4.4 性能实测

在测试机上跑出来的数据挺能说明问题。CPU 方案是指主线程里用循环逐粒子更新，GPU 方案就是用上面的 compute 流程。两者都包含最终渲染。

数据来自我手上的测试工程，机型不同数值会有差异，但趋势是稳定的：粒子数量越大，CPU 方案损耗增长速度越猛，GPU 方案则基本稳住。50 万粒子时 CPU 方案已经掉到 20 帧以下，GPU 方案还能稳定在 60 帧。这就是 GPU Compute 的价值：它不是让快的东西更快，而是让 CPU 干不了的事情变得可能。

对比两张图的 GPU 时间占比会发现，compute pass 本身的耗时在 10 万粒子时只有不到 1ms。真正的瓶颈反而在渲染阶段——10 万个实例化的绘制开销、纹理采样和混合。这算是一个常见的认知偏差：很多人以为 compute 慢，其实在实例化渲染下，compute 往往不是瓶颈。

5. 踩坑记录与调优：从数据错乱到帧率崩盘

5.1 同步缺失导致的数据错乱

第一个坑就出在同步上。我最初把 compute pass 的 dispatch 和 render pass 的 draw 写在同一帧里，结果在 Metal 设备上一切正常，换到一台 GLES 安卓机上，粒子每隔几帧就会闪一下“回到初始位置”。一开始我以为是粒子逻辑写错了，排查了大半天。

后来我把问题拆开逐步验证。第一步检查 buffer 的 usage flag，确认是Storage | Vertex，没问题。第二步在 compute pass 和 render pass 之间手动调memoryBarrier，问题消失。这时候才反应过来：Metal 因为有 hazard tracking，会自动插入隐式依赖，所以看起来“不用管同步也能跑”；GLES 驱动不管这些，写后读全部要靠用户显式 barrier。

这个教训很典型：在功能验证阶段，千万别拿能力最强的后端（Metal）当作唯一参考实现。你在 Metal 上跑得顺，不代表抽象层真的做对了。从第一天起就要在 GLES 设备上同步测，否则等到集成后期再发现同步缺失，改起来会牵扯到一大片资源管理代码。

排查同步类问题的完整链路，我建议按这个顺序走：

1. 先确认 buffer 的 usage 是否同时包含读写需要的 flag；
2. 再确认 compute pass 结束后是否显式或隐式插入了 barrier；
3. 然后确认 ping-pong 双缓冲是否真的在交替写入，而不是两个 pass 写同一块内存；
4. 最后检查是不是线程组数量算错，导致部分粒子根本没被更新。

每步都给一段对照日志，定位速度会快很多。

5.2 分支过多导致的并行效率崩塌

第二个坑是性能层面的。把粒子系统加上“按区域随机重置”的逻辑后，我在 update 里加了这样一个分支：

if (p.position.y < -100.0f) { resetParticle(p); } else { updatePhysics(p); }

看起来人畜无害，但实测帧率直接掉了四成。原因很简单：一个线程组内有 256 个线程，GPU 以 wavefront 为单位执行，同一个 wavefront 内所有线程必须执行相同的指令。当分支出现分歧时，GPU 会把两条路径都执行一遍，然后各自掩蔽。粒子系统里粒子的状态高度不均，一个 warp 里可能一半粒子要重置、一半要更新，等于两条指令流都跑满，计算量翻倍。

解决方式不是去掉分支，而是改成“无分歧”的写法：用算术方式做选择，或者把重置逻辑从主循环里拆出去，单独用一个 dispatch 处理边界粒子。我这里用了更简单的做法：把重置条件里的分支去掉，改为用mix和step组合出选择结果。代价是有些粒子即便不需要重置也会多算几下，但因为不需要执行第二条指令流，整体反而快得多。

5.3 CPU 回读的代价

第三个坑来自调试习惯。为了验证粒子内存里到底写了什么，我一开始每帧把 particle buffer 回读到 CPU，用printf打印几个粒子的位置。结果帧率从 60 直接掉到 15。

在移动端，GPU buffer 回读是非常昂贵的操作。哪怕只有 10 万个粒子，同步回读也可能让 GPU 管线直接停顿等待 CPU。之后我换成了两个方案：一是只在特定调试帧里回读，用按键触发而不是每帧触发；二是利用 RHI 的异步读回接口，把结果放进一个队列，等 GPU 完成那一帧的工作后再在后台取回。实测显示，异步回读对帧率的影响几乎可以忽略。

如果你真的需要每帧回读小数据量，也尽量不要回读整个粒子 buffer，而是单独维护一个“统计 buffer”，只让 GPU 把归约结果写进去，再回读那几百个字节。颗粒度越小，开销越小。

5.4 线程组大小与寄存器占用

最后一个调优点，是线程组大小（local_size）的选择。我用同一份粒子 kernel，在不同设备上调了 64、128、256、512 四档，结果很有意思：

• 64：部分移动 GPU 上利用率偏低，因为 wavefront 数量太少；
• 256：绝大多数设备表现稳定，兼容性最好；
• 512：在个别 Adreno 设备上有提升，但在 Mali 上寄存器溢出导致性能倒退。

寄存器溢出是个值得展开说的坑。每个线程可用的寄存器数量有限，当 kernel 里的临时变量太多、线程组又很大时，寄存器放不下，数据会 spill 到本地内存。本地内存虽然在线程组内共享，但访问速度远慢于寄存器，性能直接腰斩。Mobile GPU 上这个问题更明显。

经验做法是：先用 256 跑通，然后用图形调试工具看寄存器占用和 spill 情况。如果 spill 严重，就调低线程组大小，或者精简 kernel 的临时变量。线程组大小和寄存器占用是一对矛盾：线程越多，每线程分到的寄存器越少。找到平衡点的方法只有一个——在你的目标设备矩阵上挨个测。

调试线程组问题时，我推荐直接看 RHI 暴露的maxThreadsPerGroup能力值。把它和实测数据放一起对比，能很快判断某个设备是硬件限制还是驱动 bug。

这次 RHI 从图形渲染扩展到 GPU Compute，给我最大的感受是：引擎层真正开始把 GPU 当作一个“可编程的数据处理器”来对待，而不仅仅是一个画图设备。对于想接 GPU 粒子和后处理的开发者来说，建议第一步不要直接上流体模拟那种自带迭代复杂度的项目，而是先拿一个后处理特效（比如高斯模糊或颜色分级）练手。后处理的输入输出都是纹理，链条短，出了问题最好排查；跑通了再逐步往粒子、物理这类有状态迭代的场景扩展。等你把资源生命周期、同步屏障、后端差异这一整套手感磨出来，再回头看这次 RHI 升级的设计意图，会发现每一步取舍都有它存在的道理。