如果你把 GPU 只当作一块“画画的”硬件那你大概率还没体会到它一半的价值。这句话是我在把粒子系统从 CPU 循环改成 GPU Compute 之后最真实的感慨。Axmol 这次把 RHIRendering 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 决定。要支撑这个模型四件事是绕不开的独立的 compute pipeline不再绑定顶点布局、混合状态、深度状态这些渲染专属属性独立的 command encoder能够在 CommandBuffer 中插入计算 pass可写可读的存储缓冲区storage buffer让 shader 能直接修改数据明确的内存屏障和同步语义保证计算与渲染、计算与计算之间的先后关系不会乱。这套划分思路在 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的 bufferGLES 上需要生成一个 SSBOVulkan 上则是带 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 的完整矩阵搬进来而是提供了一层更加语义化的 barriercommandBuffer-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 不强行统一而是让上层引擎再封装一层。一个典型的取舍案例是读回readbackCPU 想读 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 版本为例平台支持情况大致如下平台图形 APICompute 支持主要限制与备注iOS / macOSMetal支持线程组内存上限 32KB需按 GPU family 判断具体能力AndroidGLES 3.1支持驱动质量参差中低端设备性能不稳定AndroidVulkan实验性支持能力最完整但引擎集成成熟度还在爬坡WebWebGL2不支持需要回退到 CPU 或旧的 fragment shader 方案WebWebGPU支持目前主要覆盖 Chrome 系仍在逐步完善这个矩阵对上层决策很有用。引擎层拿到isComputeSupported()之后最好还能拿到“支持强度”信息比如最高线程组大小、是否支持 storage texture。这样上层系统可以做多级回退最强的走完整 compute中等的只做局部计算最弱的直接 CPU。4. 实战在 Axmol 里跑一个十万粒子系统4.1 粒子数据的 Buffer 设计理论讲再多不如一个能跑的例子。我挑粒子系统来演示是因为它最能体现 compute 的价值数据量大、每帧都要更新、更新逻辑简单、渲染端用实例化就能画完。而且粒子系统的迭代状态天然存在特别适合展示“读改写同一块 buffer”的用法。粒子数据我用一个结构体表示位置和速度各三个 floatstruct ParticleData { float position[3]; float velocity[3]; };场景里预分配 10 万个粒子的 bufferusage 同时带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 Compute 更新耗时1 万1.8 ms0.3 ms10 万8.6 ms0.9 ms50 万40 ms2.8 ms数据来自我手上的测试工程机型不同数值会有差异但趋势是稳定的粒子数量越大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 设备上同步测否则等到集成后期再发现同步缺失改起来会牵扯到一大片资源管理代码。排查同步类问题的完整链路我建议按这个顺序走先确认 buffer 的 usage 是否同时包含读写需要的 flag再确认 compute pass 结束后是否显式或隐式插入了 barrier然后确认 ping-pong 双缓冲是否真的在交替写入而不是两个 pass 写同一块内存最后检查是不是线程组数量算错导致部分粒子根本没被更新。每步都给一段对照日志定位速度会快很多。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 升级的设计意图会发现每一步取舍都有它存在的道理。