关于webgl:WebGPU-计算管线计算着色器通用计算入门案例2D-物理模拟

原文译名：WebGPU – 专一于解决外围（GPU Cores），而不是绘图画布（Canvas）

原文公布于 2022 年 3 月 8 日，传送门 https://surma.dev/things/webgpu

这篇货色十分长，不计代码字符也有 1w 字，能比拟好了解 WebGPU 的计算管线中的各个概念，并应用一个简略的 2D 物理模拟程序来了解它，本篇重点是在计算管线和计算着色器，绘图局部应用 Canvas2D 来实现。

WebGPU 是行将推出的 WebAPI，你能够用它拜访图形处理器（GPU），它是一种底层接口。

原作者对图形编程没有多少教训，他是通过钻研 OpenGL 构建游戏引擎的教程来学习 WebGL 的，还在 ShaderToy 上学习 Inigo Quilez 的例子来钻研着色器。因而，他能在 PROXX 中创立背景动画之类的成果，然而他示意对 WebGL 并不太称心。别急，下文马上会解释。

当作者开始留神 WebGPU 后，大多数人通知他 WebGPU 这货色比 WebGL 多很多条条框框。他没思考这些，曾经预感了最坏的状况，他尽可能找了一些教程和标准文档来看，尽管彼时并不是很多，因为他找的时候 WebGPU 还在晚期制订阶段。不过，他深刻之后发现 WebGPU 并没有比 WebGL 多所谓的“条条框框”，反而是像见到了一位老朋友一样相熟。

所以，这篇文章就是来分享学到的货色的。

作者明确指出，他不会在这里介绍如何应用 WebGPU 绘制图形，而是要介绍 WebGPU 如何调用 GPU 进行它自身最原始的计算（译者注：也就是通用计算）。

他感觉曾经有很多材料介绍如何用 WebGPU 进行绘图了，例如 austin 的例子，或者他思考之后也写一些绘图方面的文章。

他在这里会探讨得比拟深刻，心愿读者能正确、无效地应用 WebGPU，然而他不保障你读完就能成为 GPU 性能专家。

絮絮叨叨完结后，筹备发车。

WebGL 是 2011 年公布的，迄今为止，它是惟一能在 Web 拜访 GPU 的底层 API，实际上它是 OpenGL ES 2.0 的繁难封装版以便能在 Web 中应用。WebGL 和 OpenGL 都是科纳斯组标准化的，这个工作组是图形界的 W3C，能够这么了解。

OpenGL 自身是一个颇具历史的 API，按明天的规范看，它不算是一个很好的 API，它以外部全局状态对象为核心。这种设计能够最大限度缩小特定调用的 GPU 的 IO 数据量。然而，这种设计有很多额定的开销老本。

上图：WebGL 外部全局状态对象的可视化，源自 WebGL Fundamentals

外部状态对象，说白了，大多数都是一些指针。调用 OpenGL API 会扭转这些指针的指向，所以扭转状态的程序相当重要，这导致了形象和写库的艰难水平大大增加。你必须十分分明地晓得你当初马上要进行的 API 调用须要筹备什么状态，调用完了还得复原到之前的值。

他说，他常常会看到一个彩色的画布（因为 WebGL 报错大多数时候就这样），而后得狂躁地找没调用哪些 API 没有正确设置全局状态。

他抵赖，他不晓得 ThreeJS 是如何做到状态治理架构的，然而确实做的不错，所以大多数人会应用 ThreeJS 而不是原生 WebGL，这是次要的起因了。

“不能很好认同 WebGL”这只是对原作者他本人说的，而不是读者们。他示意，比他聪慧的人用 WebGL 和 OpenGL 曾经做了不少 nice 的货色，然而他始终不称心罢了。

随着机器学习、神经网络以及加密货币的呈现，GPU 证实了它能够干除了画三角形之外的事件。应用 GPU 进行任意数据的计算，这种被称为 GPGPU，然而 WebGL 1.0 的目标并不在于此。如果你在 WebGL 1.0 想做这件事，你得把数据编码成纹理，而后在着色器中对数据纹理进行解码、计算，而后从新编码成纹理。WebGL 2.0 通过 [转移反馈]() 让这摊子事件更容易了一些，然而直到 2021 年 9 月，Safari 浏览器才反对 WebGL 2.0（大多数浏览器 2017 年 1 月就反对了），所以 WebGL 2.0 不算是好的抉择。

尽管如此，WebGL 2.0 依然没有扭转 WebGL 的实质，就是全局状态。

在 Web 畛域外，新的图形 API 曾经逐步成型。它们向内部裸露了一套拜访显卡的更底层的接口。这些新的 API 改进了 OpenGL 的局促性。

次要就是指 DirectX 12、Vulkan、Metal

一方面来说，当初 GPU 哪里都有，甚至挪动设施都有不错的 GPU 了。所以，古代图形编程（3D 渲染、光追）和 GPGPU 会越来越广泛。

另一方面来看，大多数设施都有多核处理器，如何优化多线程与 GPU 进行交互，是一个重要的课题。

WebGPU 规范制定者留神到了这些现状，在预加载 GPU 之前要做好验证工作，这样能力给 WebGPU 开发者以更多精力专一于压迫 GPU 的性能。

下一代最受欢迎的 GPU API 是：

