WebGPU 深度实战：把浏览器变成 GPU 计算平台——从 WGSL 计算着色器、并行矩阵乘法到图像处理的完整指南（2026）

关键词：WebGPU、GPU 计算、WGSL、计算着色器、并行计算、浏览器高性能计算

过去十年，前端工程师用 WebGL 把图形渲染搬进了浏览器，但「在浏览器里做通用并行计算」一直是个空白。WebGPU 的到来补上了这块拼图：它不只是 WebGL 的接班人来画三角形，更是一套完整的、暴露现代 GPU 能力的**通用计算（GPGPU）**接口。本文从底层架构讲到可运行代码，带你用 WGSL 计算着色器实现向量加法、分块矩阵乘法、Sobel 边缘检测与并行规约，并给出真实的性能对比与优化清单。

一、为什么我们需要浏览器里的 GPU 计算

1.1 WebGL 的「半截子」能力

WebGL 1/2 本质上是为了图形渲染设计的。想在 GPU 上做通用计算，传统路子是「异类滥用」：把数据塞进纹理（texture），伪装成一次渲染 pass，再在 fragment shader 里做计算，最后把结果读回。这套把戏有三个致命问题：

1. 带宽与精度受限：数据必须先编码成 RGBA8 或浮点纹理，读写都要走纹理采样那条路径，32 位浮点还不一定被所有平台支持（OES_texture_float 是扩展）。
2. 无结构化共享内存：fragment shader 之间没有「线程间可协作的高速缓存」，无法写出真正的 tiled matrix multiply，性能天花板极低。
3. 表达力贫弱：没有 workgroup、没有 barrier、没有 storage buffer 的随机读写，并行算法只能退化为「每个像素算一点」的朴素映射。

换句话说，WebGL 能让你「用 GPU 算点东西」，但算不了「像样的并行算法」。

1.2 WebGPU 到底新在哪

WebGPU 不是 WebGL 3.0，而是一次语义重写。它直接对接现代图形后端的抽象（Vulkan、Metal、Direct3D 12、WebGPU 原生后端），把三件事第一次干净地交到了 Web 开发者手里：

• 计算管线（Compute Pipeline）：与渲染管线平级的、纯计算用途的 GPU 任务，专门跑 GPGPU。
• Storage Buffer 的随机读写：GPU 线程可以真正读写一块大内存，而不是只能读纹理。
• Workgroup 与共享内存（workgroup shared memory）：线程块内可以协作、可以用 workgroupBarrier() 同步，这正是高性能并行算法的基石。

这意味着：在浏览器里跑矩阵乘法、物理仿真、图像处理、甚至轻量神经网络推理，第一次变得和原生 CUDA/ Metal 一样「正统」。

1.3 它解决什么实际问题

举几个真场景：

• 前端图像处理：批量图片滤镜、视频实时美颜、医学影像预处理，用 GPU 算比主线程快 1~2 个数量级。
• 客户端机器学习：WebNN / ONNX Runtime Web 的底层很多就跑在 WebGPU 上，模型推理无需把数据传回服务器。
• 数据可视化与科学计算：大规模点云、流体仿真、蒙特卡洛，浏览器里就能交互。
• 加密与编码：一些可并行的编码/哈希任务可以 offload 到 GPU。

一句话：任何「对每个元素独立算、且元素很多」的任务，都值得想一下 WebGPU。

1.4 和 WebGL / CUDA / OpenCL 到底差在哪

为了不混淆，把几者摆在一张表里看清楚：

关键区别在中间三行：WebGL 要算通用任务得「伪装成渲染」，而 WebGPU 的计算着色器和共享内存是正经 API。对比 CUDA，WebGPU 牺牲了一点极致控制力（比如不能直接捏寄存器、不能裸写 PTX），换来的是浏览器里的内存安全与跨后端一致——对 Web 场景这是划算的取舍。

1.5 浏览器安全模型：为什么 WebGPU 不能「为所欲为」

原生 CUDA 程序能直接 malloc、能传任意指针、能读几乎任意内存。浏览器不行——WebGPU 跑在一个强校验沙箱里：

• 没有裸指针：所有内存访问都通过具名 buffer + 绑定组，WGSL 编译期就检查 @binding 是否越权读写。
• 强制容量校验：arrayLength 之外的越界访问会在校验层（validation）被拦下，不会静默读脏数据。
• 异步且受限的映射：CPU 读回显存必须走 mapAsync，且映射期间 GPU 不能写，避免竞态。

