WebGPU Subgroups：SIMD 级并行与 subgroup 操作

说明：原文正文仅有“TBD”（待定）占位符，无实质段落可逐句编译。以下内容为基于 WebGPU / WGSL 规范及公开资料的深度解读，并保留核心代码示例，帮助读者理解 Subgroups 这一 GPU 特性。

---

作者与出处

• 英文原作者：webgpufundamentals Contributors
• 原始链接：https://webgpufundamentals.org/webgpu/lessons/webgpu-subgroups.html

---

什么是 Subgroups？

在 GPU 计算模型中，Compute Shader 将大量线程组织为 Workgroup。而 Subgroup 是比 Workgroup 更细粒度的执行单元，通常对应硬件的 SIMD 宽度（例如 32 或 64 个线程）。同一 Subgroup 内的线程可以高效地交换数据、进行归约操作或广播，而无需通过共享内存和屏障。

WebGL 与 WebGPU 的差异
WebGL 本身并没有 Subgroup 概念（甚至 WebGL 2 也不支持 Compute Shader）。Subgroup 是 Vulkan / OpenGL / DirectX 中较新硬件的特性，WebGPU 将其作为可选特性引入。GLSL 中通过扩展 GL_KHR_shader_subgroup 提供 gl_SubgroupInvocationID、subgroupAdd 等。

---

WebGPU 中的 Subgroups 支持

Subgroups 在 WebGPU 中不是默认开启的，需要通过设备特性请求：

const adapter = await navigator.gpu.requestAdapter();
if (!adapter.features.has('subgroups')) {
  console.warn('Subgroups not supported');
}
const device = await adapter.requestDevice({
  requiredFeatures: ['subgroups'],
});

在 WGSL 代码中，必须显式启用：

enable subgroups;

启用后，可使用以下内置输入和函数：

• @builtin(subgroup_size)：当前子组的大小（如 32、64）
• @builtin(subgroup_invocation_id)：当前调用在子组内的索引
• subgroupAdd(value) / subgroupMin(value) / subgroupMax(value) / subgroupMul(value)：子组内归约操作
• subgroupBroadcast(value, id)：将某个调用的值广播到整个子组
• subgroupShuffle(value, id)：与指定通道交换数据
• subgroupBallot(predicate)：生成一个位掩码，指示哪些调用满足条件

---

核心代码示例

下面是一个使用 Subgroup 进行高效求和的 Compute Shader 示例：

enable subgroups;

@group(0) @binding(0) `var<storage, read>` inputs: `array<f32>`;
@group(0) @binding(1) `var<storage, write>` outputs: `array<f32>`;

@compute @workgroup_size(64)
fn main(
  @builtin(subgroup_size) subSize: u32,
  @builtin(subgroup_invocation_id) subInvocationId: u32,
  @builtin(local_invocation_index) localIndex: u32,
) {
  // 每个调用读取一个元素
  let value = inputs[localIndex];

  // 子组内全部元素求和
  let sum = subgroupAdd(value);

  // 每个子组中第一个调用将结果写出
  if (subInvocationId == 0u) {
    let subgroupId = localIndex / subSize;
    outputs[subgroupId] = sum;
  }
}

注意：上述代码只负责子组内归约，最终跨子组的归约仍需在 Workgroup 层面继续处理，或者使用第二个 Pass。

---

三个核心优化要点

1. 用 Subgroup 代替 Shared Memory + Barrier 进行归约

传统 Workgroup 归约通常需要 workgroupBarrier() 和 arrays，这会增加同步开销和内存延迟。利用 subgroupAdd() 等内置函数可以在一个时钟周期内完成子组内归约，大幅减少共享内存访问次数，尤其适合直方图、前缀和、粒子系统等算法。

2. 避免子组内分支发散（Divergence）

GPU 以 SIMD 方式执行子组，如果子组内不同调用走向不同分支，硬件会串行执行每个分支（屏蔽其他通道），导致性能下降。应尽量保持子组内所有调用走相同控制流。例如，处理边界条件时使用算术操作（select、mix）而非 if/else，或者将分支条件设计为“全真/全假”。

3. 适配硬件子组大小，避免硬编码

不同 GPU 的 Subgroup 大小不同：NVIDIA 通常为 32，AMD 通常为 64，Intel 可能为 8、16 或 32。假设固定大小会导致非最佳占用率或资源浪费。应使用运行时查询 @builtin(subgroup_size) 动态调整 Workgroup 大小或算法路径，例如：

let numSubgroups = (workgroupSize + subSize - 1u) / subSize;

这样代码可在不同硬件上保持良好性能。

---

GLSL / WebGL 专业术语对照检查

英文	中文推荐	说明
Subgroup	子组	比 Workgroup 更小的 SIMD 执行单元
Workgroup	工作组	WebGPU Compute Shader 中的线程组
Invocation	调用 / 实例	着色器的一次执行线程
Compute Shader	计算着色器	用于通用并行计算的着色器阶段
Shader	着色器	GPU 程序通用术语
Storage Buffer	存储缓冲区	WGSL 中可读写的大型数据缓冲区
Barrier	屏障	同步多个调用/线程的执行
Divergence	分支发散	同一子组内线程走向不同分支的现象
Quad	四方块	像素着色器中 2×2 的片段块
Ballot	投票掩码	表示哪些线程满足条件的位掩码
Shuffle	穿梭操作	子组内线程间直接交换数据
GL_KHR_shader_subgroup	GL_KHR_shader_subgroup	OpenGL / GLSL 的 Subgroup 扩展
gl_SubgroupInvocationID	gl_SubgroupInvocationID	GLSL 中子组内调用索引
subgroupAdd	subgroupAdd	子组内加法归约函数

注意：WebGPU 使用 WGSL 而非 GLSL；Subgroup 特性在 WebGL 中不可用，仅用于类比 Vulkan/OpenGL 概念。

---

结语

WebGPU Subgroups 为 GPU 并行计算带来了更细粒度的通信能力，能够显著降低归约、扫描等算法的延迟。但因其是可选特性，开发时务必做好适配和降级方案。理解硬件执行模型，善用子组操作，才能真正发挥新硬件的性能潜力。