科纳斯组的 Vulkan
苹果的 Metal
微软的 DirectX 12

为了把这些技术交融并带到 Web，WebGPU 就诞生了。

WebGL 是 OpenGL 的一个浅层封装，然而 WebGPU 并没这么做。它引入了本人的抽象概念体系，吸取上述 GPU API 的长处，而不是继承自这些更底层的 API.

起因很简略，这三个 API 并不是全部都是全平台通用的，而且有一些他们本人的十分底层的概念，对于 Web 这个畛域来说显得不那么正当。

相同，WebGPU 的设计让人感觉“哇，这就是给 Web 设计的”，然而它的的确确又基于你以后机器的 GPU API，形象进去的概念被 W3C 标准化，所有的浏览器都得实现。因为 WebGPU 相对来说比拟底层，它的学习曲线会比拟平缓，然而作者示意会尽可能地合成。

最开始接触到的 WebGPU 抽象概念是适配器（Adapter）和设施（Device）。

上图：形象层，从物理 GPU 到逻辑设备。

物理设施就是 GPU 自身，有内置的 GPU（核芯显卡）和内部 GPU（独立显卡）两种。通常，某个设施个别只有一个 GPU，然而也有两个或者多个的状况。例如，微软的 Surface 笔记本就具备双显卡，以便操作系统在不同的状况进行切换。

操作系统应用显卡厂商提供的驱动程序来拜访 GPU；反过来，操作系统也能够用特定的 API（例如 Vulkan 或者 Metal）向外裸露 GPU 的性能。

GPU 是共享资源，它不仅要被各种程序调用，还要负责向显示器上输入。这看起来须要一个货色来让多个过程同时应用 GPU，以便每个过程把本人的货色画在屏幕上。

对于每个过程来说，仿佛看起里他们对 GPU 有惟一的控制权，然而那只是表象，实际上这些简单逻辑是驱动程序和操作系统来实现调度的。

适配器（Adapter）是特定操作系统的 API 与 WebGPU 之间的中介。

然而，因为浏览器又是一个能够运行多个 Web 程序的“迷你操作系统”，因而，在浏览器层面仍须要共享适配器，以便每个 Web 程序感觉上就像惟一管制 GPU 一样，所以，每个 Web 程序就取得了再次形象的概念：逻辑设备（Logical Device）。

要拜访适配器对象，请调用 navigator.gpu.requestAdapter()，在写本文时，这个办法的参数比拟少，能让你选申请的是高性能的适配器（通常是高性能独显）还是低功耗适配器（通常是核显）。

译者注：本篇探讨 WebGPU 的代码，没非凡指明，均为浏览器端的 WebGPU JavaScript API.

软渲染：一些操作系统（诸如小众 Linux）可能没有 GPU 或者 GPU 的能力有余，会提供“后备适配器（Fallback Adapter）”，实际上这种适配器是纯软件模仿进去的，它可能不是很快，可能是 CPU 模仿进去的，然而能根本满足零碎运作。

若能申请到非空的适配器对象，那么你能够持续异步调用 adapter.requestDevice() 来申请逻辑设备对象。上面是示例代码：

if (!navigator.gpu) throw Error("WebGPU not supported.");

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) throw Error("Couldn’t request WebGPU adapter.");

const device = await adapter.requestDevice();
if (!device) throw Error("Couldn’t request WebGPU logical device.");

如果没有任何申请设施的参数，那么 requestDevice() 会返回一个不匹配任何设施性能要求的设施，即 WebGPU 团队认为是正当且对于所有 GPU 都通用的设施对象。

申请设施对象过程中的“限度”见标准。

举个例子，即便我的 GPU 能够轻易解决 4GB 的数据，返回的设施对象也只容许最大 1GB 的数据，你申请再多也只会返回最大容许 1GB，这样就算你切换到别的机器上跑代码，就不会有太多问题。

你能够拜访 adapter.limits 查看物理 GPU 的理论限度状况。也能够在申请设施对象时，传递你所须要测验的更高限度参数。

如果你用过 WebGL，那么你应该相熟顶点着色器和片元（片段）着色器。其实也没多简单，惯例技术路线就是上载三角形缓冲数据到 GPU，通知 GPU 缓冲数据是如何形成三角形的。顶点缓冲的每个顶点数据形容了顶点的地位，当然还包含色彩、纹理坐标、法线等其它辅助内容。每个顶点都要通过顶点着色器解决，以实现平移、旋转、透视变形等操作。

让原作者感到困惑的是“着色器”这个词，因为它除了着色之外还有别的作用。然而在很久以前（1980 年代前期）来看，这个词十分适合，它在 GPU 上的性能就是计算出像素的色彩值。而现在，它泛指在 GPU 上运行的任何程序。

GPU 会对三角形进行光栅化解决，计算出每个三角形在屏幕上占据的像素。每个像素，则交由片段着色器解决，它能获取像素坐标，当然也能够退出一些辅助数据来决定该像素的最终着色。如果应用切当，就能绘制出令人惊叹的 3D 成果。

将缓冲数据传递到顶点着色器，而后继续传送到片段着色器，最终输入到屏幕上这一过程，能够简略的称之为管道（或管线，Pipeline），在 WebGPU 中，必须明确定义 Pipeline.

目前，WebGPU 反对两大管线：

渲染管线
计算管线