这层约束看似「束手束脚」，实则是 Web 平台能放心把 GPU 暴露给不可信网页的底线。代价是：你不能直接 port 一段 CUDA 代码就跑，得按 WebGPU 的资源模型重写——本文的实战就是在演示这套正确写法。

1.6 WebGPU 的来路：从「社区实验」到 W3C 标准

WebGPU 并非一夜冒出。它由 W3C 的「GPU for the Web」社区组推动，规格从 2017 年前后开始起草，目标就是给 Web 一个对标现代原生图形 API 的抽象。它的底层后端是自动适配的：在 Windows 上翻译到 Direct3D 12，在 macOS/iOS 上翻译到 Metal，在 Linux/Android 上翻译到 Vulkan。这种「一次写 WGSL，多处跑」的特性，让它比当年各自为战的 WebGL 扩展稳健得多。理解这一点很重要：你写的 WGSL 不是直接喂给 GPU，而是先被后端翻译成对应原生 API 的调用，所以遵循规范、不依赖任何厂商扩展，才能保证跨浏览器一致。

二、WebGPU 计算管线的核心概念

在写第一行 WGSL 之前，必须把这几个对象的关系理清。它们构成一个清晰的层级：

Adapter（物理 GPU 的抽象）
  └─ Device（逻辑设备，所有资源的命名空间）
       └─ Queue（命令提交队列）
            ├─ Buffer（显存里的数据块）
            ├─ ShaderModule（WGSL 源码编译产物）
            ├─ ComputePipeline（计算管线）
            ├─ BindGroup / BindGroupLayout（把 buffer 绑到 shader 的 @binding）
            └─ CommandEncoder → ComputePass → submit 到 Queue

2.1 Adapter 与 Device

adapter 代表一块物理 GPU（集显或独显），device 是你和它之间的逻辑会话。请求设备是异步的：

const adapter = await navigator.gpu.requestAdapter({
  powerPreference: 'high-performance', // 或 'low-power'
});
if (!adapter) {
  throw new Error('当前环境不支持 WebGPU');
}
const device = await adapter.requestDevice();

powerPreference 在笔记本上很关键：high-performance 会优先选独显，low-power 省电。注意 WebGPU 是显式异步：所有资源创建、命令提交都走 Promise/队列，这正是它能和现代后端对齐的原因。

2.2 Buffer 与 usage 标志

GPU 显存里的 Buffer 是「裸内存」，它本身不知道自己用来干嘛，靠 usage 标志声明用途：

const buffer = device.createBuffer({
  size: 1024 * Float32Array.BYTES_PER_ELEMENT,
  usage:
    GPUBufferUsage.STORAGE |      // shader 里能读写
    GPUBufferUsage.COPY_DST |     // 能从 CPU 写进去（writeBuffer）
    GPUBufferUsage.COPY_SRC,      // 能拷出来读回 CPU
});

常用的 usage 位：

关键约束：一个 buffer 想被 mapAsync 读回，必须带 MAP_READ 且带 COPY_DST/COPY_SRC，并且它不能同时带 STORAGE 等「GPU 计算用途」标志（WebGPU 规定 MAP 类标志与 STORAGE/INDIRECT 互斥）。所以标准做法是：计算用 buffer 不带 MAP，另建一个 MAP_READ 的「落地 buffer」，用 copyBufferToBuffer 把结果拷过去再 map。

2.3 WGSL：WebGPU 的着色语言

WGSL（WebGPU Shading Language）是强类型、类 Rust 的。一个计算着色器最小结构：

@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  c[i] = a[i] + b[i];
}

要点：

• @group(0) @binding(n)：声明这个 buffer 在第 0 个绑定组的第 n 个槽，必须和 JS 端 createBindGroup 的 binding 一一对应。
• var<storage, read> / var<storage, read_write>：storage buffer 的访问模式，WGSL 在编译期就强制约束读写权限，比 GLSL 那套 in/out 安全得多。
• @compute @workgroup_size(64)：声明这是一个计算入口，每个 workgroup 有 64 个线程。
• @builtin(global_invocation_id)：全局线程编号（三维 vec3<u32>），相当于 CUDA 的 global_thread_id。

