想象你走进一家生意火爆的 AI 餐厅：招牌菜——图像分类、语音识别、超分辨率——都由总店统一配送料包，服务员下单几分钟就能上桌。可有一天，一位熟客点了一道菜单上没有的菜：“我要一个动态量化版的层归一化，再融合一个带泄漏的 GELU 激活。” 领班愣住了——后厨只有标准料包，这道菜做不了。 很多前端 AI 团队就卡在这个瞬间：浏览器里预置的算子库覆盖不了那些特殊算子。有人请来一位 WASM 全能主厨——他什么菜都会做，但只有一双手，忙不过来；而 GPU 是一座拥有成千上万名帮厨的中央厨房——只要你能把菜谱写清楚，整个后厨就能以惊人的速度并行把这道菜做出来。WebGPU，就是这座厨房的菜谱系统。 这篇文章将带你从一张空白菜谱开始，一步步写出属于自己的 AI 算子：从为什么要自己写、菜谱语言长什么样，到食材怎么递进厨房、活儿怎么分给帮厨，再到整条流水线怎么搭、怎样让厨房更快，以及最后如何稳稳落地。读完后你会发现：点新菜不再需要等总店，你自己就是 AI 餐厅的主厨。

岗位关联

掌握 WebGPU 自定义算子，你就拥有了在浏览器中“自造算力”的能力——当标准算子库覆盖不到特殊算子时，你能自己写出高速 GPU 版本。这是前端 AI 工程师从“调框架”走向“造引擎”的分水岭，也是端侧推理、实时视觉特效等场景中突破性能天花板的关键钥匙。

目录
  1. 可视化演示
  2. 1. 为什么自定义算子
  3. 2. WGSL 着色器编写
  4. 3. 缓冲绑定
  5. 4. 调度组设计
  6. 5. 流水线构建
  7. 6. 性能调优
  8. 7. 工程落地
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

可视化展示WGSL着色器编写、GPU缓冲绑定与调度组设计、计算调度与工作组网格分发、以及性能调优策略对比，交互调节参数体验GPU加速效果。

本节配有交互动效（进入视口后自动加载，约 3KB）。

着色器与缓冲绑定 Uniform缓冲存储权重参数 Storage缓冲存储输入输出张量 @group(0) @binding(0) var 调度组将资源绑定到着色器

1. 为什么自定义算子

AI 餐厅总店（TensorFlow.js、ONNX Runtime Web）的料包相当齐全：卷积、池化、矩阵乘、ReLU，这些招牌算子开箱即用。可真实业务永远在菜单之外——一个新的激活函数、一种自定义的量化方案、一个融合了三个步骤的复合算子——等你找框架要时，得到的回答往往是“该算子不支持”。

更麻烦的是性能。想绕开料包自己硬算？WASM 这位全能主厨可以手动实现任何算子，但他毕竟只有一双手（单线程为主，即使起 Worker 也多核协同麻烦）。当数据量到百万级，CPU 上的逐元素循环就像一个人切一千斤土豆，再快也快不过一群帮厨。

WebGPU 的计算管线（compute pipeline）提供了第二种选择：它不关心你写的算子叫什么，只关心你给它一份“菜谱”（WGSL 着色器）。浏览器拿到菜谱后，会在 GPU 上派出成百上千的并行线程同时执行。任何能用“对每个元素做同样的事”描述的算子，都能被它加速。比如下面这个公式，就是自定义算子的“火力账本”：

所以，“自定义算子”的本质不是发明新菜，而是把菜谱写清楚、交给正确的厨房。当你掌握了这套菜谱系统，标准库覆盖不到的角落就不再是痛点，而是你的加分项——别人等总店发货，你已经在自己后厨开火了。