顾名思义，渲染管线绘制某些货色，它后果是 2D 图像，这个图像不肯定要绘制到屏幕上，能够间接渲染到内存中（被称作帧缓冲）。计算管线则更加通用，它返回的是一个缓冲数据对象，意味着能够输入任意数据。

在本文的其它局部会专一于计算管线的介绍，因为作者认为渲染管线算是计算管线的一种非凡状况。

当初开始算开历史倒车，计算管线原来其实是为了创立渲染管线而先做进去的“根底”，这些所谓的管线在 GPU 中其实就是不同的物理电路罢了。

基于上述了解，假使将来向 WebGPU 中增加更多类型的管线，例如“光追管线”，就显得天经地义了。

应用 WebGPU API，管线由一个或多个可编程阶段组成，每个阶段由一个着色器模块和一个入口函数定义。计算管线领有一个计算着色阶段，渲染管线有一个顶点着色阶段和一个片段着色阶段，如下所示是一个计算着色模块与计算管线：

const module = device.createShaderModule({
  code: `
    @stage(compute) @workgroup_size(64)
    fn main() {// ...}
  `,
})

const pipeline = device.createComputePipeline({
  compute: {
    module,
    entryPoint: "main",
  },
})

这是 WebGPU 的着色语言（WGSL，发音 /wig-sal/）的首次退场。

WGSL 给作者的初印象是 Rust + GLSL，它有很多相似 Rust 的语法，也有相似 GLSL 一样的全局函数（如 dot()、norm()、len() 等），以及类型（vec2、mat4x4 等），还有 swizzling 语法（例如 some_vec.xxy）。

浏览器会把 WGSL 源码编译成底层零碎的着色器目标程序，可能是 D3D12 的 HLSL，也可能是 Metal 的 MSL，或者 Vulkan 的 SPIR-V.

SPIR-V：是科纳斯组标准化进去的开源、二进制两头格局。你能够把它看作并行编程语言中的 LLVM，它反对多种语言编译成它本人，也反对把本人翻译到其它语言。

在下面的着色器代码中，只创立了一个 main 函数，并应用 @stage(compute) 这个个性（Attribute，WGSL 术语）将其标记为计算着色阶段的入口函数。

你能够在着色器代码中标记多个 @stage(compute)，这样就能够在多个管线中复用一个着色器模块对象了，只需传递不同的 entryPoint 抉择不同的入口函数即可。

然而，@workgroup_size(64) 个性是什么？

GPU 以提早为代价优化了数据吞吐量。想深刻这点必须看一下 GPU 的架构，然而作者没信念讲好这块，所以倡议看一看 Fabian Giesen 的文章。

家喻户晓，GPU 有十分多个外围形成，能够进行大规模的并行运算。然而，这些外围不像 CPU 并行编程一样绝对独立运作。首先，GPU 解决外围是分层分组的，不同厂商的 GPU 的设计架构、API 不尽一致。Intel 这里给了一个不错的文档，对他们的架构进行了高级的形容。

在 Intel 的技术中，最小单元被称作“执行单元（Execution Unit，EU）”，每个 EU 领有 7 个 SIMT 内核 —— 意思是，它有 7 个以“锁步”（Lock-step）的形式运行同一个指令的并行计算核。每个内核都有本人的寄存器和调度缓存的指针，只管执行着雷同的操作，然而数据能够是不同的。

所以有时候不举荐在 GPU 上执行 if/else 判断分支，是因为 EU 的起因。因为 EU 遇到分支逻辑的时候，每个内核都要进行 if/else 判断，这就失去了并行计算的劣势了。

对于循环也是如此。如果某个外围提前完成了计算工作，那它不得不伪装还在运行，期待 EU 内其它外围实现计算。

只管内核的计算频率很高，然而从内存中加载数据或者从纹理中采样像素的工夫显著要更长 —— Fabian 同志说，这起码要消耗几百个时钟周期。这些工夫显然能够拿来算货色。为了充分利用这些时钟周期，每个 EU 必须负重前行。

EU 闲暇的时候，譬如在等内存的食物过去的时候，它可不会就始终闲上来，它会立马投入到下一个计算中，只有这下一个计算再次进入期待时，才会切换回来，切换的过程十分十分短。

GPU 就是以这样的技术为代价换来吞吐量的优化的。GPU 通过调度这些工作的切换机制，让 EU 始终处于繁忙状态。

上图：Intel 锐炬 Xe 显卡芯片架构。它被分成 8 个子块，每个子块有 8 个 EU；每个 EU 领有 7 个 SIMT 内核。

不过，依据上图来看，EU 只是 Intel 显卡设计架构层级最低的一个，多个 EU 被 Intel 分为所谓的“子块（SubSlice）”，子块中所有的 EU 都能够拜访共有的部分缓存（Shared Local Memory，SLM），大略是 64KB，如果所运行的程序有同步指令，那么就必须在同一个子块中运行，因为这样能力共享内存。

再往上，子块就形成了块（Slice），形成 GPU；对于集成在 CPU 中的 GPU，大概有 170 ~ 700 个内核。对于独立显卡，则会有 1500 或以上个内核。

其它厂商兴许会用其它的术语，然而架构基本上能够这么类比了解。

为了充分利用 GPU 的架构劣势，须要专门写程序调用，这样就能够最大限度地压迫 GPU 的性能。所以，图形 API 得向外裸露相似的线程模型来调用计算工作。