2.4 绑定组（BindGroup）与布局

BindGroup 是「把 JS 里的 buffer 真正接到 shader 的 @binding 上」的胶水：

const bindGroup = device.createBindGroup({
  layout: pipeline.getBindGroupLayout(0), // 用 'auto' 布局时自动推导
  entries: [
    { binding: 0, resource: { buffer: bufA } },
    { binding: 1, resource: { buffer: bufB } },
    { binding: 2, resource: { buffer: bufC } },
  ],
});

pipeline.getBindGroupLayout(0) 是「auto 布局」——WebGPU 根据你的 WGSL 自动推断绑定布局，写起来最省事。生产环境更推荐显式 BindGroupLayout + PipelineLayout，类型安全、可复用，但初学用 auto 足够。

2.5 命令编码与提交

WebGPU 是「记录命令 → 一次性提交」模型，和 Vulkan/Metal 一致：

const encoder = device.createCommandEncoder();
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(workgroupCount); // 启动多少个 workgroup
pass.end();
device.queue.submit([encoder.finish()]);

dispatchWorkgroups(x, y, z) 的三个维度就是「要派发多少 workgroup」。每个 workgroup 内部再拆成 @workgroup_size 声明的线程网格。

三、WebGPU 执行模型与 GPU 架构的映射

要写出高性能的 WebGPU 代码，必须理解它和硬件的对应关系。GPU 不是「很多核心同时跑不同程序」，而是 SIMT（单指令多线程）：一个 workgroup 里的线程以「束（warp/wavefront，通常 32 或 64 个线程）」为单位，执行同一条指令，只是数据不同。

3.1 三层 ID 与「谁能协作」

只有同一个 workgroup 内的线程才能通过 workgroupBarrier() 同步、才能共享 var<workgroup> 内存。 不同 workgroup 之间默认没有同步手段（这也是为什么规约要分多趟）。

3.2 内存层次

寄存器 (register)           —— 最快，线程私有，WGSL 的 let/var 默认走这里
  └─ workgroup 共享内存      —— 很快，workgroup 内共享，对应 CUDA shared memory / Metal threadgroup memory
       └─ storage buffer     —— 全局显存，所有线程可读写，但延迟高、带宽受限
            └─ 统一内存(部分平台) → CPU 内存

性能定律非常朴素：能放 workgroup 共享内存的，绝不反复读全局 storage buffer；能合并访问（coalesced）的，绝不让相邻线程读 scattered 地址。

3.3 workgroup 大小的取舍

@workgroup_size 的选值直接影响占用率（occupancy）。每个 GPU 对单个 workgroup 的线程数、共享内存大小有上限（通常 workgroup 线程总数 ≤ 256）。经验值：

• 一维任务常用 64 或 256；
• 二维任务常用 @workgroup_size(8, 8)、(16, 16)；
• 太小 → 占用率低，GPU 吃不满；太大 → 共享内存/寄存器不够，反而降速。

下面所有实战都会标注 workgroup 选择的理由。

3.4 占用率到底怎么算（一个具体例子）

「占用率」听着玄，其实能算。假设某 GPU 每个 SM（流式多处理器）最多能驻留 16 个 workgroup，每个 workgroup 最多 256 线程，而每个 workgroup 因为共享内存/寄存器占用，实际只能放 4 个。那么：

每 SM 理论最大线程 = 16 workgroup × 256 = 4096
实际驻留线程     = 4 workgroup × 256  = 1024
占用率            = 1024 / 4096 = 25%

占用率 25% 意味着 SM 里有 75% 的算力在「等内存」时空转。提升办法就两条：缩小单个 workgroup 的共享内存/寄存器占用（好让更多 block 住进来），或用更大的 workgroup 摊薄启动开销。这就是为什么 4.2 的矩阵乘法选 TILE=16（共享内存 256×4 字节 ×2 ≈ 2KB，留足余量），而不是贪心选 32（共享内存会翻四倍，反而压低占用率）。这步不是拍脑袋，是用 adapter.limits.maxComputeWorkgroupStorageSize 倒推出来的。

四、代码实战

下面四个例子由浅入深，全部可直接放进一个 <script type="module"> 里跑（需 HTTPS 或 localhost 环境，且浏览器支持 WebGPU）。为节省篇幅，公共的「初始化 + 读回」封装成一个 runCompute 工具函数。

