Chapter 25: Compute Shaders
Compute Shader 把 GPU 上可编程执行单元从固定图形阶段中抽出来,让开发者直接提交一组并行线程去读写 buffer、texture 和共享内存。本章要建立的能力是:看到一个图形任务时,能判断它是否适合放进 Compute Shader,能把 dispatch、work group、资源读写和渲染 pass 之间的同步关系追踪清楚,并能根据可观察证据定位错误或性能瓶颈。
本章使用一个贯穿案例:一帧粒子渲染先由 compute pass 更新 100000 个粒子的位置和速度,再由 render pass 把同一个 particle buffer 当作顶点或 instance 数据绘制出来。这个案例足够小,可以用几十行 shader 表达;它也覆盖 Compute Shader 最常见的工程问题:线程如何映射到数据、group size 如何选择、shared memory 何时有收益、写完 buffer 后何时能被图形管线读取。
官方资料对这个模型的描述在不同 API 中名称不同,但核心结构稳定。Direct3D 用 Dispatch 和 HLSL numthreads 描述 group 数量与 group 内线程数量,Microsoft Learn 的 Compute Shader Overview 把它归为可并行执行的通用计算阶段;WGSL 使用 @compute 和 @workgroup_size 标记 compute entry point,W3C WGSL Candidate Recommendation Draft, 2026-06-05 明确给出 global_invocation_id 与 workgroup 网格的关系;Vulkan 把 compute pipeline、shader stage 和同步放入显式命令模型,Vulkan 1.4.353 Specification, 2026-06-04 需要开发者自己表达跨 pass 的资源依赖。
读完本章后,Compute Shader 应该被理解成一条“可调度的并行数据处理路径”。它的输入是 buffer、texture、uniform、push constant 或 descriptor;处理过程是大量 invocation 按 work group 组织执行;输出是写回 GPU 资源;影响会在后续 render pass、copy pass、readback 或下一次 compute dispatch 中表现出来。
贯穿案例的数据路径如下图。图中只展示本章讨论的最小闭环:CPU 提交参数,compute pass 更新粒子,barrier 或 resource transition 建立可见性,render pass 消费结果。
这条路径给本章提供了统一判断标准:每个问题都回到“哪段资源被哪个 pass 读写、哪个 dispatch 负责多少数据、哪些线程需要协作、结果何时对后续管线可见”。
25.1 Compute Shader 的定位与用途边界
Compute Shader 的工作定义是:由图形 API 调度、在 GPU shader 执行单元上运行、以 thread grid 为执行空间、主要通过 storage buffer、storage texture、UAV 或等价资源读写数据的可编程阶段。它和 vertex shader、fragment shader 使用相近的 GPU 执行硬件,但执行触发方式由 dispatch 决定,输出也通常写入显式资源。
在粒子案例中,Compute Shader 负责的任务是“把上一帧粒子状态更新为当前帧粒子状态”。输入资源包括 particle buffer、时间步长、重力、噪声参数和边界范围。输出资源仍然是 particle buffer。render pass 只读取这个结果,并把每个粒子变成屏幕上的点、sprite 或 mesh instance。这个拆分让计算路径和绘制路径形成清楚的读写关系。
Compute Shader 适合承担数据并行程度高、每个元素处理规则相近、输入输出可以放进 GPU 资源的任务。典型例子包括粒子更新、GPU culling、prefix sum、buffer compaction、image filter、tile-based light list、simulation step、procedural data 生成。它们共同的结构是:一批元素可以映射为一批 invocation,并且大部分工作发生在 GPU 内存中。
Compute Shader 适合度低的任务也有稳定特征:任务依赖复杂的 CPU 对象图、每一步都需要和系统 I/O 交互、分支路径高度不均匀、结果马上需要 CPU 逐项读取、输入规模太小导致 dispatch 和同步开销占主导。一个只更新 12 个 UI 控件位置的小任务放到 GPU 上,CPU 提交、资源转换和同步开销通常会超过计算本身。一个包含大量不规则字符串解析的任务也很难映射到统一的 GPU 线程网格。
判断 Compute Shader 边界时,先看任务形态,再看资源路径。任务形态回答“是否存在大量同构元素”;资源路径回答“数据是否已经在 GPU 上,输出是否继续被 GPU 使用”。粒子更新同时满足这两个条件:100000 个粒子共享同一套更新公式,更新后的结果直接进入 render pass。CPU 逐粒子更新会消耗大量主机计算和上传带宽;compute pass 让数据停留在 GPU buffer 中,减少跨 CPU/GPU 边界的移动。
下列简化 WGSL 片段展示了这个边界。代码的目的只在于表达执行粒度和资源读写关系,省略随机数、碰撞和生命周期管理。
struct SimParams {
dt: f32,
gravity: vec2<f32>,
count: u32,
};
struct Particle {
position: vec2<f32>,
velocity: vec2<f32>,
};
@group(0) @binding(0)
var<storage, read_write> particles: array<Particle>;
@group(0) @binding(1)
var<uniform> params: SimParams;
@compute @workgroup_size(256)
fn update_particles(@builtin(global_invocation_id) gid: vec3<u32>) {
let index = gid.x;
if (index >= params.count) {
return;
}
var p = particles[index];
p.velocity = p.velocity + params.gravity * params.dt;
p.position = p.position + p.velocity * params.dt;
particles[index] = p;
}
这段代码对应的工程事实很直接:global_invocation_id.x 决定当前 invocation 处理哪个粒子,@workgroup_size(256) 决定每个 work group 内最多有 256 个 invocation,params.count 是越界保护,particles[index] 是本次 compute pass 的读写资源。读者在 frame capture 中复查时,应先找 compute pipeline 绑定的 storage buffer,再看 dispatch 数量是否覆盖 count,最后看 render pass 是否读取同一份 buffer。
Compute Shader 的定位由“输入资源 → 并行计算 → 输出资源 → 后续消费者”闭合。只要这个闭合关系说不清,Compute Shader 就会变成一个模糊的加速标签;一旦闭合关系明确,后续 group size、同步和调试都可以沿同一条路径展开。
25.2 并行计算模型与线程分组(Work Group)设计
Compute Shader 的执行空间由两层组成:dispatch 提交 work group 的数量,shader 声明每个 work group 内的 invocation 数量。总 invocation 数等于 dispatch_group_count_x * workgroup_size_x,二维或三维任务再把 y、z 维度也乘进去。Direct3D 的 numthreads(X, Y, Z) 和 WGSL 的 @workgroup_size(X, Y, Z)都在表达 group 内尺寸;API 的 Dispatch、dispatchWorkgroups 或 vkCmdDispatch 表达 group 网格尺寸。
粒子更新是一维数据任务,最自然的映射是一个 invocation 处理一个粒子。假设粒子数量为 100000,work group size 选 256,则 dispatch group count 应为向上取整后的 391,因为 391 * 256 = 100096。多出来的 96 个 invocation 通过 index >= count 返回。这种越界保护是 GPU compute 常规写法,因为 dispatch 网格通常按整组提交。
uint32_t groupSize = 256;
uint32_t particleCount = 100000;
uint32_t groupCount = (particleCount + groupSize - 1) / groupSize;
Dispatch(groupCount, 1, 1);
这段伪代码对应的判断点是:CPU 提交的是 group 数量,shader 声明的是 group 内线程数量,shader 内部拿到的是每个 invocation 的 id。很多 dispatch 错误来自把这三者混成一个数字。若把 Dispatch(100000, 1, 1) 和 @workgroup_size(256) 搭配使用,实际 invocation 数会变成 25600000,shader 内的越界保护会掩盖错误,但会造成大量空转线程。
work group size 的选择要同时考虑四类约束。第一类是 API 或设备 limit,例如 Direct3D cs_5_0 常见上限为每组 1024 个线程,WebGPU 也通过设备 limits 暴露最大 workgroup 尺寸。第二类是硬件 wave 或 subgroup 粒度,选择 64、128、256 这类倍数通常能减少空 lane。第三类是寄存器、shared memory 和 occupancy 的关系,每个线程资源占用越高,同一计算单元上同时驻留的 group 数量越少。第四类是内存访问模式,连续线程访问连续元素时更容易形成合并内存访问。
粒子更新的主路径只读写自己的 particles[index],线程之间没有数据交换,shared memory 收益很低。此时 group size 主要服务调度效率和内存访问连续性。256 是常见起点,因为它能覆盖多种硬件的 wave 粒度,又不会让一个 group 过大到降低 occupancy。最终选择仍需以 profiler 的 GPU time、occupancy、memory transaction 和 stall 指标为准。
shared memory 适合线程组内部重复使用同一批数据的任务。图像模糊是更典型的例子:一个 group 处理一块 tile,先把 tile 和边缘 halo 读入 shared memory,再让多个线程重复读取邻域像素。粒子更新如果加入局部邻域碰撞,也可能使用 shared memory 缓存空间分区中的粒子;普通重力积分则没有这个必要。
下面的 Mermaid 图把 group、invocation 和 particle index 的关系压缩成一个可复查模型。
这个模型的调试顺序也很固定。先确认 groupCount 是否由元素数量和 group size 向上取整得到;再确认 shader 使用 global_invocation_id 生成线性 index;然后确认每个写入 index 落在 buffer 有效范围内;最后观察空转比例是否过高。空转比例来自末尾不足一个 group 的元素,通常可接受;若大多数 invocation 都因越界返回,dispatch 规模就需要重新计算。
Compute Shader 的并行模型强调“用 id 映射数据”,并让 group 内线程在必要时协作。读懂一个 compute pass 时,不应从 shader 数学公式开始,而应先把 dispatch 尺寸、workgroup 尺寸和数据规模对齐。对齐之后,公式才有明确执行范围。
25.3 通用计算任务在 Graphics API 中的实现
通用计算任务进入 Graphics API 后,会被拆成四个工程对象:compute pipeline、descriptor 或 bind group、dispatch command、输出资源。pipeline 描述执行哪个 compute entry point;descriptor 描述 shader 能访问哪些 buffer 和 texture;dispatch command 描述启动多少 work group;输出资源决定后续 pass 如何消费结果。这个结构在 Vulkan、Direct3D、Metal 和 WebGPU 中名称不同,读写关系相同。
粒子更新的最小 API 路径可以这样理解:创建 particle buffer,创建包含 compute shader 的 pipeline,把 particle buffer 和参数 buffer 绑定到 pipeline layout,录制 dispatch 命令,随后在 render pass 中绑定同一个 particle buffer 作为 vertex buffer、storage buffer 或 instance 数据。工具复查时,compute draw event 周围应能看到 pipeline、resource binding、dispatch size 和后续 render event。
这张图说明了通用计算在图形工程中的位置:Compute Shader 本身只负责计算,API 命令负责把资源绑定到正确阶段,barrier 或等价同步负责让写入结果对后续读取可见。把这三个职责分开,排查时就不会把 shader 公式、descriptor 绑定和资源状态混成同一个错误。
prefix sum 是 Compute Shader 中最常见的基础任务之一,它把一组标记值转换为写入位置。以 GPU culling 为例,每个对象先被计算为可见或不可见,prefix sum 把可见标记累计成紧凑索引,后续 compaction pass 把可见对象写入 indirect draw buffer。这个流程能把“每个对象独立判断”变成“只绘制可见对象”的 GPU 内部路径。
buffer compaction 是 prefix sum 的直接消费者。粒子系统中也会使用 compaction:每个粒子先根据生命周期写出 alive = 0/1,prefix sum 计算存活粒子位置,compaction pass 把存活粒子移动到紧凑 buffer。后续 render pass 只绘制紧凑后的数量。这里的性能收益来自减少后续渲染读取和绘制的元素数量,额外成本来自多个 dispatch、临时 buffer 和同步。
image filter 把二维 dispatch 映射到 texture 坐标。一个常见任务是后处理模糊、降噪或亮度提取。每个 invocation 处理一个像素或一小块像素,读取 source texture,写入 storage texture。二维任务的关键是用 global_invocation_id.xy 对应像素坐标,并处理边界采样。若 filter 需要邻域重复读取,shared memory tile 能降低重复 texture fetch;若 filter 只是逐像素颜色变换,直接读写通常更简单。
simulation 任务通常需要更多状态。粒子、布料、流体和 flocking 都会把位置、速度、约束、邻域结构放入多个 buffer。工程上要先把状态拆成稳定布局,再安排每个 pass 的读写职责。布料模拟常见拆分是外力积分、约束迭代、碰撞修正和法线更新。每个 pass 都应能说清输入、输出和依赖,否则多次迭代会产生读写冲突或结果抖动。
culling 任务连接 Compute Shader 和间接绘制。输入是对象 bounds、camera planes、LOD 参数和上一阶段可见性信息;输出可能是 visible index buffer、draw count buffer 或 indirect command buffer。这个路径的收益来自减少后续 vertex work、fragment overdraw 和 CPU draw submission。代价是 compute pass 自身的扫描、写入和同步成本。小场景或低 draw count 场景中,CPU culling 往往已经足够;大场景、实例数量高或数据已经在 GPU 上时,GPU culling 更容易形成收益。
这些任务的实现都可以套用同一个检查表:数据是否按 buffer 或 texture 清晰布局;每个 invocation 的 index 是否由 dispatch id 推导;写入是否存在竞争;竞争是否用 atomic、prefix sum、分组归约或多 pass 分解处理;输出是否被后续 pass 读取;读取前是否声明资源依赖。这个检查表比记住某个 API 的函数名更可迁移,因为函数名会变化,资源读写和执行依赖的结构稳定。
贯穿案例如果扩展到“只绘制存活粒子”,就会形成三段 compute path:第一段更新粒子并写出 alive 标记,第二段 prefix sum 生成存活索引,第三段 compaction 写入 visible particle buffer。render pass 读取 visible particle buffer 和 draw count。此时 frame capture 中应出现多个 compute dispatch,且每个 dispatch 的输出都是下一段的输入。若画面粒子数量错误,排查顺序应从 alive 标记开始,再查 prefix sum,最后查 compaction 和 draw count。
25.4 结合渲染管线的混合同步模式
Compute Shader 和 render pass 共享资源时,同步负责表达两个事实:前一个阶段的写入已经完成,后一个阶段能看到这些写入。显式 API 中,这通常体现为 resource state transition、memory barrier、pipeline barrier、queue ownership transfer 或 frame graph dependency。隐式 API 会隐藏一部分细节,但资源读写顺序仍然存在。
粒子案例中最小依赖是:compute pass 对 particle buffer 执行 storage write,render pass 读取同一个 buffer 作为 vertex input 或 shader resource。若缺少依赖,render pass 可能读取上一帧数据、部分写入数据或未定义结果。可见症状包括粒子延迟一帧、闪烁、局部位置跳变、某些平台稳定而另一些平台错误。显式同步把这种依赖写成可验证的 API 状态。
// 伪代码:表达 compute write -> graphics read 的资源依赖。
cmd.bindComputePipeline(updateParticlesPipeline);
cmd.bindResources(particleBuffer, paramsBuffer);
cmd.dispatch(groupCount, 1, 1);
cmd.resourceBarrier({
.resource = particleBuffer,
.beforeAccess = ShaderStorageWrite,
.afterAccess = VertexOrShaderRead,
.beforeStage = ComputeShader,
.afterStage = VertexShader,
});
cmd.beginRenderPass(framebuffer);
cmd.bindRenderPipeline(particleRenderPipeline);
cmd.bindVertexBuffer(particleBuffer);
cmd.draw(particleCount);
cmd.endRenderPass();
这段伪代码不绑定某个具体 API。它表达的是同步语义:compute 阶段写入 particle buffer,后续 vertex 阶段读取 particle buffer。Vulkan 可能使用 VkBufferMemoryBarrier2 和 vkCmdPipelineBarrier2 表达;Direct3D 12 可能使用 resource barrier;Metal 可能通过 encoder 边界、resource usage 和 fence/event 组合表达;WebGPU 由 pass 编码和实现层验证一部分资源使用规则。跨 API 对齐时,先对齐语义,再翻译函数。
resource state transition 解决“资源以哪种用途被访问”的问题。一个 texture 作为 storage texture 写入后,又作为 sampled texture 被 fragment shader 读取,需要从 unordered/storage write 语义过渡到 shader read 语义。一个 buffer 作为 storage buffer 写入后,又作为 indirect command buffer 使用,需要过渡到 indirect argument read 语义。状态名随 API 变化,读写方向和消费者阶段是稳定维度。
barrier 的粒度会影响性能。过宽的 barrier 会让 GPU 在更多阶段之间等待,降低并行度;过窄的 barrier 会留下数据可见性缺口。粒子更新只需要让 particle buffer 的 storage write 对后续 vertex read 可见,barrier 不应顺手阻塞所有资源、所有阶段和所有队列。实际工程中,frame graph 会根据 pass 的读写声明自动生成依赖,手写命令时则需要开发者明确写出资源、访问类型和阶段范围。
async compute 允许 compute queue 和 graphics queue 在合适条件下重叠执行。它的收益来自把独立 compute work 安排到 graphics work 的空档中,例如前一帧后处理、粒子模拟、light list 构建或 culling 与其他 graphics pass 重叠。收益前提是资源依赖允许重叠、GPU 有可用执行资源、queue 切换和同步成本小于重叠收益。若 compute pass 的结果马上被下一个 draw 读取,异步空间很小;若它服务后续较晚的 pass,调度空间更大。
queue ownership transfer 处理资源跨队列使用的所有权问题。某些 API 或平台中,compute queue 写完的资源交给 graphics queue 前,需要 semaphore、fence 或 ownership transfer。这个步骤表达两个层面:执行顺序上 graphics queue 等待 compute queue 完成;内存可见性上 graphics queue 能读取 compute 写入结果。把这两个层面分清,才能判断问题来自“还没执行完”还是“执行完但读取视图没有建立”。
frame graph 是组织混合同步的工程工具。每个 pass 声明读写资源和访问方式,系统根据声明生成 barrier、aliasing 决策和队列调度。粒子案例中,SimulateParticles pass 声明写 particleBuffer,DrawParticles pass 声明读 particleBuffer。frame graph 可以自动插入 compute-write 到 vertex-read 的依赖。调试时,开发者应检查 pass 声明是否准确,而不只检查最终生成的 barrier。
混合同步的排查顺序可以固定为五步。第一步,列出每个 pass 对资源的读写方向。第二步,标出写入者和第一个读取者。第三步,确认二者之间存在执行顺序依赖。第四步,确认内存访问类型覆盖写入和读取。第五步,观察 barrier 是否过宽造成 GPU bubble。这个顺序能同时处理正确性和性能:先保证结果可见,再缩小同步范围。
Compute Shader 与渲染管线结合时,最容易出错的位置通常不在计算公式,而在资源生命周期和同步声明。公式决定“算出什么”,barrier 决定“谁能在什么时候看到结果”。本章贯穿案例的最终闭环是:dispatch 更新 particle buffer,依赖声明让写入可见,render pass 读取同一 buffer 并产生当前帧图像。
最小自检任务
给定一个粒子系统:particleBuffer 中有 65537 个粒子,Compute Shader 使用 @workgroup_size(128) 更新位置,render pass 随后把同一个 buffer 作为 vertex buffer 绘制粒子。请完成三件事:计算 dispatch group count;说明 shader 内为什么需要边界检查;写出从 compute pass 到 render pass 的最小资源依赖语义。可以使用伪代码或自然语言,不需要依赖特定 API。
答案要点
dispatch group count 应为 (65537 + 128 - 1) / 128 = 513。实际启动的 invocation 数为 513 * 128 = 65664,多出的 127 个 invocation 需要在 shader 中通过 if (index >= particleCount) return; 退出,从而让写入范围保持在 buffer 有效元素内。资源依赖语义应表达:compute stage 对 particleBuffer 执行 storage write,render pass 的 vertex stage 或等价读取阶段随后读取同一个 buffer;两者之间需要执行顺序依赖和内存可见性依赖。若用伪代码表示,可以写成 ComputeShader/StorageWrite -> VertexShader/VertexOrShaderRead,资源对象限定为 particleBuffer。
本章知识点总结
- 执行定位:Compute Shader 是由 dispatch 启动的 GPU 并行数据处理阶段,核心输出通常写回 buffer 或 texture。
- 任务边界:适合 Compute Shader 的任务具有大量同构元素、GPU 常驻数据和后续 GPU 消费者。
- 资源闭环:分析 compute pass 时,应追踪输入资源、执行网格、输出资源和后续 pass 消费路径。
- 线程映射:dispatch 提交 work group 数量,shader 声明 work group 内 invocation 数量,
global_invocation_id映射到数据 index。 - 越界保护:元素数量通常无法整除 group size,shader 需要用 count 检查保护末尾多出的 invocation。
- Group Size:work group size 需要同时考虑设备 limit、subgroup 粒度、寄存器和 shared memory 对 occupancy 的影响。
- 共享内存:shared memory 适合 group 内重复读取同一批数据的任务,普通逐元素粒子积分通常收益有限。
- 通用任务:prefix sum、compaction、culling、image filter 和 simulation 都可拆成 pipeline、binding、dispatch 和输出资源。
- 写入竞争:多个 invocation 写同一输出位置时,需要 atomic、prefix sum、分组归约或多 pass 分解来稳定结果。
- 同步语义:compute 写入被 render pass 读取前,需要表达执行顺序和内存可见性。
- 状态转换:resource state transition 的稳定维度是资源对象、写入访问、读取访问和消费者阶段。
- Barrier 粒度:barrier 应覆盖必要资源和必要阶段,过宽会制造等待,过窄会留下可见性缺口。
- 异步计算:async compute 的收益依赖资源独立性、队列调度空间和同步成本。
- Frame Graph:frame graph 通过 pass 的读写声明生成依赖,调试时应先检查声明是否准确。
- 排查顺序:Compute Shader 问题应先对齐 dispatch 与数据规模,再检查资源绑定、写入竞争和跨 pass 同步。