在 WebGPU API 中，这种线程模型就叫做“工作组（Workgroup）”。

每个顶点都会被顶点着色器解决一次，每个片元则会被片元着色器解决一次（当然，这是简略说法，疏忽了很多细节）。

而在 GPGPU 中，与顶点、片元相似的概念是须要开发者本人定义的，这个概念叫做 计算项，计算项会被计算着色器解决。

一组计算项就形成了“工作组”，作者称之为“工作负载”。工作组中的每个计算项会被同时运行的计算着色器作用。在 WebGPU 中，工作组能够设想成一个三维网格，最小层级的是计算项，计算项形成稍大级别的是工作组，再往上就形成规模更大的工作负载。

上图：这是一个工作负载，其中红色小立方体由 4³ 个红色小立方体形成，红色小立方是计算项，而红色小立方体则由这 64 个红色小立方形成，即工作组。

基于上述概念，就能够探讨 WGSL 中的 @workgroup_size(x, y, z) 个性了，它的作用很简略，就是通知 GPU 这个计算着色器作用的工作组有多大。用下面的图来说，其实就是红色小立方的大小。x*y*z 是每个工作组的计算项个数，如果不设某个维度的值，那默认是 1，因而，@workgroup_size(64) 等同于 @workgroup_size(64, 1, 1).

当然，理论 EU 的架构当然不会是这个 3D 网格外面的某个单元。应用这个图来形容计算项的目标是凸显出一种部分性质，即假如相邻的工作组大概率会拜访缓存中类似的区域，所以依次运行相邻的工作组（红色小立方）时，命中缓存中已有的数据的几率会更高一些，而无需在再跑去显存要数据，节俭了十分多工夫周期。

然而，大多数硬件仍旧是程序执行工作组的，所以设置 @workgroup_size(64) 和 @workgroup_size(8, 8) 的两个不同的着色器实际上差别并不是很大。所以，这个设计上略显冗余。

工作组并不是有限维度的，它受设施对象的限度条件束缚，打印 device.limits 能够获取相干的信息：

console.log(device.limits)

/*
{
  // ...
  maxComputeInvocationsPerWorkgroup: 256,
  maxComputeWorkgroupSizeX: 256,
  maxComputeWorkgroupSizeY: 256,
  maxComputeWorkgroupSizeZ: 64,
  maxComputeWorkgroupsPerDimension: 65535,
  // ...
}
*/

能够看到，每个维度上都有最大限度，而且累乘的积也有最大限度。

提醒：防止申请每个维度最大限度数量的线程。尽管 GPU 由操作系统底层调度，但如果你的 WebGPU 程序霸占了 GPU 太久的话，零碎有可能会卡死。

那么，适合的工作组大小倡议是多少呢？这须要具体问题具体分析，取决于工作组各个维度有什么指代含意。作者认为这答案很含混，所以他援用了 Corentin 的话：“用 64 作为工作组的大小（各个维度累乘后），除非你非常分明你须要调用 GPU 干什么事件。”

64 像是个比拟稳当的线程数，在大多数 GPU 上跑得还能够，而且能让 EU 尽可能跑满。

到目前为止，曾经写好了着色器并设置好了管线，剩下的就是要调用 GPU 来执行。因为 GPU 能够是有本人内存的独立显卡，所以能够通过所谓的“指令缓冲”或者“指令队列”来管制它。

指令队列，是一块内存（显示内存），编码了 GPU 待执行的指令。编码与 GPU 自身严密相干，由显卡驱动负责创立。WebGPU 裸露了一个“CommandEncoder”API 来对接这个术语。

const commandEncoder = device.createCommandEncoder()
const passEncoder = commandEncoder.beginComputePass()
passEncoder.setPipeline(pipeline)
passEncoder.dispatch(1)
passEncoder.end()
const commands = commandEncoder.finish()
device.queue.submit([commands])

commandEncoder 对象有很多办法，能够让你把某一块显存复制到另一块，或者操作纹理对应的显存。它还能够创立 PassEncoder（通道编码器），它能够配置管线并调度编码指令。

在上述例子中，展现的是计算管线，所以创立的是计算通道编码器。调用 setPipeline() 设置管线，而后调用 dispatch() 办法通知 GPU 在每个维度要创立多少个工作组，以备进行计算。

换句话说，计算着色器的调用次数等于每个维度的大小与该维度调用次数的累积。

例如，一个工作组的三个维度大小是 2, 4, 1，在三个维度上要运行 4, 2, 2 次，那么计算着色器一共要运行 2×4 + 4×2 + 1×2 = 18 次。

顺便说一下，通道编码器是 WebGPU 的抽象概念，它就是文章最开始时作者埋怨 WebGL 全局状态机的良好替代品。运行 GPU 管线所需的所有数据、状态都要通过通道编码器来传递。

形象：指令缓冲也只不过是显卡驱动或者操作系统的钩子，它能让程序调用 GPU 时不会互相烦扰，确保互相独立。指令推入指令队列的过程，其实就是把程序的状态保留下来以便待会要用的时候再取出，因为硬件执行的速度十分快，看起来就是各做各的，没有受到其它程序的烦扰。