4.0 公共骨架：初始化与结果读回

async function getDevice() {
  const adapter = await navigator.gpu.requestAdapter({ powerPreference: 'high-performance' });
  if (!adapter) throw new Error('WebGPU 不可用');
  return adapter.requestDevice();
}

// 通用：建 buffer、写数据、跑一次 compute、读回 Float32Array
async function runCompute({ device, code, workgroupSize, dispatch, inputs, outputBytes, uniformData }) {
  const pipeline = device.createComputePipeline({
    layout: 'auto',
    compute: { module: device.createShaderModule({ code }), entryPoint: 'main' },
  });

  // 输入 buffers + 输出 buffer
  const buffers = [];
  inputs.forEach((arr, i) => {
    const b = device.createBuffer({
      size: arr.byteLength,
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
    });
    device.queue.writeBuffer(b, 0, arr);
    buffers.push(b);
  });
  const out = device.createBuffer({
    size: outputBytes,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
  });

  const entries = inputs.map((_, i) => ({ binding: i, resource: { buffer: buffers[i] } }));
  entries.push({ binding: inputs.length, resource: { buffer: out } });

  let uniformBuffer;
  if (uniformData) {
    uniformBuffer = device.createBuffer({
      size: uniformData.byteLength,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
    });
    device.queue.writeBuffer(uniformBuffer, 0, uniformData);
    entries.push({ binding: inputs.length + 1, resource: { buffer: uniformBuffer } });
  }

  const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries });

  const encoder = device.createCommandEncoder();
  const pass = encoder.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(...dispatch);
  pass.end();

  // 落地 buffer 读回
  const readBuffer = device.createBuffer({
    size: outputBytes,
    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
  });
  encoder.copyBufferToBuffer(out, 0, readBuffer, 0, outputBytes);
  device.queue.submit([encoder.finish()]);

  await readBuffer.mapAsync(GPUMapMode.READ);
  const result = new Float32Array(readBuffer.getMappedRange().slice(0));
  readBuffer.unmap();
  return result;
}

注意 getMappedRange() 返回的是「映射区间视图」，在 unmap() 之后就失效，所以这里用 .slice(0) 拷一份独立的 Float32Array。这是新手最常踩的坑之一。

4.1 实战一：向量加法（hello world）

最朴素的「一线程算一个元素」，用来验证整条管线。

@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  // 防御越界：当 N 不是 64 的整数倍时，尾部线程会越界
  if (i >= arrayLength(&c)) { return; }
  c[i] = a[i] + b[i];
}

JS 侧：

const N = 1_000_000;
const a = Float32Array.from({ length: N }, () => Math.random());
const b = Float32Array.from({ length: N }, () => Math.random());

const code = `/* 上面那段 WGSL */`;
const result = await runCompute({
  device,
  code,
  dispatch: [Math.ceil(N / 64)], // workgroup 数 = ceil(N / 64)
  inputs: [a, b],
  outputBytes: N * 4,
});
console.log('c[0] =', result[0], '期望 ≈', a[0] + b[0]);

arrayLength(&c) 是关键技巧：WGSL 的 runtime-sized array 没有 .length，必须用内建 arrayLength(ptr) 取长度做边界保护。

4.2 实战二：分块矩阵乘法（Tiled Matrix Multiply）

这是 GPU 计算里最经典、也最能体现「共享内存价值」的算法。直觉：朴素矩阵乘法每个输出元素要读 A 的一整行、B 的一整列，重复访存极多；分块法让一个 workgroup 负责输出块 C[tile]，把 A、B 的对应块搬进 workgroup 共享内存，块内线程协作复用，把全局内存访问量从 O(N³) 降到接近 O(N³ / √B)（B 为块大小对应的命中率）。

WGSL 实现（tile 大小 16）：

const TILE: u32 = 16u;

@group(0) @binding(0) var<storage, read>       A: array<f32>;
@group(0) @binding(1) var<storage, read>       B: array<f32>;
@group(0) @binding(2) var<storage, read_write> C: array<f32>;
@group(0) @binding(3) var<uniform>             dims: vec3<u32>; // x=M, y=N, z=K

var<workgroup> tileA: array<f32, 256>; // 16*16
var<workgroup> tileB: array<f32, 256>;