threadstotal=workgroup_size×workgroup_count\displaystyle \text{threads}_{\text{total}} = \text{workgroup\_size} \times \text{workgroup\_count}
公式通俗说明：GPU 这波同时有多少帮厨在干活，等于“每个灶台的人数”乘以“灶台的数量”。这是理解 WebGPU 并行模型的第一笔账。
  1. workgroup_size（工作台尺寸）每个工作台内同时开工的线程数，相当于每张灶台上帮厨的人数，在着色器里用 @workgroup_size(256) 指定。
    workgroup_size\text{workgroup\_size}
  2. workgroup_count（工作台数量）主厨一次派发的灶台总数，运行时通过 dispatchWorkgroups 决定。
    workgroup_count\text{workgroup\_count}
  3. 线程总数（threads_total）两者相乘就是本波并行执行的线程总数，也就是 GPU 这次最多能同时料理多少份食材。
    threads_total=workgroup_size×workgroup_count\text{threads\_total} = \text{workgroup\_size} \times \text{workgroup\_count}

2. WGSL 着色器编写

菜谱用什么语言写？WebGPU 不用 GLSL、不用 HLSL，而是用一门专为它量身定做的 WGSL（WebGPU Shading Language）。它比 JavaScript 更像一份“操作清单”：每条指令都在告诉帮厨“手里的这份食材该怎么处理”。

一份计算着色器菜谱分两半：头顶的“开工声明”和中间的“处理步骤”。开工声明用 @compute @workgroup_size(256) 告诉大家：这是一道需要 256 名帮厨同时开工的大锅菜。处理步骤写在 main 函数里，每个帮厨拿着自己的工号牌跑同一套流程。

工号牌是 WGSL 内置变量 global_invocation_id。为什么需要它？因为没有它，256 个人做的是完全相同的动作，菜品会重复。有了工号牌，帮厨 A 处理第 0 份食材、帮厨 B 处理第 1 份食材——每个人的活儿才各不相同。真正自定义的算子逻辑就是 main 函数里的那几行：取出自己的元素，做变换，写回结果。

上面的菜谱实现了这样一个自定义元素级算子——它可以是一个融合了线性变换与激活函数的新算子：

回到厨房：一份菜谱写得好不好，决定帮厨们会不会手忙脚乱。WGSL 菜谱的每一行都“只说一件事”，剩下的交给 GPU 的千人后厨去跑——这正是自定义算子的第一块基石：把算法翻译成帮厨能看懂的分步指令。

yi=max⁡(0, 2xi+1)\displaystyle y_i = \max(0, \, 2x_i + 1)
公式通俗说明：把每个元素先“翻倍再调味”，然后像温度计一样，低于 0 的读数一律记为 0。这就是一个典型的自定义激活算子。
  1. x_i（食材样本）输入数组中的第 i 个元素，对应工号牌为 i 的帮厨手里的那份食材。
    xix_i
  2. 线性变换（2x_i + 1）先乘 2 再加 1，对应菜谱里“每份食材先翻倍再调味”的预处理步骤。
    2xi+12x_i + 1
  3. ReLU 截断max(0, ·) 把结果压到非负区间，如同温度计读数低于 0 一律记为 0，这是一类经典激活函数。
    yi=max⁡(0, 2xi+1)y_i = \max(0, \, 2x_i + 1)