跑起代码，因为 workgroup_size 个性显式指定了 64 个工作组，且在这个维度上调用了 1 次，所以最终生成了 64 个线程，尽管这个管线啥事儿都没做（因为没写代码），然而至多起作用了，是不是很酷炫？

随后，咱们搞点数据来让它起作用。

如文章结尾所言，作者没打算间接用 WebGPU 做图形绘制，而是打算拿它来做物理模仿，并用 Canvas2D 来简略的可视化。尽管叫是叫物理模仿，实际上就是生成一堆圆几何，让它们在立体范畴内随机静止并模仿他们之间互相碰撞的过程。

为此，要把一些模仿参数和初始状态传递到 GPU 中，而后跑计算管线，最初读取后果。

这能够说是 WebGPU 最头皮发麻的的一部分，因为有一堆的数据术语和操作要学。不过作者认为恰好是这些数据概念和数据的行为模式造就了 WebGPU，使它成为了高性能的且与设施无关的 API.

为了与 GPU 进行数据交换，须要一个叫绑定组的布局对象（类型是 GPUBindGroupLayout）来裁减管线的定义。

首先要说说绑定组（类型是 GPUBindGroup），它是某种管线在 GPU 执行时各个资源的几何，资源即 Buffer、Texture、Sampler 三种。

而先于绑定组定义的绑定组布局对象，则记录了这些资源的数据类型、用处等元数据，使得 GPU 能够提前晓得“噢，这么回事，提前通知我我能够跑得更快”。

下列创立一个绑定组布局，简略起见，只设置一个存储型（type: "storage"）的缓冲资源：

const bindGroupLayout = device.createBindGroupLayout({
  entries: [{
    binding: 1,
    visibility: GPUShaderStage.COMPUTE,
    buffer: {type: "storage",}
  }]
})

// 紧接着，传递给管线
const pipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module,
    entryPoint: 'main'
  }
})

binding 这里设为了 1，能够自在设定（当然得按程序），它的作用是在 WGSL 代码中与雷同 binding 值的 buffer 变量绑定在一起。

@group(0) @binding(1)
var<storage, write> output: array<f32>;

type 字段是 "storage"，即阐明这个 Buffer 的类型是存储型。它还能够设置为其它的选项，其中 "read-only-storage" 即“只读存储型”，即着色器只能读，然而不能写这个 Buffer，只读型缓冲能够优化一些读写同步的问题；而 "uniform" 则阐明 Buffer 类型是对立数据（Uniform），作用和存储型差不多（在着色器中值都一样）。

至此，绑定组布局对象创立结束，而后就能够创立绑定组了，这里就不写进去了；一旦创立好了对应的绑定组和存储型 Buffer，那么 GPU 就能够开始读取数据了。

然而，在此之前，还有一个问题要探讨：暂存缓冲区。

这个大节的内容略长，请急躁浏览。

作者再次强调：GPU 以提早为代价，高度优化了数据 IO 性能。GPU 须要相当快的速度向内核提供数据。在 Fabian 他 2011 年的博客中做了一些计算，得出的论断是 GPU 须要维持 3.3 GB/s 的速度能力运行 1280×720 分辨率的纹理的采样计算。

为了满足当初的图形需要，GPU 还要再快。只有 GPU 的内核与缓冲存储器高度集成能力实现，这意味着也就难以把这些存储区交由 CPU 来读写。

咱们都晓得 GPU 有本人的内存，叫显存，CPU 和 GPU 都能够拜访它，它与 GPU 的集成度不高，个别在电路板的旁边，它的速度就没那么快了。

暂存缓冲区（Staging buffers），是介于显存和 GPU 之间的缓存，它能够映射到 CPU 端进行读写。为了读取 GPU 中的数据，要先把数据从 GPU 内的高速缓存先复制到暂存缓冲区，而后把暂存缓冲区映射到 CPU，这样能力读取回主内存。对于数据传递至 GPU 的过程则相似。

回到代码中，创立一个可写的 Buffer，并增加到绑定组，以便计算着色器能够写入它；同时还创立一个大小一样的 Buffer 以作为暂存。创立这些 Buffer 的时候，要用位掩码来告知其用处（usage），GPU 会依据参数申请、创立这些缓冲区，如果不合乎 WebGPU 规定，则抛出谬误：

const BUFFER_SIZE = 1000
const output = device.createBuffer({
  size: BUFFER_SIZE,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
})
const stagingBuffer = device.createBuffer({
  size: BUFFER_SIZE,
  usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
})

const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [{
    binding: 1,
    resource: {buffer: output}
  }]
})

留神，createBuffer() 返回的是 GPUBuffer 对象，不是 ArrayBuffer，创立完 Buffer 后还不能马上写入或者读取。为了实现读写 Buffer，须要有独自的 API 调用，而且 Buffer 必须有 GPUBufferUsage.MAP_READ 或 GPUBufferUsage.MAP_WRITE 的用处能力读或写。

TypeScript 提醒：在各开发环境还未退出 WebGPU API 时，想要取得 TypeScript 类型提醒，还须要装置 Chrome WebGPU 团队保护的 @webgpu/types 包到你的我的项目中。

到目前为止，不仅有绑定组的布局对象，还有绑定组自身，当初须要批改通道编码器局部的代码以应用这个绑定组，随后还要把 Buffer 中计算好的数据再读取回 JavaScript：