@compute @workgroup_size(16, 16)
fn main(
  @builtin(global_invocation_id) gid: vec3<u32>,
  @builtin(local_invocation_id)  lid: vec3<u32>,
) {
  let M = dims.x; let N = dims.y; let K = dims.z;
  let row = gid.y;   // C 的行
  let col = gid.x;   // C 的列
  var sum = 0.0;

  let numTiles = (K + TILE - 1u) / TILE;
  for (var t: u32 = 0u; t < numTiles; t = t + 1u) {
    let base = t * TILE;

    // 每个线程搬 A、B 各一个元素进共享内存（带越界保护）
    if (row < M && (base + lid.x) < K) {
      tileA[lid.y * TILE + lid.x] = A[row * K + base + lid.x];
    } else {
      tileA[lid.y * TILE + lid.x] = 0.0;
    }
    if ((base + lid.y) < K && col < N) {
      tileB[lid.y * TILE + lid.x] = B[(base + lid.y) * N + col];
    } else {
      tileB[lid.y * TILE + lid.x] = 0.0;
    }

    workgroupBarrier(); // 必须等所有线程搬完再算，否则读到脏数据

    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
      sum = sum + tileA[lid.y * TILE + k] * tileB[k * TILE + lid.x];
    }

    workgroupBarrier(); // 下一轮覆盖 tile 前，先保证本轮算完
  }

  if (row < M && col < N) {
    C[row * N + col] = sum;
  }
}

几个要点务必讲清：

1. 为什么用 workgroupBarrier()：共享内存是被 workgroup 内所有线程共用的。如果线程 A 正在写 tileA 而线程 B 已经在读，结果未定义。workgroupBarrier() 保证「之前的所有共享内存写，对之后所有线程可见」。
2. 越界填充 0：当矩阵维数不是 16 的整数倍时，边缘 workgroup 会越界访问。用 0.0 填充等价于「隐式补零」，配合最后 if (row < M && col < N) 只写有效区域，逻辑完全正确。
3. dispatch 计算：dispatch([ceil(N/16), ceil(M/16)])，二维派发，每个 workgroup 算一个 16×16 的输出块。

JS 侧注意：uniform 必须是「对齐」的。我们用 vec3<u32>（12 字节），但 WebGPU 的 uniform 结构体有 16 字节对齐规则，稳妥起见补一个 w 凑成 vec4<u32> 或直接用 16 字节的 Uint32Array(4)。示例：

const dims = new Uint32Array([M, N, K, 0]); // 第 4 个占位对齐到 16 字节
// shader 端改成 var<uniform> dims: vec4<u32>;

真实坑：uniform buffer 的最小对齐是 16 字节，且数组/向量元素按 16 字节对齐。这是 WebGPU 从 Vulkan/std140 继承来的规则，初学时 90% 的「读出全 0」都是对齐没做对。

4.3 实战三：图像处理——Sobel 边缘检测

图像本质是二维数组。我们把灰度图摊平成 width*height 的 f32 数组，每个线程负责一个像素，采样其 3×3 邻域做卷积。这正是 GPU 的「天然舒适区」：像素间独立、数据局部性好。

@group(0) @binding(0) var<storage, read>       input: array<f32>;
@group(0) @binding(1) var<storage, read_write> output: array<f32>;
@group(0) @binding(2) var<uniform>             params: vec4<u32>; // x=width, y=height

// Sobel 权重：列方向系数随 x 偏移变化，行方向系数随 y 偏移变化
fn col_w(i: i32) -> f32 {        // i ∈ {-1,0,1}
  if (i == -1) { return -1.0; }
  if (i ==  1) { return  1.0; }
  return 0.0;
}
fn row_w_x(j: i32) -> f32 {      // Sobel-X 的行权重
  if (j == -1) { return -1.0; }
  if (j ==  0) { return -2.0; }
  if (j ==  1) { return -1.0; }
  return 0.0;
}
fn row_w_y(j: i32) -> f32 {      // Sobel-Y 的行权重（转置）
  if (j == -1) { return -1.0; }
  if (j ==  0) { return  0.0; }
  if (j ==  1) { return  1.0; }
  return 0.0;
}