wgsl6 行
@compute @workgroup_size(256)fn main(@builtin(global_invocation_id) gid : vec3<u32>) {    let i = gid.x;    if (i >= arrayLength(&input)) { return; }   // 越界保护    output[i] = relu(input[i] * 2.0 + 1.0);     // 自定义算子本体}

3. 缓冲绑定

菜谱写好了，食材怎么递进厨房？帮厨们不能直接从你手里接食材——中间隔着厚厚的隔墙。WebGPU 的答案是：把食材放进托盘（buffer），通过传菜窗口（bind group）递进去。托盘就是 GPU 里的一块内存，CPU 和 GPU 都认识它。

传菜窗口的规格不是随便定的：每个窗口有多少个槽位、每个槽位放什么类型的托盘，要提前画一张“窗口布局图”（bind group layout）。这张图决定了 CPU 侧往窗口的 0 号槽放输入托盘、1 号槽放输出托盘；着色器里则用 @group(0) @binding(0) 认领对应槽位。

数据要来回走：CPU 先把数据写进输入托盘（writeBuffer 或创建时映射），着色器从 0 号窗口读 input、往 1 号窗口写 output，最后把输出托盘拷贝到一个可映射缓冲，再 mapAsync 读回 CPU。整条路像是点餐—出菜—顾客取餐的闭环：

托盘的尺寸必须精确计算，多一分浪费、少一分装不下：

传菜窗口设计得好，后厨就能稳定供菜；设计得乱，帮厨会在窗口前挤成一团。绑定组就是把“哪盘菜放哪个窗口”固定成契约，让 CPU 和 GPU 永远对得上号——这也是 WebGPU 与旧式 WebGL 最本质的区别之一。

byteLength=N×sizeof(f32)=4N\displaystyle \text{byteLength} = N \times \text{sizeof}(\text{f32}) = 4N
公式通俗说明：一份 f32 数据占 4 字节，N 份数据就要准备 4N 字节的托盘空间，这就是 createBuffer 的 size 参数。
  1. N（元素个数）数组里有多少份数据，决定托盘要多长。
    NN
  2. sizeof(f32)（单元素宽度）每个 32 位浮点数占 4 字节，就像每份食材要占一格托盘的规格。
    sizeof(f32)=4\text{sizeof}(\text{f32}) = 4
  3. 总字节数两者相乘得到缓冲区总大小，直接传给 createBuffer 的 size 参数。
    byteLength=N×4=4N\text{byteLength} = N \times 4 = 4N
js10 行
const bufSize = N * 4; // f32 每元素 4 字节const inputBuf = device.createBuffer({ size: bufSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });device.queue.writeBuffer(inputBuf, 0, inputData);const bgLayout = device.createBindGroupLayout({  entries: [{ binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }]});const bindGroup = device.createBindGroup({  layout: bgLayout,  entries: [{ binding: 0, resource: { buffer: inputBuf } }]});

4. 调度组设计

菜和窗口都备好了，主厨开始排班。WebGPU 的排班单位是工作台（workgroup）：每个工作台里有固定数量的帮厨，台数由你派发。一桌菜分几个台、每台几个人，直接决定出菜速度。

帮厨的工号牌由三段组成：global_invocation_id 的 x、y、z 三个分量。对一维数组计算，我们通常只用 x：第几个帮厨就处理第几份数据。而 local_invocation_id 是台内编号，workgroup_id 是台号——三者组合起来，能定位到任何一个帮厨手里的任何一份食材。

现实数据的长度未必恰好是工作台尺寸的整数倍。比如 1000 份食材、每台 256 人，4 台就是 1024 个位子，最后 24 个位子是空的。空位上的帮厨不能瞎忙，必须在菜谱里加一句“越界保护”：if (i >= arrayLength(&input)) { return; }，让多余的人原地待命，而不是去越界写坏别人的菜。

排班是门手艺：台太少浪费 GPU，台太多调度开销变大；尺寸一般取 64~256，并且是 32 的倍数最稳（硬件线程束的天然对齐）。记住主厨的经验：宁可让个别帮厨闲着，也不能让任何一份食材没人处理——边界永远要保护。

workgroup_count=⌈Nworkgroup_size⌉=⌊N+workgroup_size−1workgroup_size⌋\displaystyle \text{workgroup\_count} = \left\lceil \frac{N}{\text{workgroup\_size}} \right\rceil = \left\lfloor \frac{N + \text{workgroup\_size} - 1}{\text{workgroup\_size}} \right\rfloor
公式通俗说明：无论 N 能不能被整除，派发的台数都要“向上凑整”，保证最后一台哪怕只有一个人，也必须把它派出去——宁多勿漏。
  1. 向上取整（ceil）即使最后一台只有一个人，也必须派发一台，确保每份数据都有人处理。
    ⌈Nworkgroup_size⌉\left\lceil \frac{N}{\text{workgroup\_size}} \right\rceil
  2. 补数技巧分子先加 size−1 再向下取整，得到与向上取整等价的结果，更适合整数运算的代码实现。
    ⌊N+workgroup_size−1workgroup_size⌋\left\lfloor \frac{N + \text{workgroup\_size} - 1}{\text{workgroup\_size}} \right\rfloor
  3. 两式等价两种写法数学上完全一样，实际代码里任选其一即可。
    ⌈Ns⌉=⌊N+s−1s⌋\left\lceil \frac{N}{s} \right\rceil = \left\lfloor \frac{N + s - 1}{s} \right\rfloor

5. 流水线构建

菜谱、食材、排班就绪后，还要把整条后厨流水线搭起来。WebGPU 的流水线有固定的七步，少一步菜都出不来：请主厨、发工牌、挂菜谱、设灶台、下单、传菜、上桌。

第一步 requestAdapter 是问浏览器“这台机器有没有 GPU 厨房”。第二步 requestDevice 从厨房拿到一把总钥匙，之后所有操作都通过它发起。第三步把 WGSL 编译成 module，第四步把它和入口函数绑成计算管线——灶台正式点火。

下单环节用 CommandEncoder 记账：beginComputePass 开一单，setPipeline 指定灶台，setBindGroup 指定传菜窗口，dispatchWorkgroups 喊一嗓子“开工”，end 结单。最后把账单交给 queue.submit 提交，GPU 才开始真正干活。整条链路都异步友好：requestDevice 要 await，读回要 await mapAsync——厨房忙碌时，你只管排队等叫号。

这条流水线的骨架，可以用一条链式公式表达：

记住一个画面：浏览器不是直接给 GPU 下命令，而是先写好一张“施工单”（command buffer），凑齐一批再一起提交。就像主厨不会每炒一个菜都跑到食材市场进货，而是把一桌菜单一次性发到后厨——这就是 WebGPU 流水线高效的根本原因。

adapter→requestDevicedevice→createShaderModulemodule→createComputePipelinepipeline\displaystyle \text{adapter} \xrightarrow{\text{requestDevice}} \text{device} \xrightarrow{\text{createShaderModule}} \text{module} \xrightarrow{\text{createComputePipeline}} \text{pipeline}
公式通俗说明：从 GPU 适配器一步步拿到设备、模块、管线，每一环都是上一环的产物，像请主厨、发工牌、挂菜谱、设灶台一样环环相扣。
  1. requestAdapter向浏览器询问“这台机器有没有 GPU 可用”，相当于先找到一位能拍板的主厨。
    requestAdapter→adapter\text{requestAdapter} \rightarrow \text{adapter}
  2. requestDevice从适配器申请一个设备句柄，所有 WebGPU 调用都通过它发起，相当于给主厨发一张长期有效的工牌。
    requestDevice→device\text{requestDevice} \rightarrow \text{device}
  3. 创建模块与管线把 WGSL 菜谱编译成模块，再绑定入口与布局生成计算管线，至此灶台完全就绪。
    createShaderModule→createComputePipeline→pipeline\text{createShaderModule} \rightarrow \text{createComputePipeline} \rightarrow \text{pipeline}
js14 行
const adapter = await navigator.gpu.requestAdapter();const device = await adapter.requestDevice();const module = device.createShaderModule({ code: wgslCode });const pipeline = device.createComputePipeline({  layout: 'auto',  compute: { module, entryPoint: 'main' }});const encoder = device.createCommandEncoder();const pass = encoder.beginComputePass();pass.setPipeline(pipeline);pass.setBindGroup(0, bindGroup);pass.dispatchWorkgroups(Math.ceil(N / 256));pass.end();device.queue.submit([encoder.finish()]);

6. 性能调优

流水线通了，客人却投诉上菜慢。查一查发现：大部分时间根本没花在“炒菜”上，而是花在传菜员来回跑冷库——也就是 CPU 和 GPU 之间的数据搬运与同步。

最常见的坑是“每个中间结果都想读回看一眼”。mapAsync 会让 CPU 阻塞等待 GPU 干完活，一次两次还好，循环里每次都读，等于每道菜都去冷库进货。正确做法是：把多轮计算塞进同一个 command encoder，最后只读一次最终结果。把多轮计算合并成一批，正是“批量合并”的核心思想：

另外两个细节决定天花板：一是内存对齐，WGSL 里结构体成员建议按 16 字节对齐，未对齐的访问会被 GPU 拆成多次读写；二是回读频率，把 N 次 mapAsync 合并成一次，性能往往是数量级提升。厨房优化的本质，就是“让每个帮厨一直有活干、让传菜员少跑路”——你的自定义算子就能从 100 秒跑到 1 秒，快得让点菜的人以为自己看错了。

U=TactiveTmax×100%\displaystyle U = \frac{T_{\text{active}}}{T_{\text{max}}} \times 100\%
公式通俗说明：占用率 U 衡量 GPU 忙不忙——活跃线程占硬件最大线程数的比例。占用率太低说明灶台空着，太高则可能资源拥挤；50%~75% 是实践中的甜点区间。
  1. T_active（活跃线程）当前真正在跑的线程数，对应这一波同时开工的帮厨人数。
    TactiveT_{\text{active}}
  2. T_max（硬件上限）GPU 单次最多容纳的线程数，超出会排队，太少则灶台闲置。
    TmaxT_{\text{max}}
  3. 占用率 U两者比值衡量 GPU 的忙碌程度，是判断 workgroup 尺寸是否合理的风向标。
    U=TactiveTmax×100%U = \frac{T_{\text{active}}}{T_{\text{max}}} \times 100\%

7. 工程落地

会写菜谱、会搭流水线，不等于能开餐厅。真正工程化还要回答两个问题：什么时候开 GPU 厨房？万一客人用的设备没有 GPU 厨房怎么办？

选择标准其实很朴素：数据并行、规模大、访存规则的活儿交给 WebGPU；控制流复杂、强顺序依赖、数据量小的活儿留在 WASM。比如逐元素激活、矩阵乘、归约这类“千篇一律”的计算，正是 GPU 的舒适区；而解析字符串、递归遍历这类“各有各的走法”的逻辑，交给 CPU 反而更快。

落地时用渐进增强：先探测 if (navigator.gpu)，有的设备走 WebGPU，没有就降级到 WASM 甚至普通 JS，保证功能永远可用。调试时多看 Chrome DevTools 的 WebGPU 面板和 WGSL 编译报错，它能精确告诉你菜谱第几行写错了。与框架集成时，ONNX Runtime Web 这类推理引擎开放了自定义 Execution Provider 的接入点，可以把你的算子挂进去。

好的 AI 餐厅从不赌单一路线：GPU 厨房做爆款大菜，WASM 小灶做精细小炒，兼容降级是永远的后备方案。当你把“什么时候用什么厨房”想清楚，自定义算子才真正从 Demo 变成了生产力。

strategy={WebGPUif data-parallel and N is largeWASM / JSotherwise\displaystyle \text{strategy} = \begin{cases} \text{WebGPU} & \text{if data-parallel and } N \text{ is large}\\ \text{WASM / JS} & \text{otherwise} \end{cases}
公式通俗说明：选厨房不是越高级越好，而是“数据并行且规模大”走 GPU 千人厨房，否则回到 WASM 小馆用灵活换省事。
  1. data-parallel（数据并行）算子在每个元素上做相同操作、互不依赖，这时 GPU 的千人后厨优势最大。
    data-parallel\text{data-parallel}
  2. N 足够大数据量大到能摊薄数据搬运与调度开销时，GPU 才划算，否则会被传菜时间拖垮。
    N≫0N \gg 0
  3. 分支决策不满足条件就退回 WASM/JS，像小馆只做少量精菜一样，用最合适的厨房做最合适的菜。
    strategy=WebGPU or WASM / JS\text{strategy} = \text{WebGPU or WASM / JS}

AI 应用场景

场景一：端侧实时视频特效

阶段：推理部署与性能优化

解决的问题：传统的视频滤镜用 Canvas/WebGL 或标准算子拼装，遇到运动模糊、夜间降噪、人像分割后处理等特殊算子，标准库表达不了，只能逐帧丢给 CPU 计算，掉帧严重。自定义算子把“对每个像素做同样的事”的滤镜逻辑写进 WGSL，GPU 一次并行处理整帧百万像素。

具体案例：某在线视频会议产品接入 WebGPU 自定义降噪与美颜算子，把 720p 画面的滤波计算从单帧 32ms 压到 3ms，主线程不再卡顿，低端笔记本也能流畅开美颜。

场景二：浏览器内大语言模型推理

阶段：模型推理优化

解决的问题：LLM 需要 RMSNorm、量化/反量化、KV-cache 管理等特殊算子。这些算子在 Python 里随手可得，但在浏览器里没有对应算子库；用 WASM 逐元素计算太慢。把“归一化 + 量化”融合成一个自定义算子，正好发挥 GPU 数据并行的优势。

具体案例：团队在浏览器运行 7B 量化模型时，把 RMSNorm 与 per-token 量化合并成一个 WGSL 算子，同时把 KV-cache 的更新逻辑用自定义算子并行化，单 token 生成延迟下降约 40%，浏览器端 LLM 应用变得可商用。

场景三：Web 科学计算与物理模拟

阶段：计算加速

解决的问题：N 体引力模拟是 O(N²) 的计算，几千粒子用 JS 双重循环就要数十秒；直方图统计、矩阵归约等数据分析任务同样受 CPU 单核限制。自定义算子把“每一对粒子之间的力”分给不同线程并行计算，让浏览器里的科学计算接近原生性能。

具体案例：某数据可视化工具用 WebGPU 自定义累加算子，把百万级数据点的直方图统计从秒级降到毫秒级，用户拖拽筛选时图表实时响应，体验从“转圈等待”变成“丝滑跟手”。

Python 演示

python40 行
import mathimport time  def custom_op(x):    # 模拟 WGSL 着色器里的自定义算子: y = relu(2x + 1)    return max(0.0, 2.0 * x + 1.0)  def gpu_style_dispatch(data, workgroup_size):    # 模拟 WebGPU 派发模型: 数据按工作台分块, 台内逐元素处理    n = len(data)    workgroup_count = math.ceil(n / workgroup_size)    result = [0.0] * n    for wg in range(workgroup_count):        start = wg * workgroup_size        end = min(start + workgroup_size, n)        for i in range(start, end):            result[i] = custom_op(data[i])    return result  def main():    n = 1_000_000    workgroup_size = 256    data = [float(i) / n * 6.0 - 3.0 for i in range(n)]     t0 = time.perf_counter()    out = gpu_style_dispatch(data, workgroup_size)    t1 = time.perf_counter()     expected = [max(0.0, 2.0 * x + 1.0) for x in data]    assert all(a == b for a, b in zip(out, expected))    print("数据量 %d, 工作台尺寸 %d, 派发 %d 个工作台"          % (n, workgroup_size, math.ceil(n / workgroup_size)))    print("耗时 %.4f 秒, 结果校验通过" % (t1 - t0))  if __name__ == "__main__":    main()
目的：用 Python 模拟 WebGPU 的“按工作台派发 + 越界保护”模型，演示自定义算子的并行逻辑，并校验与朴素实现结果一致。
  1. workgroup_count = math.ceil(n / workgroup_size)计算需要派发多少个工作台，向上取整保证覆盖全部数据。
  2. for wg in range(workgroup_count):逐个工作台派发，模拟 dispatchWorkgroups 的循环。
  3. end = min(start + workgroup_size, n)最后一个工作台的“越界保护”，防止处理超出数组末尾。
  4. result[i] = custom_op(data[i])每个“帮厨”只处理自己工号对应的那份数据，这正是着色器 main 函数做的事。
  5. assert all(a == b for a, b in zip(out, expected))校验并行派发结果与朴素实现完全一致，像 CI 一样守住正确性。

常见误区

误区我听过 WebGPU 都是画三角形、跑 PBR 材质的，它就是个更高级的 WebGL，怎么能用来算 AI 算子呢？

正解WebGPU 是并行计算和图形渲染的统一接口，渲染管线只是它的一半。它还有完整的计算管线（compute pipeline），专门用于通用并行计算——AI 推理、物理模拟、视频处理都在这条管线上跑。实现自定义算子时，你创建的是 compute pipeline 而不是 render pipeline，两者在 WebGPU 中是并列的兄弟。

为何常见：早期 WebGL 只提供图形能力，先入为主的印象让人忽略了 WebGPU 对计算管线的扩展，就像只知道餐厅有小炒窗口，不知道后面还有千人后厨。

误区我把一个 100 元素的循环改写成 WGSL，结果反而更慢了，WebGPU 是不是吹过头了？

正解GPU 的优势来自海量并行，但数据要从 CPU 搬运到 GPU、跑完还要搬回来，再加上调度开销。数据量小时，这些固定开销远超并行收益。经验上，元素数至少上万、且算子本身是数据并行时，GPU 才明显占优。

为何常见：“并行=快”的朴素直觉忽略了数据传输和调度成本——只看到千人后厨的声势，没看到备货送货也要时间。

误区我把 `@workgroup_size` 写成 1024，以为人数越多出菜越快，结果浏览器直接报错。

正解WebGPU 对 workgroup 每维尺寸有硬件上限（通常是 256），超过会校验失败。即使不超标，过大的尺寸也可能导致寄存器压力上升、占用率下降。实践中最常取 64 或 256，并与线程束（32 的倍数）对齐。

为何常见：“人多力量大”的直觉让人忽略了硬件约束与资源竞争——一万个帮厨塞进一个灶台，反而挤得谁也动不了。

误区我在 JS 里直接 `new Float32Array(outputBuffer)`，以为能像读普通数组一样拿到 GPU 算好的数据。

正解WebGPU 缓冲对 CPU 默认不可见，必须通过 mapAsync 建立映射，再用 getMappedRange 读取，读完还要 unmap。而且只有带 MAP_READ 标记（或先拷贝到可映射缓冲）的缓冲才能读。想当然地直接读，拿到的只会是空数据或报错。

为何常见：习惯了 JS 数组的零成本读写，容易忽略 GPU 内存与 CPU 内存是两套物理存储——跨墙取货必然要“按门铃”。

误区我照着 GLSL 的写法把代码抄过来，结果编译报错一堆，WGSL 不就是改个名吗？

正解WGSL 借鉴了 GLSL 的语法外形，但资源绑定模型完全不同：WGSL 用 @group/@binding 显式声明存储类（storage/uniform），内建变量用 @builtin 标注，类型系统也更严格（数组越界、对齐规则都有硬性要求）。照搬 GLSL 就像把中餐菜谱直接拿到西餐厨房执行，火候和工具都对不上。

为何常见：两者语法相似度造成“差不多”的错觉，直到编译器和运行时用报错教你认识差异。

面试题

Q1. WebGPU 的计算着色器语言是？

单选题 · easy

  • A.GLSL
  • B.WGSL← 正确答案
  • C.HLSL
  • D.CUDA

答案：B。WGSL 是 WebGPU 官方指定的着色器语言，专为浏览器设计；GLSL 用于 WebGL/OpenGL，HLSL 是 DirectX 的语言，CUDA 是 NVIDIA 的通用计算框架，后三者都无法直接运行在 WebGPU 上。就像不同的厨房各有各的菜谱符号体系，WebGPU 后厨只认 WGSL。

Q2. WGSL 中计算入口用什么属性标记？

单选题 · easy

  • A.@vertex
  • B.@compute← 正确答案
  • C.@fragment
  • D.@uniform

答案：B。@compute 标记计算着色器入口，@vertex 和 @fragment 属于渲染管线，@uniform 是资源地址空间而非入口标记。类比：一张菜谱只有盖了“计算”章，主厨才把它送进计算厨房而非绘画厨房。

Q3. 处理 N 个元素，每工作台 256 线程，应派发多少个工作台？

单选题 · mid

  • A.按 N 除以 256 向下取整
  • B.按 N 除以 256 向上取整← 正确答案
  • C.按 N 乘以 256 再加一
  • D.按 N 减去 256 直接发

答案：B。派发工作台数 = ceil(N / workgroup_size)，向上取整保证最后一份数据也有人处理；向下取整会丢失尾部数据，乘法和减法更是错误操作。类比：一桌 13 人、每桌坐 4 人，必须开 4 桌而不是 3 桌。

Q4. 把 CPU 数据写入 GPU 缓冲，缓冲需要标记哪种 usage？

单选题 · mid

  • A.GPUBufferUsage.COPY_SRC
  • B.GPUBufferUsage.COPY_DST← 正确答案
  • C.GPUBufferUsage.VERTEX
  • D.GPUBufferUsage.INDEX

答案：B。向 GPU 缓冲写入数据用 COPY_DST（拷贝目标）；COPY_SRC 表示“能作为拷贝源”，是 GPU 结果往外拷贝时用的；VERTEX/INDEX 是顶点/索引缓冲标记，用于渲染。类比：COPY_DST 是“能从 CPU 收货”的托盘，COPY_SRC 是“能把货送出去”的托盘，用途不能混。

Q5. 关于调度，正确的说法有？

多选题 · mid

  • A.工作台尺寸越大性能越好
  • B.全局线程数等于两者乘积← 正确答案
  • C.数据不整除时要做越界保护← 正确答案
  • D.工作台尺寸取质数效果最好

答案：BC。全局线程数 = workgroup_size × workgroup_count，这是并行模型的核心；数据不能整除时必须做越界保护，否则多余线程会越界写坏内存。而“尺寸越大越好”和“取质数效果最好”都是误传，实际取 64/256 且为 32 的倍数最稳。

Q6. 从 GPU 读回计算结果，正确的调用是？

单选题 · mid

  • A.调用缓冲区的异步映射接口← 正确答案
  • B.直接读取缓冲区的内部内存
  • C.调用设备队列的提交接口
  • D.强制主线程休眠等待结果

答案：A。mapAsync 是读取 GPU 缓冲结果的正确姿势，它会建立 CPU 可访问的映射；直接读内存拿不到数据，调用 submit 只是提交执行而非读取，休眠等待纯属碰运气。类比：隔墙取货要先按门铃（mapAsync），不能直接穿墙。

Q7. 减少 CPU 与 GPU 同步开销的做法有？

多选题 · hard

  • A.频繁回读中间结果
  • B.一次编码多个调度← 正确答案
  • C.批量读取最终结果← 正确答案
  • D.每次同步等待完成

答案：BC。减少同步开销的三个方向：一次编码多个调度、批量读取最终结果、避免频繁回读。反过来，频繁回读中间结果和每次同步等待都会让 CPU 阻塞在 GPU 的进度上，把流水线卡死。类比：传菜员把一桌菜单集中送一次，而不是每道菜跑一趟冷库。

Q8. 关于存储缓冲的内存对齐，WGSL 建议怎么做？

单选题 · hard

  • A.一字节对齐可节省空间
  • B.十六字节对齐提升效率← 正确答案
  • C.三十二字节对齐浪费资源
  • D.任意对齐方式均可使用

答案：B。WGSL 存储缓冲的结构体通常按 16 字节对齐，这是 GPU 硬件读取的最优粒度，能避免未对齐访问被拆成多次读写。1 字节对齐虽省空间但访问慢，32 字节浪费，任意对齐更是错误。类比：食材按标准箱规格装箱，搬运队才搬得最快。

Q9. WebGPU 与 WASM 计算对比，正确说法是？

单选题 · mid

  • A.GPU 方案适合数据并行← 正确答案
  • B.CPU 方案适合数据并行
  • C.两种方案性能完全相同
  • D.GPU 方案无法做通用计算

答案：A。WebGPU（GPU 方案）适合数据并行、规模大的计算；WASM 跑在 CPU 上，擅长控制流复杂、强顺序依赖的任务；两者性能不可能完全相同；WebGPU 的计算管线恰恰就是做通用计算的。类比：千人厨房做大锅菜，小灶主厨做精细小炒，各有主场。

Q10. 哪些情况会导致自写算子性能不升反降？

多选题 · extreme

  • A.数据量小且频繁回读← 正确答案
  • B.每帧重复创建管线← 正确答案
  • C.未做内存对齐访问← 正确答案
  • D.开启越界保护检查

答案：ABC。数据量小且频繁回读，搬运开销吃掉并行收益；每帧重复创建管线浪费 GPU 资源；未做内存对齐导致多次访问。而开启越界保护是正确习惯，只会增加一次条件判断，不会明显拖慢性能。类比：厨房低效往往不是炒菜慢，而是进货太勤、灶台重搭、货箱不整。