const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(pipeline)
passEncoder.setBindGroup(0, bindGroup)
passEncoder.dispatch(1)
passEncoder.dispatch(Math.ceil(BUFFER_SIZE / 64))
passEncoder.end()
commandEncoder.copyBufferToBuffer(
  output,
  0, // 从哪里开始读取
  stagingBuffer,
  0, // 从哪里开始写
  BUFFER_SIZE
)
const commands = commandEncoder.finish()
device.queue.submit([commands])

await stagingBuffer.mapAsync(
  GPUMapMode.READ,
  0, // 从哪里开始读，偏移量
  BUFFER_SIZE // 读多长
 )
const copyArrayBuffer = stagingBuffer.getMappedRange(0, BUFFER_SIZE)
const data = copyArrayBuffer.slice()
stagingBuffer.unmap()
console.log(new Float32Array(data))

稍前的代码中，管线对象借助管线布局增加了绑定组的局对象，所以如果在通道编码的时候不设置绑定组就会引起调用（dispatch）失败。

在计算通道 end() 后，指令编码器紧接着触发一个缓冲拷贝办法调用，将数据从 output 缓冲复制到 stagingBuffer 缓冲，最初才提交指令编码的指令缓冲到队列上。

GPU 会沿着队列来执行，没法揣测什么时候会实现计算。然而，能够异步地提交 stagingBuffer 缓冲的映射申请；当 mapAsync 被 resolve 时，stagingBuffer 映射胜利，然而 JavaScript 仍未读取，此时再调用 stagingBuffer.getMappedRange() 办法，就能获取对应所需的数据块了，返回一个 ArrayBuffer 给 JavaScript，这个返回的缓冲数组对象就是显存的映射，这意味着如果 stagingBuffer 的状态是未映射时，返回的 ArrayBuffer 也随之没有了，所以用 slice() 办法来拷贝一份。

显然，能够在控制台看到输入成果：

上图：对付，然而阐明了一个问题，那就是从 GPU 显存中把这堆 0 给拿下来了

或者，制作点 0 之外的数据会更有说服力。在进行高级计算之前，先搞点人工数据到 Buffer 中，以证实计算管线的确按预期在运行：

@group(0) @binding(1)
var<storage, write> output: array<f32>;

@stage(compute) @workgroup_size(64)
fn main(@builtin(global_invocation_id)
  global_id : vec3<u32>,

  @builtin(local_invocation_id)
  local_id : vec3<u32>,

) {output[global_id.x] =
    f32(global_id.x) * 1000. + f32(local_id.x);
}

前两行申明了一个名为 output 的模块范畴的变量，它是一个 f32 元素类型的数组。它的两个个性申明了起源，@group(0) 示意从第一个（索引为 0）绑定组中获取第 1 个绑定资源。output 数组是动静长度的，会主动反射对应 Buffer 的长度。

WGSL 变量：与 Rust 不同，let 申明的变量是不可变的，如果心愿变量可变，应用 var 申明

接下来看 main 函数。它的函数签名有两个参数 global_id 和 local_id，当然这两个变量的名称随你设定，它们的值取决于对应的内置变量 global_invocation_id、local_invocation_id，别离指的是 工作负载 中此着色器调用时的全局 x/y/z 坐标，以及 工作组 中此着色器调用时的部分 x/y/z 坐标。

上图：三个计算项，a、b、c，用绿色字母标注。

这张图中应用的工作组大小是 @workgroup_size(4, 4, 4)，应用图中的坐标轴程序，那么对于图中的 a、b、c 计算项：

a：local_id = (x=0, y=0, z=0)，global_id = (x=0, y=0, z=0)
b：local_id = (x=0, y=0, z=0)，global_id = (x=4, y=0, z=0)
c：local_id = (x=1, y=1, z=0)，global_id = (x=5, y=5, z=0)

而对于咱们的例子来说，工作组的大小被设为 @workgroup_size(64, 1, 1)，所以 local_id.x 的取值范畴是 0 ~ 63. 为了能查看 local_id 和 global_id，作者把这两个值进行编码，合成一个数字；留神，WGSL 类型是严格的，local_id 和 global_id 都是 vec3<u32>，因而要显式地转换为 f32 类型能力写入 output 缓冲区。

上图：GPU 写入的理论值，留神 local_id 是 63 为循环的起点，而 global_id 则仍旧在持续编码

上图证实了计算着色器的确向缓冲区输入了值，然而很容易发现这些数字看似是没什么程序的，因为这是成心留给 GPU 去做的。

你可能会留神到，计算通道编码器的调度办法调度次数 Math.ceil(BUFFER_SIZE / 64) * 64 这个值，算进去就是 1024：

passEncoder.dispatch(Math.ceil(BUFFER_SIZE / 64))

这间接导致着色器代码中 global_id.x 的取值能取到 1024，大于 Buffer 的长度 1000.

不过还好，WGSL 是有爱护超出数组索引范畴的机制的，即一旦产生对数组索引越界的写入，那么总是会写入最初一个元素。这样尽管能够防止内存拜访谬误，然而仍有可能会生成一些有效数据。譬如，你把 JavaScript 端返回的 Float32Array 的最初 3 个元素打印进去，它们是 247055、248056、608032；如何防止因数组索引越界而可能产生的有效数据问题呢？能够用卫语句提前返回：