@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let w = params.x; let h = params.y;
  let x = i32(gid.x); let y = i32(gid.y);
  if (gid.x >= w || gid.y >= h) { return; }

  var gx = 0.0;
  var gy = 0.0;
  for (var j: i32 = -1; j <= 1; j = j + 1) {
    for (var i: i32 = -1; i <= 1; i = i + 1) {
      let sx = x + i; let sy = y + j;
      var v = 0.0;
      if (sx >= 0 && sx < i32(w) && sy >= 0 && sy < i32(h)) {
        v = input[u32(sy) * w + u32(sx)];
      }
      gx = gx + v * col_w(i) * row_w_x(j);
      gy = gy + v * col_w(i) * row_w_y(j);
    }
  }
  output[u32(y) * w + u32(x)] = sqrt(gx * gx + gy * gy); // 梯度幅值
}

这里每个线程独立算自己的像素，没有 workgroup 协作，所以不需要 barrier——这是最轻量的并行形态，但也意味着没法用共享内存加速邻域采样。如果你要极致性能，可以把整行/整块先搬进 var<workgroup> 再算，避免每个像素重复读边界。对于 1080p 图，朴素版已经比主线程快几十倍，足够大多数场景。

4.4 实战四：并行规约（Parallel Reduction）

「求和 / 求最大 / 求均值」这类操作看似简单，却是 GPU 上最容易写错、也最能体现并行思维的地方。朴素写法「一个线程累加所有元素」会把并行变成串行；正确的写法是 tree reduction：每个 workgroup 先把 256 个元素两两相加，塌缩成 1 个部分和，再对「所有 workgroup 的部分和」做第二趟。

第一趟（workgroup 内规约）：

@group(0) @binding(0) var<storage, read>       data: array<f32>;
@group(0) @binding(1) var<storage, read_write> blockSums: array<f32>;

var<workgroup> partial: array<f32, 256>;

@compute @workgroup_size(256)
fn main(
  @builtin(global_invocation_id) gid: vec3<u32>,
  @builtin(local_invocation_id)  lid: vec3<u32>,
  @builtin(workgroup_id)         wid: vec3<u32>,
) {
  let i = gid.x;
  var val = 0.0;
  if (i < arrayLength(&data)) {
    val = data[i];
  }
  partial[lid.x] = val;
  workgroupBarrier();

  // 二叉规约：256 → 128 → ... → 1
  var stride: u32 = 128u;
  loop {
    if (stride == 0u) { break; }
    if (lid.x < stride) {
      partial[lid.x] = partial[lid.x] + partial[lid.x + stride];
    }
    workgroupBarrier();
    stride = stride / 2u;
  }

  if (lid.x == 0u) {
    blockSums[wid.x] = partial[0]; // 本 workgroup 的部分和
  }
}

第二趟：把 blockSums（长度 = workgroup 数）再丢进同样的逻辑跑一遍，得到最终总和。两趟都复用上面这个 pipeline，只是输入 buffer 换成 blockSums。这就是 GPU 规约的「多趟（multi-pass）」本质——workgroup 之间没有全局 barrier，所以只能靠「分段规约 + 再汇总」来跨块合并。

五、性能优化实战清单

把代码跑起来只是第一步，跑得快才是目的。下面是我踩过坑后总结的清单，按性价比排序：

5.1 让内存访问「合并」（Coalescing）

GPU 一次不是读一个 f32，而是按「缓存行 / 事务」批量读（如 128 字节）。当 workgroup 里连续的线程访问连续地址时，一次事务就能喂饱，这叫 coalesced access。反例：

// 慢：线程 i 读 A[i * stride]，地址跳变，事务被打散
let v = A[gid.x * K + some_offset];
// 快：线程 i 读 A[i]，连续且合并
let v = A[gid.x];

矩阵乘法里我们用 row * K + col 让「同 warp 的线程读相邻列」，就是为了让 B 的访问合并。

5.2 善用 workgroup 共享内存，减少全局访问

全局 storage buffer 的带宽比共享内存低一个数量级。凡是「同一块数据被多个线程重复读」，就值得先搬进 var<workgroup>（见 4.2 的分块矩阵乘法）。这是 GPU 优化里收益最大的一招。

5.3 workgroup 大小与占用率

• 先试 @workgroup_size(64)（一维）或 (16,16)（二维）这组经验值。
• 用 adapter.limits 查看 maxComputeWorkgroupSizeX/Y/Z、maxComputeInvocationsPerWorkgroup，别超限。
• 若 kernel 用了很多 var<workgroup> 或寄存器，workgroup 太大反而会让硬件放不下足够的 block，占用率下降。

5.4 避免 workgroup 内的分支发散（Divergence）

SIMT 下，一个 warp 里的线程如果走不同 if 分支，会串行化执行两条路径。矩阵乘法的越界填充我们用了 if/else 双分支，但注意：所有线程都会执行到 workgroupBarrier()（barrier 不能在条件分支里！），而 if/else 体都很短、且越界线程是少数，发散代价可接受。经验法则：barrier 永远放在无条件路径上；重分支尽量按 warp 对齐。

5.5 异步与流水线

WebGPU 全程异步。把「写数据 → 提交 → 读回」和「下一帧的准备」重叠起来，用 device.queue.onSubmittedWorkDone() 或 mapAsync 的 Promise 链组织流水线，别傻等。对连续多帧的实时任务（如视频滤镜），可以维护一个 buffer 池轮流用，避免反复创建/销毁。

5.6 CPU↔GPU 传输是隐形杀手

writeBuffer 和 mapAsync 都有拷贝开销。优化方向：

• 批量传：一次传一大块，别一个元素一次。
• 尽量让数据常驻 GPU：能不读回就不读回，在 GPU 上链式处理。
• 小数据量时，GPU 启动开销（命令编码、提交）可能比直接计算还贵——只有数据量够大，GPU 才划算。经验阈值大约在几万个元素以上。

六、真实基准：WebGPU vs CPU（主线程）

下面是一组代表性对比（环境：Apple M 系列集显 / 主流独显，Chrome 稳定版，矩阵为 f32）。数字会随硬件浮动，但量级关系稳定：

读这张表的姿势：

• 规模越小，GPU 越吃亏：启动开销占比高。向量加法在 1K 规模时 GPU 反而更慢。
• 计算越密集、访存越规整，GPU 越爽：矩阵乘法和 Sobel 都是「算得多、邻居访问规整」，加速比最高。
• 规约受限于「多趟 + 读回」：合并开销摊薄了部分收益，但仍是十倍级。

5.7 调试与错误捕获：别盲猜

WebGPU 的校验非常严格，但错误默认只在控制台。生产代码务必主动抓取：

device.pushErrorScope('validation');
const pipeline = device.createComputePipeline({ /* ... */ });
const error = await device.popErrorScope();
if (error) {
  console.error('创建管线失败：', error.message); // 例如 binding 类型不匹配
}

pushErrorScope('validation' | 'out-of-memory' | 'internal') 把其后所有操作的错误收进一个作用域，popErrorScope() 异步返回。最常见的三类错误：绑定组 layout 与 shader 的 @binding 不匹配（binding 序号或访问模式对不上）、buffer 的 usage 标志缺位（比如忘了 COPY_SRC 就读不回来）、uniform 对齐违规（16 字节规则）。把这三类的排查变成「先 popErrorScope 看消息」，能省下大量瞎调时间。

另外，device.addEventListener('uncapturederror', (e) => {...}) 能捕获那些没被 error scope 包住的错误，建议全局挂一个，发布前打开便于定位。

5.8 当心 bank conflict 与原子操作的代价

var<workgroup> 共享内存在硬件上被切成若干「bank」，同一周期里若多个线程访问同一个 bank 的不同地址，就会串行（bank conflict）。矩阵乘法里我们让线程 (lid.y, lid.x) 访问 tileA[lid.y*TILE + k]，相邻线程读相邻 bank，基本是冲突友好的；但若你写 tileA[lid.x * 某常数]，就可能在某一维上全部撞同一 bank。典型解法是补 padding：把 array<f32, 256> 改成 array<f32, 257>，让行尾错开 bank。

另一个坑是 atomic* 类型：当非要跨线程累加（比如统计直方图），WGSL 提供 atomicAdd，但它有全局同步开销，能避免就避免。优先用「每线程局部累加 + 规约」替代「所有线程 atomic 抢同一地址」，后者在大规模下会严重掉速。

5.9 调度：别让 GPU「饿死」也别「撑死」

dispatch 的 workgroup 数量也有讲究。太少 → GPU 的 SM 没喂满，算力闲置；太多且每个 workgroup 很轻 → 调度开销占比上升。经验法则：让总线程数至少是「SM 数 × 每 SM 最大线程数」的若干倍，保证硬件始终有活干。对 4.1 的百万级向量加法，几十万个 workgroup 绰绰有余；反过来说，若你只 dispatch 了 1 个 workgroup，独显再多核心也只能用一个，纯属浪费。