fn main(/* ... */) {if (global_id.x >= arrayLength(&output)) {return;}
  
  output[global_id.x] = f32(global_id.x) * 100. + f32(local_id.x)
}

若读者感兴趣，能够运行这个例子看成果。

还记得指标吗？是在 2D 的 Canvas 中挪动一些圆，并让他们激情地碰撞。

所以，每个圆都要有一个半径参数和一个坐标参数，以及一个速度矢量。能够持续用 array<f32> 来示意上述数据，例如第一个数字是 x 坐标，第二个数字是 y 坐标，以此类推。

然而，这看起来有点蠢，WGSL 是容许自定义构造体的，把多条数据关联在一个构造内。

留神：如果你晓得什么是内存对齐，你能够跳过本大节；如果你不晓得，作者也没打算认真解释，他打算间接展现为什么要这么做。

因而，定义一个构造体 Ball，示意 2D 中的圆，并应用 array<Ball> 示意一系列的 2D 圆球。

应用构造体，就不得不探讨内存对齐问题。

struct Ball {
  radius: f32;
  position: vec2<f32>;
  velocity: vec2<f32>;
}

@group(0) @binding(1)
var<storage, write> output: array<Ball>;

@stage(compute) @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>,
  @builtin(local_invocation_id) local_id: vec3<u32>,
) {let num_balls = arrayLength(&output);
  if (global_id.x >= num_balls) {return;}
  
  output[global_id.x].radius = 999.;
  output[global_id.x].position = vec2<f32>(global_id.xy);
  output[global_id.x].velocity = vec2<f32>(local_id.xy);
}

你能够运行这个代码，关上控制台能够看到：

上图：因为内存对齐的起因，这个 TypedArray 有显著的数据填充景象

着色器代码首先把数据 999.0 写入到构造体的第一个字段 radius 中，以便于察看两个构造的分隔界线；然而，这个打印的 Float32Array 中，两个 999 数字之间，实际上逾越了 6 个数字，譬如上图中 0~5 位数字是 999, 0, 0, 0, 0, 0，紧随其后的 6~11 位数字是 999, 0, 1, 0, 1, 0，这就意味着每个构造体都占据了 6 个数字，然而 Ball 构造体明明只须要 5 个数字即可存储：radius、position.x、position.y、velocity.x 和 velocity.y. 很显著，每个 radius 前面都塞多了一个 0，这是为什么呢？

起因就是内存对齐。每一种 WGSL 中的数据类型都要严格执行对齐要求。

若一个数据数据类型的对齐尺度是 N（字节），则意味着这个类型的数据值只能存储在 N 的倍数的内存地址上。举个例子，f32 的对齐尺度是 4（即 N = 4），vec2<f32> 的对齐尺度是 8（即 N = 8）.

假如 Ball 构造的内存地址是从 0 开始的，那么 radius 的存储地址能够是 0，因为 0 是 4 的倍数；紧接着，下个字段 position 是 vec2<f32> 类型的，对齐尺度是 8，问题就呈现了 —— 它的前一个字段 radius 闲暇地址是第 4 个字节，并非 position 对齐尺度 8 的倍数，为了对齐，编译器在 radius 前面增加了 4 个字节，也就是从第 8 个字节开始才记录 position 字段的值。这也就阐明了控制台中看到 999 之后的数字为什么总是 0 的起因了。

当初，晓得构造体在内存中是如何散布字节数据的了，能够在 JavaScript 中进行下一步操作了。

咱们曾经从 GPU 中读取到数据了，当初要在 JavaScript 中解码它，也就是生成所有 2D 圆的初始状态，而后再次提交给 GPU 运行计算着色器，让它“动起来”。初始化很简略：

let inputBalls = new Float32Array(new ArrayBuffer(BUFFER_SIZE))
for (let i = 0; i < NUM_BALLS; i++) {inputBalls[i * 6 + 0] = randomBetween(2, 10) // 半径
  inputBalls[i * 6 + 1] = 0 // 填充用
  inputBalls[i * 6 + 2] = randomBetween(0, ctx.canvas.width) // x 坐标
  inputBalls[i * 6 + 3] = randomBetween(0, ctx.canvas.height) // y 坐标
  inputBalls[i * 6 + 4] = randomBetween(-100, 100) // x 方向速度重量
  inputBalls[i * 6 + 5] = randomBetween(-100, 100) // y 方向速度重量
}

小技巧：如果你当前的程序用到了更简单的数据结构，应用 JavaScript 拼凑这些字节码会十分麻烦，你能够用 Google 的 buffer-backed-object 库去创立简单的二进制数据（相似序列化）。

还记得如何把 Buffer 传递给着色器吗？不记得的回去看看上文。只须要调整一下计算管线的绑定组布局即可接管新的 Buffer：

const bindGroupLayout = device.createBindGroupLayout({
  entries: [
    {
      binding: 0,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {type: 'read-only-storage'}
    },
    {
      binding: 1,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {type: 'storage'}
    }
  ]
})

而后创立一个新的绑定组来传递初始化后的 2D 圆球数据：

const input = device.createBuffer({
  size: BUFFER_SIZE,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
})

const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [
    {
      binding: 0,
      resource: {buffer: input // 输出初始化数据}
    },
    {
      binding: 1,
      resource: {buffer: output}
    }
  ]
})

就像读取数据一样，从技术角度来看，为了输出初始化的 2D 圆球数据，要创立一个可映射的暂存缓冲区 input，作为着色器读取数据的容器。

WebGPU 提供了一个简略的 API 便于咱们把数据写进 input 缓冲区：

device.queue.writeBuffer(input, 0, inputBalls)

就是这么简略，并不需要指令编码器 —— 也就是说不须要借助指令缓冲，writeBuffer() 是作用在队列上的。

device.queue 对象还提供了一些不便操作纹理的 API.

当初，在着色器代码中要用新的变量来与这个新的 input 缓冲资源绑定：

// ... Ball 构造体定义 ...

@group(0) @binding(0)
var<storage, read> input: array<Ball>;

// ... output Buffer 的定义

let TIME_STEP: f32 = 0.016;

@stage(compute) @workgroup_size(64)
fn main(@builtin(global_invocation_id)
  global_id: vec3<u32>
) {let num_balls = arrayLength(&output);
  if (global_id.x >= num_balls) {return;}
  
  // 更新地位
  output[global_id.x].position = 
    input[global_id.x].position +
    input[global_id.x].velocity * TIME_STEP;
}

心愿大部分着色器代码你能看得懂。

最初要做的，只是把 output 缓冲再次读取回 JavaScript，写一些 Canvas2D 的可视化代码把 Ball 的静止成果展现进去（须要用到 requestAnimationFrame），你能够看示例成果：demo

3.5 大节最初演示的代码只是能让 Ball 静止起来，还没有特地简单的计算。在进行性能观测之前，要在着色器中加一些适当的物理计算。

作者就不打算解释物理计算了，写到这里，博客曾经很长了，然而他简略的阐明了物理成果的外围原理：每个 Ball 都与其它的 Ball 进行碰撞检测计算。

如果你非常想晓得，能够看看最终的演示代码：final-demo，在 WGSL 代码中你还能够找到物理计算的材料连贯。

作者并未优化物理碰撞算法，也没有优化 WebGPU 代码，即便是这样，在他的 MacBook Air（M1 处理器）上体现得也很不错。

当超过 2500 个 Ball 时，帧数才掉到 60 帧以下，然而应用 Chrome 开发者工具去观测性能信息时，掉帧并不是 WebGPU 的问题，而是 Canvas2D 的绘制性能有余 —— 应用 WebGL 或 WebGPU 绘图就不会呈现这个问题了。

上图：即便是 14000 个 Ball，WebGPU 在 M1 处理器的 MBA 笔记本上也才用了 16 毫秒的单帧计算工夫

作者敞开了 Canvas2D 绘图，退出 performance.measure() 办法来查看 16 毫秒之内到底能够模仿多少个 Ball 的物理计算。

这性能体现还是没有优化过的，曾经让作者为之沉醉。

WebGPU 曾经开发了蛮久了，作者认为制订标准的人心愿 API 是稳固的。

话是这么说没错，然而 WebGPU API 目前只能跑在 Chrome 类浏览器和 FireFox 浏览器上，对 Safari 放弃乐观态度 —— 尽管写本文时，Safari TP（技术预览）还没什么货色能看。

在稳定性体现上，即便是写文章的这段时间里，也是有变动的。

例如，WGSL 着色器代码的个性语法，从单方括号改为 @ 符号：

[[stage(compute), workgroup_size(64)]]
↓
@stage(compute) @workgroup_size(64)

对通道编码器完结的办法，Firefox 浏览器依然是 endPass()，而 Chrome 类浏览器曾经改为最新的 end().

标准中还有一些内容也并不是齐全实现在所有浏览器上的，用于挪动设施的 API 以及局部着色器常量就是如此。

基本上，WebGPU 进入 stable 阶段后，不排除会产生很多重大变动。

“在 Web 上能间接应用 GPU”这种古代的 API 看起来很好玩。在经验过最后的平缓学习曲线后，作者认为真的能够应用 JavaScript 调用 GPU 进行大规模并行运算了。

wgpu 是应用 Rust 实现的 WebGPU，你能够在浏览器之外应用 Rust 语言调用 WebGPU 标准的 API；wgpu 还反对编译到 WebAssembly，你甚至能够应用 Rust 的 wgpu 编写 wasm，而后再放到浏览器运行高性能的代码。

还有个乏味的货色：Deno 借助 wgpu，内置了 WebGPU 的反对。

如果你有啥问题，你能够去 WebGPU Matrix 频道（国内可能拜访不太通顺）发问，那里有一些 WebGPU 的用户、浏览器工程师和制订标准的人。

感激 Brandon Jones 校对本文，感激 WebGPU Matrix 频道解惑。

也感激原作者分享这篇长文。

关于webgl:WebGPU-计算管线计算着色器通用计算入门案例2D-物理模拟

1. WebGL

2. WebGPU

2.1. 适配器（Adapter）和设施（Device）

2.2. 着色器（Shaders）

2.3. 管线（Pipeline）

2.4. 并行（Parallelism）

2.5. 工作组（Workgroup）

2.6. 指令（Command）

3. 数据交换

3.1. 绑定组的布局（GPUBindGroupLayout）

3.2. 暂存缓冲区（Staging Buffer）

3.3. 适度调度

3.4. 麻烦的构造体（内存地址对齐问题）

3.5. 输入输出

4. 性能

5. 稳定性与可用性

总结