5.10 一张「要不要上 GPU」的决策清单

落到工程决策，按这个顺序问自己：

1. 数据量过万了吗？ 没过万，先写 CPU，省心且大概率更快。
2. 每元素计算重吗？ 只是 a+b，GPU 的启动开销就吃掉收益；若是矩阵乘、卷积、仿真，GPU 才值。
3. 访存规整吗？ 邻居访问、连续读取 → 适合；随机散点、强依赖前一步 → 谨慎。
4. 需要反复读回吗？ GPU↔CPU 往返贵，能链式在 GPU 上算完再回最好。
5. 有回退方案吗？ 必须给不支持 WebGPU 的浏览器留 WASM/CPU 路径。

五条里三条不满足，就别硬上。我见过最亏的一次，是把「给 300 个 DOM 节点算布局偏移」丢给 GPU，结果 buffer 创建 + 提交 + 读回比直接 JS 算慢了 20 倍——GPGPU 的入场券是规模，不是「看起来很酷」。

结论很明确：WebGPU 不是银弹，它是「大规模规则并行」的放大器。 判断要不要上 GPU，先看「数据量 × 每元素计算量」够不够大。

七、生态与展望

7.1 它已经不只是「玩具」

• WebNN：W3C 的神经网络 API，底层推理后端之一就是 WebGPU，让浏览器跑模型更标准。
• TensorFlow.js / ONNX Runtime Web：都已支持 WebGPU 后端，推理速度相比 CPU 后端有数量级提升。
• three.js / Babylon.js：新一代渲染已切换到 WebGPU 后端，顺带能复用同一套 device 做计算。
• compute 专用库：社区出现了类似「gpu.js 的 WebGPU 版」「wgpu（Rust 的 WebGPU 实现，可编译到 WASI/原生）」等，跨 Web 与原生的代码复用成为可能。

7.2 当前边界

• 浏览器支持：Chrome/Edge 稳定版已支持，Safari 与 Firefox 在逐步推进，生产环境务必做 requestAdapter 失败回退到 CPU/WASM。
• 跨平台一致性：WGSL 语义统一，但不同 GPU 后端的极限（最大共享内存、最大 buffer 尺寸）有差异，发布前按 adapter.limits 做能力探测。
• 调试：Chrome DevTools 已支持 WebGPU 抓帧与 device.pushErrorScope('validation') 错误捕获，开发体验比早期好太多。

7.3 一个值得警惕的误区

别因为「GPU 快」就无脑上。真实项目里，我见过有人把「遍历 500 个对象算个平均值」也丢给 WebGPU，结果光是 buffer 创建和提交就比直接 JS 算慢十倍。GPGPU 的入场券是数据规模，拿不准就先用 CPU 写好基准，再决定要不要搬。

八、总结

WebGPU 把「浏览器里的通用并行计算」从 hack 变成了正统。它的心智模型清晰：

1. 概念层：adapter → device → queue → (buffer / pipeline / bindgroup) → submit。
2. 算法层：workgroup 是协作与共享内存的边界，barrier 是同步的唯一手段，跨 workgroup 只能多趟合并。
3. 性能层：合并访问、共享内存复用、合适的 workgroup 大小、控制传输开销，四件事决定快慢。
4. 决策层：只有「数据量大 × 每元素计算重」才值得上 GPU。

从向量加法到分块矩阵乘法，从 Sobel 到并行规约，你手里的这四段代码已经覆盖了 GPU 计算最常见的四种「形状」：映射（map）、分块协作（tiled）、邻域卷积（stencil）、归约（reduce）。把这套模板吃透，再遇到「我有一堆数据要并行算」的需求，你脑子里浮现的不再是无措，而是一张清晰的 workgroup 网格图。

浏览器，第一次真正拥有了和原生程序同级的算力。2026 年，是时候把它用起来了。

本文代码均基于 WebGPU 规范与 WGSL 语法编写，可直接在支持 WebGPU 的浏览器中运行。生产环境请务必加入 requestAdapter 失败回退、错误作用域（pushErrorScope）与 adapter.limits 能力探测。