19. WebGPU —计算着色器(compute shader)

WebGPU 是即将推出的 Web API，可提供对GPU的底层控制并用于通用目的计算任务 。

我对计算机图形不是很有经验。我通过阅读有关如何使用 OpenGL 构建游戏引擎的教程了解了 WebGL 的点点滴滴，并通过观看 Inigo Quilez 在 ShaderToy 上仅使用着色器而不使用任何 3D 网格或模型来完成令人惊叹的事情，从而了解了有关着色器的更多信息。这让我足以在 PROXX 中构建诸如背景动画之类的东西，但我对 WebGL 一直不满意，我将很快解释原因。

当 WebGPU 进入我的视野时，我想进一步了解它，但很多人警告我 WebGPU 的样板文件比 WebGL 还要多。我没有被吓倒，但预料到了最坏的情况，我将能找到的所有教程和规范拼凑在一起，其中并不多，因为 WebGPU 还处于早期阶段。我试了一下，我并没有发现 WebGPU 比 WebGL 有更多的样板文件，但实际上作为一个 API，我更容易适应。

所以我们在这里相遇。我想分享我在研究 GPU 和 WebGPU 时学到的东西。这篇博文的目标是让 Web 开发人员可以访问 WebGPU。但提前提醒一下：我不会使用 WebGPU 生成图形。相反，我将使用 WebGPU 访问 GPU 提供的原始计算能力。也许我会写一篇后续博客文章如何使用 WebGPU 渲染到您的屏幕，但那里已经有相当多的内容了。我将尽可能深入地了解 WebGPU，并希望能让您有效地使用它——但不一定有效。我不能让你成为 GPU 性能专家；主要是因为我自己也不是一个 GPU 性能专家。

足够的免责声明。这是一个很长的文章。系好安全带！

1. WebGL

WebGL 发布于 2011 年，到目前为止，它是唯一可以从web访问 GPU 的底层 API。 WebGL 的 API 实际上只是 OpenGL ES 2.0，做了一些简单封装和工具程序以使其与 Web 兼容。 WebGL 和 OpenGL 都是由 Khronos Group 标准化的，它基本上是 3D 图形的 W3C。

OpenGL 的 API 本身可以追溯到更早的时候，按照今天的标准，它并不是一个很好的 API。该设计以内部全局状态对象为中心。从这个角度来看，这种设计是有意义的，因为它最大限度地减少了任何给定调用需要传输到 GPU 和从 GPU 传输的数据量。然而，它也引入了大量的精神开销。

内部状态对象基本上是指针的集合。对 API 调用会影响状态对象指向的对象，也会影响状态对象本身。因此，API 调用的顺序非常重要，我总觉得这使得构建抽象和库变得困难。必须非常仔细地清理所有可能干扰将要进行的 API 调用的指针和状态项，同时还要将指针和值恢复到它们调用前的值，以便抽象能够正确组合。我经常发现自己盯着一块黑色的画布（因为这几乎是您在 WebGL 中的错误报告方面得到的全部内容）并狂躁的查找哪个破指针当前没有指向正确的地方。老实说，我不知道 ThreeJS 是如何变得如此健壮的，但它确实以某种方式做到了。我认为这是大多数人使用 ThreeJS 而不是直接使用 WebGL 的主要原因之一。

不是你，是我：明确地说，我无法认同 WebGL 可能是我自己的缺点。比我聪明的人已经能够使用 WebGL（和web 外的 OpenGL）构建令人惊叹的东西，但它从来没有真正让我满意。

随着机器学习、神经网络和加密货币的出现，GPU 已经证明它们不仅可以用于在屏幕上绘制三角形。使用 GPU 进行任何类型的计算通常被称为通用 GPU 或 GPGPU，而 WebGL 1 在这方面并不擅长。如果你想在 GPU 上处理任意数据，必须将其编码为纹理，在着色器中解码，进行计算，然后将结果重新编码为纹理。 WebGL 2 通过 Transform Feedback 使这一切变得容易得多，但直到 2021 年 9 月，Safari 才支持 WebGL2（而大多数其他浏览器自 2017 年 1 月以来就支持 WebGL2），因此它并不是一个真正的选择。即便如此，WebGL2 的某些限制仍然让它感觉有些笨拙。

2. WebGPU

在 web 之外，新一代图形 API 已经建立起来，它们向外部公开了一个更底层图形卡接口。这些新的 API 升级了设计 OpenGL 时不存在的新用例和约束。一方面，GPU 现在几乎无处不在。甚至移动设备都内置了功能强大的 GPU。因此，现代图形编程（3D 渲染和光线追踪）和 GPGPU 用例越来越普遍。另一方面，大多数设备都有多核处理器，因此能够从多线程与 GPU 交互可能是一个重要的优化方向。在 WebGPU 人员参与其中的同时，他们还重新审视了一些以前的设计决策，并预先加载了 GPU 必须完成的大量验证工作，从而使开发人员能够从 GPU 中榨取更多性能。

最流行的下一代 GPU API 是 Khronos Group 的 Vulkan、Apple 的 Metal 和 Microsoft 的 DirectX 12。为了将这些新功能带到web上，WebGPU 应运而生。比起 WebGL 只是 OpenGL 的简单封装，但 WebGPU 选择了不同的方法。它引入了自己的抽象并且不直接简单复制任何这些本机 API。这部分是因为没有单一的 API 在所有系统上都可用，但也因为许多概念（例如极低级别的内存管理）对于面向 Web 的 API 而言并不惯用。取而代之的是，WebGPU 的设计既让人感觉“web化”，又能舒适地位于任何原生图形 API 之上，同时抽象出它们的特性。它正在 W3C 中标准化，所有主要的浏览器供应商都参与其中。由于其相对低级的特性和强大的功能，WebGPU 有一点学习曲线并且在设置上相对繁琐，但我会尽我所能将其分解。

2.1 适配器和设备(Adapters and Devices)

首先接触到的 WebGPU 抽象是适配器(adapters )和（逻辑）设备(devices)。

抽象层次，从物理 GPU 到逻辑设备。

物理设备就是 GPU 本身，通常区分为集成 GPU 和独立 GPU。通常，任何给定设备都只有一个 GPU，但也可能有两个或更多 GPU。例如，微软的 SurfaceBook 以其低功耗集成 GPU 和高性能独立 GPU 而著称，操作系统将在这两者之间按需切换。

由 GPU 厂商提供的驱动程序(driver )将以操作系统理解和期望的方式向操作系统 暴漏 GPU 的功能。操作系统反过来可以使用操作系统提供的图形 API（如 Vulkan 或 Metal）将其 暴漏 给应用程序。

GPU 是一种共享资源。它不仅可以被许多应用程序同时使用，而且还控制着您在显示器上看到的内容。需要有一些东西可以让多个进程同时使用 GPU，这样每个应用程序都可以在屏幕上显示自己的 UI，而不会干扰其他应用程序甚至恶意读取其他应用程序的数据。对于每个进程来说，看起来它们对物理 GPU 拥有唯一的控制权，但事实显然并非如此。第一层 多路复用主要由驱动程序和操作系统完成。

适配器（Adapters）是从操作系统的本机图形 API 到 WebGPU 的转换层。由于 浏览器是 操作系统级应用程序，在浏览器中 还可以运行多个 Web 应用程序，因此 第二层需要多路复用，以便每个 Web 应用程序都感觉它可以单独控制 GPU。这是在 WebGPU 中使用逻辑设备(Logical Device)概念建模的。

要访问适配器，您可以调用 navigator.gpu.requestAdapter() 。在撰写本文时， requestAdapter() 只有很少的配置选项。这些选项允许您请求高性能或低能耗适配器。

软件渲染：一些实现还为没有 GPU 或 GPU 能力不足的系统提供“后备适配器”。
后备适配器实际上是一种纯软件实现，速度不会很快，但可以保持应用程序正常运行。

如果成功，即返回的适配器是非 null ，您可以检查适配器的功能并使用 adapter.requestDevice() 从适配器请求逻辑设备(logical device)。

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 的最低公分母。详细信息在 WebGPU 标准中指定。例如，尽管我的 GPU 能够轻松处理最大 4GiB 的数据缓冲区，但返回的 device 将只允许最大 1GiB 的数据缓冲区，并且会拒绝任何更大的数据缓冲区。这可能看起来有限制，但实际上非常有用：如果您的 WebGPU 应用程序在默认设备上运行，它将在绝大多数设备上运行。如有必要，您可以通过 adapter.limits 检查物理 GPU 的实际限制，并通过将选项对象传递给 requestDevice() 来请求提高限制的 设备。

2.2 着色器（Shaders ）

如果您曾经使用过 WebGL，您可能熟悉顶点着色器和片段着色器。无需过多深入，传统设置的工作原理如下：将数据缓冲区上传到 GPU 并告诉它如何将数据解释为一系列三角形。每个顶点占据该数据缓冲区的一块，描述该顶点在 3D 空间中的位置，但可能还有辅助数据，如颜色、纹理 ID、法线和其他内容。列表中的每个顶点都由 GPU 在顶点阶段处理，在每个顶点上运行顶点着色器，这将应用平移、旋转或透视变形等变换。

着色器(Shaders)：“着色器”这个词曾经让我感到困惑，因为你可以做的不仅仅是着色。
但在过去（即 1980 年代后期！），这个术语是恰当的：它只是在 GPU 上运行的一小段代码，
用于决定每个像素应该是什么颜色，以便您可以对正在渲染的对象进行着色，
从而实现灯光和阴影的视觉效果。如今，着色器泛指在 GPU 上运行的任何程序。

随后 GPU 将三角形光栅化，这意味着 GPU 计算出每个三角形在屏幕上覆盖了哪些像素。然后每个像素由片段着色器处理，片段着色器可以访问像素坐标，也可以访问辅助数据来决定该像素应该是哪种颜色。如果使用得当，可以通过此过程创建令人惊叹的 3D 图形。

这种将数据传递给顶点着色器，然后传递给片段着色器，然后将其直接输出到屏幕上的系统称为渲染管线（pipeline），在 WebGPU 中，您必须明确定义管线。

2.3 管线(Pipelines)

当前，WebGPU 允许您创建两种类型的管线：渲染管线和计算管线。顾名思义，渲染管线渲染某些东西，这意味着它创建一个 2D 图像。该图像不一定非得是在屏幕上，也可以只渲染到内存（称为帧缓冲区）。计算管线更通用，因为它返回一个缓冲区，该缓冲区可以包含任何类型的数据。对于这篇博文的剩余部分，将重点关注计算管线，因为我喜欢将渲染管线视为计算管线的专业化/优化。现在，这在历史上是倒置的——计算管线是作为对专门构建的渲染管线的基础而构建的——而这些管线在 GPU 中只是物理上不同的电路。然而，就 API 而言，我发现这种心智模型非常有用。在未来，似乎有更多类型的管线——可能是光线追踪管线——将被添加到 WebGPU 中。

使用 WebGPU，管线由一个（或多个）可编程阶段组成，其中每个阶段由着色器和入口点定义。计算管线有一个 compute 阶段，而渲染管线有一个 vertex 和 fragment 阶段：

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

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 编译为底层系统期望的任何内容。这可能是 DirectX 12 的 HLSL、Metal 的 MSL 和 Vulkan 的 SPIR-V。

SPIR-V：SPIR-V 很有趣，因为它是一种开放的二进制中间格式，由 Khronos Group 标准化。
您可以将 SPIR-V 视为并行编程语言编译器的 LLVM，它支持将多种语言编译为 SPIR-V 
以及将 SPIR-V 编译为多种其他语言。

In the shader module above we are just creating a function called main and marking it as an entry point for the compute stage by using the @compute attribute. You can have multiple functions marked as an entry point in a shader module, as you can reuse the same shader module for multiple pipelines and choose different functions to invoke via the entryPoint options. But what is that @workgroup_size(64) attribute?

在上面的着色器模块中，我们只是创建了一个名为 main 的函数，并使用 @compute 属性将其标记为计算阶段的入口点。您可以将多个函数标记为着色器模块中的入口点，因为可以为多个管道重用相同的着色器模块，并通过 entryPoint 选项 选择不同的函数来调用。但@workgroup_size(64) 属性是什么意思？

2.4 并行（Parallelism)

GPU 以延迟为代价针对吞吐量进行了优化。要理解这一点，我们必须稍微了解一下 GPU 的架构。我不想（老实说，不能）完整地解释它。我会尽可能深入，因为我觉得有必要。如果您想了解更多，Fabian Giesen 的这个由 13 部分组成的博文系列真的很棒。

众所周知，GPU 拥有大量内核，可以进行大规模并行工作。但是，在为多核 CPU 编程时，内核并不像您习惯的那样独立。首先，GPU 内核按层次分组。层次结构 的术语在厂商和 API 之间并不一致。英特尔有一份很好的文档，对他们的架构进行了高级概述，我被告知可以安全地假设其他 GPU 至少以类似的方式工作，尽管 GPU 的确切架构是受 NDA 保护的秘密。

就英特尔而言，层次结构中的最低级别是“执行单元”(EU)，它具有多个（在本例中为七个）SIMT 内核。这意味着它有七个以锁步方式运行并始终执行相同指令的内核。但是，每个内核都有自己的一组寄存器和堆栈指针。因此，虽然他们必须执行相同的操作，但他们可以在不同的数据上执行。这也是 GPU 性能专家避免分支（如 if / else 或循环）的原因：如果 EU 遇到 if / else ，则所有内核都必须执行两个分支，除非所有内核都碰巧采取相同的分支。每个核心都可以被告知忽略它正在输入的指令，但这显然浪费了本可以用于计算的宝贵周期。这同样适用于循环！如果一个核心提前完成循环，它将不得不假装执行循环体，直到所有核心都完成循环。

Despite the core’s frequency, getting data from memory (or pixels from textures) still takes relatively long — Fabian says it takes a couple hundred clock cycles. These couple hundred cycles could be spent on computation instead. To make use of these otherwise idle cycles, each EU is heavily oversubscribed with work. Whenever an EU would end up idling (e.g. to wait for a value from memory), it instead switches to another work item and will only switch back once the new work item needs to wait for something. This is the key trick how GPUs optimize for throughput at the cost of latency: Individual work items will take longer as a switch to another work item might stop execution for longer than necessary, but the overall utilization is higher and results in a higher throughput. The GPU strives to always have work queued up to keep EUs busy at all times.

尽管核心频率很高，但从内存（或从纹理获取像素）获取数据仍然需要相对较长的时间——Fabian 说这需要几百个时钟周期。这几百个周期本可以用在计算上。为了利用这些原本空闲的周期，每个 EU 都会被严重超量分配工作(each EU is heavily oversubscribed with work)。每当 EU 空闲时（例如等待内存中的值）时，它就会切换到另一个工作项，并且只有在新工作项需要等待时才会切换回来。这是 GPU 如何以延迟为代价优化吞吐量的关键技巧(即 切换到别的 工作项 会 延迟 当前工作项，以延迟单个工作项来提高整体吞吐量)：单个工作项将花费更长的时间，因为切换到另一个工作项可能会停止执行的时间超过必要的时间，但总体利用率更高并导致更高的吞吐量。 GPU 努力让工作始终排队，以使 EU 始终保持忙碌。

Intel Iris Xe Graphics 芯片的架构。 EU 有 7 个 SIMT 核心。 SubSlice 有 8 个 EU。 8 个 SubSlices 组成一个 Slice。

EUs are just the lowest level in the hierarchy, though. Multiple EUs are grouped into what Intel calls a “SubSlice”. All the EUs in a SubSlice have access to a small amount of Shared Local Memory (SLM), which is about 64KiB in Intel’s case. If the program to be run has any synchronization commands, it has to be executed within the same SubSlice, as only they have shared memory for synchronization.

不过，EU 只是层次结构中的最低级别。多个 EU 被分组为英特尔所谓的“SubSlice”。 SubSlice 中的所有 EU 都可以访问少量本地共享内存 ( Shared Local Memory ， SLM)，在英特尔的情况下约为 64KiB。如果要运行的程序有任何同步命令，它必须在同一个 SubSlice 中执行，因为只有它们有同步共享内存。

在最后一层，多个 SubSlice 被组合成一个 Slice，形成 GPU。对于集成的 Intel GPU，您最终会得到总共 170-700 个内核。独立的 GPU 可以轻松拥有 1500 个或更多内核。同样，这里的命名取自英特尔，其他供应商可能使用不同的名称，但每个 GPU 的总体架构都是相似的。

为了充分利用这种架构的好处，需要专门为这种架构编写程序，以便可以最大限度地利用GPU。因此，图形 API 公开了一个线程模型，该模型自然允许以这种方式分解工作。在 WebGPU 中，这里重要的原语是“工作组（workgroup）”。

2.5 工作组(Workgroups)

在传统设置中，顶点着色器会为每个顶点调用一次，片段着色器会为每个像素调用一次（我知道，我在这里掩盖了一些细节）。在 GPGPU 设置中，您的计算着色器将针对您计划的**每个工作项(work item)**调用一次。但工作项是什么由您来定义 。

The collection of all work items (which I will call the “workload”) is broken down into workgroups. All work items in a workgroup are scheduled to run together. In WebGPU, the work load is modelled as a 3-dimensional grid, where each “cube” is a work item, and work items are grouped into bigger cuboids to form a workgroup.

所有工作项（我称之为“工作量workload”）的集合被分解成工作组(workgroup)。工作组中的所有工作项都安排在一起运行。在 WebGPU 中，工作量 被建模为 3 维网格，其中每个“立方体”都是一个工作项，工作项被分组到更大的长方体中以形成一个工作组。

这是一个工作量。白边立方体是一个工作项。红边长方体是一个工作组。

终于，我们有足够的信息来讨论 @workgroup_size(x, y, z) 属性，此时它甚至可能大部分是不言自明的：该属性允许您告诉 GPU 此着色器的工作组大小应该是多少。或者用上图的语言来说， @workgroup_size 属性定义了红边立方体的大小。
x*y*z 是每个工作组的工作项数。任何跳过的参数都默认为 1，因此 @workgroup_size(64) 等同于 @workgroup_size(64, 1, 1) 。

Of course, the actual EUs are not arranged in the 3D grid on the chip. The aim of modelling work items in a 3D grid is to increase locality. The assumption is that it is likely that neighboring work groups will access similar areas in memory, so when running neighboring workgroups sequentially, the chances of already having values in the cache are higher, saving a couple of hundred cycles by not having to grab them from memory. However, most hardware seemingly just runs workgroups in a serial order as the difference between running a shader with @workgroup_size(64) or @workgroup_size(8, 8) is negligible. So this concept is considered somewhat legacy.

当然，实际的EU并没有排列在芯片上的3D网格中。在 3D 网格中建模工作项的目的是增加局部性。假设相邻的工作组很可能会访问内存中的相似区域，因此当按顺序运行相邻的工作组时，缓存中已有值的可能性更高，因为不必从中获取它们而节省了数百个访存周期。然而，大多数硬件似乎只是按顺序运行工作组，因为使用 @workgroup_size(64) 或 @workgroup_size(8, 8) 运行着色器之间的区别可以忽略不计。所以这个概念被认为有点遗留。

但是，工作组在多个方面受到限制： device.limits 有一堆值得了解的属性：

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

工作组大小的每个维度的大小都受到限制，但即使 x、y 和 z 分别在限制范围内，它们的乘积 (=x×y×z ) 也可能不在范围内，因为它有自己的限制。最后，每个维度只能有这么多工作组。

专业提示：不要产生最大数量的线程。尽管 GPU 由操作系统和底层调度程序管理，
但您可能会使用长时间运行的 GPU 程序冻结整个系统。

那么合适的工作组规模是多少？这实际上取决于您分配工作项坐标的语义。我确实意识到这并不是一个真正有用的答案，所以我想给你与 Corentin 给我的相同建议：“使用 [工作组大小] 64，除非你知道你的目标是什么 GPU 或者你的工作负载需要不同的东西”这似乎是一个安全的数字，在许多 GPU 上表现良好，并允许 GPU 调度程序让尽可能多的 EU 保持忙碌。

2.6 命令（Commands ）

我们已经编写了着色器并设置了管道。剩下要做的就是实际调用 GPU 来执行这一切。由于 GPU 可以是一个完全独立的外置卡，带有自己的内存芯片，因此可以通过所谓的命令缓冲区(command buffer)或命令队列(command queue)来控制它。命令队列是一块内存，其中包含供 GPU 执行的编码命令。编码 高度依赖 GPU，由驱动程序负责。 WebGPU 公开了一个 CommandEncoder 以利用该功能。

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

commandEncoder 有多种方法允许您将数据从一个 GPU 缓冲区复制到另一个 GPU 缓冲区并操作纹理。它还允许您创建 PassEncoder ，它对管线的设置和调用进行编码。在这种情况下，我们有一个计算管线，所以我们必须创建一个计算通道，将其设置为使用我们预先声明的管线，最后调用 dispatchWorkgroups(w_x, w_y, w_z) 告诉 GPU 沿每个维度创建多少个工作组。换句话说，我们的计算着色器将被调用的次数等于
。
顺便说一下，pass 编码器是 WebGPU 的抽象，以避免我在这篇博文开始时大谈特谈的内部全局状态对象。运行 GPU 管线所需的所有数据和状态都通过传递编码器显式传递。

抽象：命令缓冲区也是驱动程序或操作系统的钩子函数(hook)，让多个应用程序在不相互干扰的情况下使用 GPU。
当您将命令排队时，下面的抽象层会将额外的命令注入队列以保存先前程序的状态并恢复您的程序状态，
这样就感觉没有其他人在使用 GPU。

运行这段代码，我们实际上在 GPU 上生成了 64 个线程，它们什么都不做。但它能运行，这很酷。让我们谈谈如何为 GPU 提供一些工作数据。

3. 交换数据 (Exchanging data)

正如承诺的那样，我不会直接将 WebGPU 用于图形，因此我认为在 GPU 上运行物理模拟并使用 Canvas2D 将其可视化会很有趣。也许我自鸣得意地称它为“物理模拟”——实际上是生成一大堆圆圈，让它们在一个平面上以随机方向滚动并让它们碰撞。

为此，我们需要将一些模拟参数和初始状态推送到 GPU，在 GPU 上运行模拟 并从 GPU 读取模拟结果。这可以说是 WebGPU 中最令人发怵的部分，因为有一堆操作数据的特殊技巧（不是说看似毫无意义的复制），但这就是允许 WebGPU 成为以最高性能水平运行的与设备无关的 API 的原因。

3.1 绑定组布局 （Bind Group Layouts ）

为了与 GPU 交换数据，我们需要使用 绑定组布局 来扩展我们的管线定义。绑定组是 GPU 实体（内存缓冲区、纹理、采样器等）的集合，可在管线执行期间访问这些实体。绑定组布局预先定义了这些 GPU 实体的类型、目的和使用方法(types, purposes and uses)，这使得 GPU 可以提前弄清楚如何最有效地运行管线。让我们在这个初始步骤中保持简单，并让管线访问单个内存缓冲区：

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 数字可以自由选择，用于将我们的 WGSL 代码中的变量绑定到绑定组布局的这个插槽中的缓冲区内容。

使用方法如下:

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

bindGroupLayout 还定义了每个缓冲区的用途，在本例中为 “storage” 。另一个选项是 “read-only-storage” ，它是只读的（duh！），并且允许 GPU 在 永远不会写入此缓冲区的基础上 进行进一步优化，因此不需要同步。缓冲区类型的还有一个可能值是 “uniform” ，在计算管线的上下文中，它在功能上基本上等同于存储缓冲区。

绑定组布局就位。现在我们可以创建绑定组本身，其中包含绑定组布局所需的 GPU 实体的实际实例。一旦带有内部缓冲区的绑定组就位，计算着色器就可以用数据填充它，我们可以从 GPU 读取它。但是有一个障碍：暂存缓冲区(Staging Buffers)。

3.2 暂存缓冲区(Staging Buffers )

再强调一遍：GPU 以延迟为代价针对吞吐量进行了大量优化。 GPU 需要能够以令人难以置信的高速率将数据馈送到内核以维持该吞吐量。 Fabian 在他 2011 年的博客文章系列中做了一些粗略计算，得出的结论是，对于以 1280x720 分辨率运行的着色器的纹理样本，GPU 需要维持 3.3GB/s。为了适应当今的图形需求，GPU 需要更快地获取数据。这只有在 GPU 的内存与内核非常紧密地集成时才有可能实现。这种紧密集成使得很难将相同的内存也暴露给主机进行读写。

相反，GPU 有额外的内存，主机和 GPU 都可以访问这些内存，但集成度不高，无法快速提供数据。暂存缓冲区是在此中间内存领域中分配的缓冲区，可以映射到主机系统进行读写。为了从 GPU 读取数据，我们将数据从内部高速缓冲区复制到暂存缓冲区，然后将暂存缓冲区映射到主机，以便可以将数据读回主内存。对于写操作，过程是相同的，但操作相反。

回到代码：我们将创建一个可写缓冲区并将其添加到绑定组，以便它可以被计算着色器写入。还将创建第二个具有相同大小的缓冲区，用作暂存缓冲区。每个缓冲区都使用 usage 位掩码创建，可以在其中声明您打算如何使用该缓冲区。然后 GPU 将确定缓冲区应该位于何处以实现所有这些用例，或者如果标志组合不合法则抛出错误。

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 。目前还不能读取或写入它们。为此，它们需要被映射，这是一个单独的 API 调用，并且只有 GPUBufferUsage.MAP_READ 或 GPUBufferUsage.MAP_WRITE 的缓冲区才可行。

TypeScript：我发现 TypeScript 在探索新 API 时非常有用。幸运的是，
Chrome 的 WebGPU 团队维护着 @webgpu/types ，因此您可以享受准确的自动完成。

现在我们不仅有绑定组布局，甚至还有实际的绑定组本身，需要更新调度代码以使用这个绑定组。之后，我们映射暂存缓冲区以将结果读回 JavaScript。

const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(pipeline);
passEncoder.setBindGroup(0, bindGroup);//设置绑定组
// passEncoder.dispatchWorkgroups(1);
passEncoder.dispatchWorkgroups(Math.ceil(BUFFER_SIZE / 64));
passEncoder.end();
commandEncoder.copyBufferToBuffer(
  output, //从哪里复制
  0, // Source offset 开始位置
  stagingBuffer, //复制到哪里
  0, // Destination offset //开始位置
  BUFFER_SIZE  //
);
const commands = commandEncoder.finish();
device.queue.submit([commands]);

await stagingBuffer.mapAsync(
  GPUMapMode.READ,
  0, // Offset
  BUFFER_SIZE // Length
 );
const copyArrayBuffer =
  stagingBuffer.getMappedRange(0, BUFFER_SIZE);
const data = copyArrayBuffer.slice();
stagingBuffer.unmap();
console.log(new Float32Array(data));

Since we added a bind group layout to our pipeline, any invocation without providing a bind group would now fail. After we define our “pass”, we add an additional command via our command encoder to copy the data from our output buffer to the staging buffer and submit our command buffer to the queue. The GPU will start working through the command queue. We don’t know when the GPU will be done exactly, but we can already submit our request for the stagingBuffer to be mapped. This function is async as it needs to wait until the command queue has been fully processed. When the returned promise resolves, the buffer is mapped, but not exposed to JavaScript yet. stagingBuffer.getMappedRange() let’s us request for a subsection (or the entire buffer) to be exposed to JavaScript as a good ol’ ArrayBuffer. This is real, mapped GPU memory, meaning the data will disappear (the ArrayBuffer will be “detached”), when stagingBuffer gets unmapped, so I’m using slice() to create JavaScript-owned copy.

由于我们在管线中添加了绑定组布局，因此任何不提供绑定组的调用现在都会失败。在定义了“pass”之后，通过命令编码器添加了一个额外的命令来将数据输出缓冲区复制到暂存缓冲区并将命令缓冲区提交到队列。 GPU 将通过命令队列开始工作。我们不知道 GPU 何时会完成，但已经可以提交对 stagingBuffer 进行映射的请求。这个函数是异步的，因为它需要等到命令队列被完全处理。当返回的 promise 解析时，缓冲区被映射，但还没有暴露给 JavaScript。 stagingBuffer.getMappedRange() 让我们请求将一个 返回（或整个缓冲区）公开给 JavaScript 作为一个 ArrayBuffer 。这真在映射 GPU 内存，这意味着数据将消失（ ArrayBuffer “分离”时），当 stagingBuffer 被取消映射时，所以使用 slice() 创建 JavaScript 独立的副本。

不是很令人兴奋，但我们从 GPU 的内存中复制了这些零。

零以外的东西可能会更有说服力。在开始在 GPU 上进行任何高级计算之前，让我们将一些精心挑选的数据放入缓冲区中，以证明管线确实按预期工作。这是新的计算着色器代码，为了清晰起见，有额外的间距。

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

@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 数组。两个属性声明数据来自何处：来自第一个（第 0 个）绑定组中的缓冲区， binding 值为 1 的条目。数组的长度将自动反映基础缓冲区的长度（向下舍入）。

变量：WGSL 与 Rust 的不同之处在于，用 let 声明的变量是不可变的。如果你想让一个变量是可变的，
他们使用的关键字是 var 。

main() 函数的签名增加了两个参数： global_id 和 local_id 。可以选择任何名称 — 它们的值由与其关联的属性决定： global_invocation_id 是一个内置值，对应于工作负载中此着色器调用的全局 x/y/z 坐标。 local_invocation_id 是该着色器在工作组中的 x/y/z 坐标。

工作负载中标记的三个工作项 a、b 和 c 的示例。

上图显示了对带有 @workgroup_size(4, 4, 4) 的工作负载的坐标系的一种可能解释。可以为您的用例定义合适的坐标系。如果我们同意上面绘制的轴，我们会看到 a、b 和 c 的以下 main() 参数：

• 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)

In our shader, we have @workgroup_size(64, 1, 1), so local_id.x will range from 0 to 63. To be able to inspect both values, I am “encoding” them into a single number. Note that WGSL is strictly typed: Both local_id and global_id are vec3, so we have to explicitly cast their values to f32 to be able to assign them to our f32 output buffer.

在着色器中，有 @workgroup_size(64, 1, 1) ，所以 local_id.x 的范围从 0 到 63。为了能够检查这两个值，将它们“编码”为一个数字。请注意，WGSL 是严格类型化的： local_id 和 global_id 都是 vec3 ，因此我们必须将它们的值显式转换为 f32 ，以便能够将它们分配给我们的 f32 输出缓冲区。

上图 :GPU 填写的实际值。请注意 local invocation ID 如何在 63 之后开始回绕，而 global invocation ID 保持不变。

这证明我们的计算着色器确实为输出内存中的每个值调用并用唯一值填充它。我们不知道这些数据的填充顺序，因为这是有意未指定的，由 GPU 的调度程序决定。

3.3 超量调度(Overdispatching)

敏锐的观察者可能已经注意到着色器调用的总数 ( Math.ceil(BUFFER_SIZE / 64) * 64 )=1024 将导致 global_id.x=1000 变得比数组的长度更大，因为每个 f32 占用 4 个字节。幸运的是，访问数组越界是受保护的，因此每次超过数组末尾的写入都将最终写入数组的最后一个元素。这避免了内存访问错误，但仍可能生成无法使用的数据。事实上，如果您检查返回缓冲区的最后 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);
}

如果需要，您可以运行此演示并检查完整源代码。

3.4 A structure for the madness

我们的目标是让大量的球在 2D 空间中移动并产生愉快的小碰撞。为此，每个球都需要有一个半径、一个位置和一个速度矢量。我们可以继续处理 array ，假设第一个浮点数是第一个球的 x 位置，第二个浮点数是第一个球的 y 位置，依此类推。那不是我所说的符合人体工程学的东西。还好，WGSL 允许定义自己的结构，将多个数据片段捆绑在一起。

老消息：如果你知道什么是内存对齐，你可以跳过这一部分（尽管一定要看一下代码示例）。如果你不知道它是什么，
我不会真正解释原因，而是向你展示它是如何表现出来的以及如何使用它。

所以用所有这些组件定义一个 struct Ball 并将我们的 array 变成 array 是有意义的。所有这一切的缺点是：我们必须谈论对齐。

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

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

@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);
}

如果你运行这个演示，你会在你的控制台中看到：

由于对齐约束，该结构在其内存布局中有一个空洞（填充）。

I put 999 the first field of the struct to make it easy to see where the struct begins in the buffer. There’s a total of 6 numbers until we reach the next 999, which is a bit surprising because the struct really only has 5 numbers to store: radius, position.x, position.y, velocity.x and velocity.y. Taking a closer look, it is clear that the number after radius is always 0. This is because of alignment.

我将 999 放在结构的第一个字段中，以便于查看结构在缓冲区中的开始位置。在我们到达下一个 999 之前总共有 6 个数字，这有点令人惊讶，因为该结构实际​​上只有 5 个数字要存储： radius 、 position.x 、 position.y 、 velocity.x 和 velocity.y .仔细一看，很明显 radius 后面的数字一直是0，这是对齐的缘故。

每种 WGSL 数据类型都有明确定义的对齐要求。如果数据类型的对齐方式为 N ，则意味着该数据类型的值只能存储在 N 的倍数的内存地址中。 f32 的对齐方式为 4，而 vec2 的对齐方式为 8。如果假设结构从地址 0 开始，那么 radius 字段可以存储在地址 0，因为 0 是 4 的倍数。结构中的下一个字段是 vec2 ，它的对齐方式是 8。但是， radius 之后的第一个空闲地址是 4，它不是 8 的倍数。为了解决这个问题，编译器添加了 4 个字节的填充到达下一个 8 的倍数的地址。这解释了我们在 DevTools 控制台中看到一个值为 0 的未使用字段。

现在知道结构在内存中是如何布局的，我们可以从 JavaScript 填充它以生成 球的初始状态，并读回它以可视化它。

3.5 输入输出（Input & Output ）

我们已经成功地从 GPU 读取数据，将它带到 JavaScript 并“解码”它。现在是时候解决另一个方向了。我们需要在 JavaScript 中生成所有球的初始状态并将其提供给 GPU，以便它可以在其上运行计算着色器。生成初始状态相当简单：

let inputBalls = new Float32Array(new ArrayBuffer(BUFFER_SIZE));
for (let i = 0; i < NUM_BALLS; i++) {
  inputBalls[i * 6 + 0] = randomBetween(2, 10); // radius
  inputBalls[i * 6 + 1] = 0; // padding
  inputBalls[i * 6 + 2] = randomBetween(0, ctx.canvas.width); // position.x
  inputBalls[i * 6 + 3] = randomBetween(0, ctx.canvas.height); // position.y
  inputBalls[i * 6 + 4] = randomBetween(-100, 100); // velocity.x
  inputBalls[i * 6 + 5] = randomBetween(-100, 100); // velocity.y
}

Buffer-backed-object：使用更复杂的数据结构，从 JavaScript 操作数据会变得非常乏味。
虽然最初是为 worker 用例编写的，但我的库 buffer-backed-object 在这里可以派上用场！

我们也已经知道如何将缓冲区暴露给着色器。只需要调整我管线绑定组布局以期待另一个缓冲区：

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

…并创建一个可以使用绑定组绑定的 GPU 缓冲区：

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,
      },
    },
  ],
});

Now for the new part: Sending data to the GPU. Just like with reading data, we technically have to create a staging buffer that we can map, copy our data into the staging buffer and then issue a command to copy our data from the staging buffer into the storage buffer. However, WebGPU offers a convenience function that will choose the most efficient way of getting our data into the storage buffer for us, even if that involves creating a temporary staging buffer on the fly:

现在是新部分：将数据发送到 GPU。就像读取数据一样，在技术上必须创建一个可以映射的暂存缓冲区，将数据复制到暂存缓冲区，然后发出命令将数据从暂存缓冲区复制到存储缓冲区。然而，WebGPU 提供了一个便利的功能，它将为我们选择最有效的方式将数据放入存储缓冲区，即使这涉及动态创建临时暂存缓冲区：

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

就这？嗯，就是这样！我们甚至不需要命令编码器。我们可以直接将这个命令放入命令队列。 device.queue 还为纹理提供了一些其他类似的便利功能。

现在需要将这个新缓冲区绑定到 WGSL 中的一个变量并对其进行一些操作：

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

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

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

const TIME_STEP: f32 = 0.016;

@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;
}

我希望此时此着色器代码的绝大部分内容不会让您感到意外。

每一帧，WebGPU 都用于更新球的位置。它们使用 Canvas2D 绘制到屏幕上。

最后，我们需要做的就是将 output 缓冲区读回 JavaScript，编写一些 Canvas2D 代码来可视化缓冲区的内容，并将其全部放入 requestAnimationFrame() 循环中。你可以看到这个演示的结果。

4. 性能(Performance)

前面的演示只是沿着它们的速度矢量移动每个球。不完全令人兴奋,计算也并不复杂。在查看性能之前，让我在着色器中进行一些适当的物理计算。我不会在这里解释它们——这篇博文已经够长了——但我会说我采用了最天真的方法：每个球都检查是否与其他球发生碰撞。如果您好奇，可以查看最终演示的源代码，其中还包含指向我用来编写 physics-y 位的资源的链接。

… 现在有了弹力墙和弹力球！

最终代码如下:

import "../../../modulepreload-polyfill.b7f2da20.js";
let v = 0;
const m = new URLSearchParams(location.search);
function l(e, a) {
    return m.has(e) ? parseFloat(m.get(e)) : a
}
const _ = l("balls", 100)
  , s = _ * 6 * Float32Array.BYTES_PER_ELEMENT
  , M = l("min_radius", 2)
  , G = l("max_radius", 10)
  , x = l("render", 1)
  , t = document.querySelector("canvas").getContext("2d");
t.canvas.width = l("width", 500);
t.canvas.height = l("height", 500);
function h(e) {
    throw document.body.innerHTML = `<pre>${e}</pre>`,
    Error(e)
}
"gpu"in navigator || h("WebGPU not supported. Please enable it in about:flags in Chrome or in about:config in Firefox.");

const P = await navigator.gpu.requestAdapter();
P || h("Couldn\u2019t request WebGPU adapter.");
const i = await P.requestDevice();
i || h("Couldn\u2019t request WebGPU device.");
const T = i.createShaderModule({
    code: `
    struct Ball {
      radius: f32,
      position: vec2<f32>,
      velocity: vec2<f32>,
    }

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

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

    struct Scene {
      width: f32,
      height: f32,
    }

    @group(0) @binding(2)
    var<storage, read> scene: Scene;

    const PI: f32 = 3.14159;
    const TIME_STEP: f32 = 0.016;

    @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;
      }
      var src_ball = input[global_id.x];
      let dst_ball = &output[global_id.x];

      (*dst_ball) = src_ball;

      // Ball/Ball collision
      for(var i = 0u; i < num_balls; i = i + 1u) {
        if(i == global_id.x) {
          continue;
        }
        var other_ball = input[i];
        let n = src_ball.position - other_ball.position;
        let distance = length(n);
        if(distance >= src_ball.radius + other_ball.radius) {
          continue;
        }
        let overlap = src_ball.radius + other_ball.radius - distance;
        (*dst_ball).position = src_ball.position + normalize(n) * overlap/2.;

        // Details on the physics here:
        // https://physics.stackexchange.com/questions/599278/how-can-i-calculate-the-final-velocities-of-two-spheres-after-an-elastic-collisi
        let src_mass = pow(src_ball.radius, 2.0) * PI;
        let other_mass = pow(other_ball.radius, 2.0) * PI;
        let c = 2.*dot(n, (other_ball.velocity - src_ball.velocity)) / (dot(n, n) * (1./src_mass + 1./other_mass));
        (*dst_ball).velocity = src_ball.velocity + c/src_mass * n;
      }

      // Apply velocity
      (*dst_ball).position = (*dst_ball).position + (*dst_ball).velocity * TIME_STEP;

      // Ball/Wall collision
      if((*dst_ball).position.x - (*dst_ball).radius < 0.) {
        (*dst_ball).position.x = (*dst_ball).radius;
        (*dst_ball).velocity.x = -(*dst_ball).velocity.x;
      }
      if((*dst_ball).position.y - (*dst_ball).radius < 0.) {
        (*dst_ball).position.y = (*dst_ball).radius;
        (*dst_ball).velocity.y = -(*dst_ball).velocity.y;
      }
      if((*dst_ball).position.x + (*dst_ball).radius >= scene.width) {
        (*dst_ball).position.x = scene.width - (*dst_ball).radius;
        (*dst_ball).velocity.x = -(*dst_ball).velocity.x;
      }
      if((*dst_ball).position.y + (*dst_ball).radius >= scene.height) {
        (*dst_ball).position.y = scene.height - (*dst_ball).radius;
        (*dst_ball).velocity.y = -(*dst_ball).velocity.y;
      }
    }
  `
})
  , w = i.createBindGroupLayout({
    entries: [{
        binding: 0,
        visibility: GPUShaderStage.COMPUTE,
        buffer: {
            type: "read-only-storage"
        }
    }, {
        binding: 1,
        visibility: GPUShaderStage.COMPUTE,
        buffer: {
            type: "storage"
        }
    }, {
        binding: 2,
        visibility: GPUShaderStage.COMPUTE,
        buffer: {
            type: "read-only-storage"
        }
    }]
})
  , A = i.createComputePipeline({
    layout: i.createPipelineLayout({
        bindGroupLayouts: [w]
    }),
    compute: {
        module: T,
        entryPoint: "main"
    }
})
  , B = i.createBuffer({
    size: 2 * Float32Array.BYTES_PER_ELEMENT,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
})
  , E = i.createBuffer({
    size: s,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
})
  , S = i.createBuffer({
    size: s,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
})
  , p = i.createBuffer({
    size: s,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
})
  , C = i.createBindGroup({
    layout: w,
    entries: [{
        binding: 0,
        resource: {
            buffer: E
        }
    }, {
        binding: 1,
        resource: {
            buffer: S
        }
    }, {
        binding: 2,
        resource: {
            buffer: B
        }
    }]
});
function R() {
    return new Promise(e=>requestAnimationFrame(e))
}
let r = new Float32Array(new ArrayBuffer(s));
for (let e = 0; e < _; e++)
    r[e * 6 + 0] = d(M, G),
    r[e * 6 + 2] = d(0, t.canvas.width),
    r[e * 6 + 3] = d(0, t.canvas.height),
    r[e * 6 + 4] = d(-100, 100),
    r[e * 6 + 5] = d(-100, 100);
let b;
i.queue.writeBuffer(B, 0, new Float32Array([t.canvas.width, t.canvas.height]));
for (; ; ) {
    performance.mark("webgpu start"),
    i.queue.writeBuffer(E, 0, r);
    const e = i.createCommandEncoder()
      , a = e.beginComputePass();
    a.setPipeline(A),
    a.setBindGroup(0, C);
    const o = Math.ceil(_ / 64);
    a.dispatchWorkgroups(o),
    a.end(),
    e.copyBufferToBuffer(S, 0, p, 0, s);
    const n = e.finish();
    i.queue.submit([n]),
    await p.mapAsync(GPUMapMode.READ, 0, s);
    const c = p.getMappedRange(0, s).slice();
    b = new Float32Array(c),
    p.unmap(),
    performance.mark("webgpu end"),
    performance.measure("webgpu", "webgpu start", "webgpu end"),
    x !== 0 ? q(b) : (v++,
    t.fillStyle = v % 2 == 0 ? "red" : "blue",
    t.fillRect(0, 0, t.canvas.width, t.canvas.height)),
    r = b,
    await R()
}
function q(e) {
    t.save(),
    t.scale(1, -1),
    t.translate(0, -t.canvas.height),
    t.clearRect(0, 0, t.canvas.width, t.canvas.height),
    t.fillStyle = "red";
    for (let a = 0; a < e.length; a += 6) {
        const o = e[a + 0]
          , n = e[a + 2]
          , f = e[a + 3]
          , c = e[a + 4]
          , U = e[a + 5];
        let u = Math.atan(U / (c === 0 ? Number.EPSILON : c));
        c < 0 && (u += Math.PI);
        const g = n + Math.cos(u) * Math.sqrt(2) * o
          , y = f + Math.sin(u) * Math.sqrt(2) * o;
        t.beginPath(),
        t.arc(n, f, o, 0, 2 * Math.PI, !0),
        t.moveTo(g, y),
        t.arc(n, f, o, u - Math.PI / 4, u + Math.PI / 4, !0),
        t.lineTo(g, y),
        t.closePath(),
        t.fill()
    }
    t.restore()
}
function d(e, a) {
    return Math.random() * (a - e) + e
}

我不想对这个实验进行任何精确的测量，因为我没有优化物理算法，也没有优化我对 WebGPU 的使用。然而，即使是这种幼稚的实现也表现得非常好（在我的 M1 MacBook Air 上）这一事实让我印象深刻。在降到 60fps 以下之前，我可以处理大约 2500 个球。然而，观察轨迹，很明显在 2500 个球时，瓶颈是 Canvas2D 试图绘制场景，而不是 WebGPU 计算。

在 14000 个球时，原始 GPU 计算时间在 M1 MBA 上达到约 16 毫秒。

为了查看这到底有多快，我禁用了渲染，而是使用 performance.measure() 来查看在耗尽 16 毫秒的帧预算之前我可以模拟多少个球。这发生在我的机器上大约 14000 个球。这种未经优化的运行速度如此之快真的让我沉迷于 WebGPU 给我提供的强大计算能力。

5. 稳定性和可用性 (Stability & Availability)

WebGPU 已经开发了一段时间，我认为标准组急于宣布 API 稳定。话虽如此，该 API 仅在 Chrome 和 Firefox 中可用。我对 Safari 发布此 API 持乐观态度，但在撰写本文时，Safari TP 中尚无任何可看的内容。

在稳定性方面，甚至在我为本文进行研究时，也发生了一些变化。例如，属性的语法已从 [[stage(compute), workgroup_size(64)]] 更改为 @compute @workgroup_size(64) 。在撰写本文时，Firefox 仍在使用旧语法。 passEncoder.end() 曾经是 passEncoder.endPass() 。规范中还有一些内容尚未在任何浏览器中实现，例如着色器常量或可在移动设备上使用的 API。

基本上我要说的是：当浏览器和标准人员处于这个 API 的 ✨稳定 ✨ 之旅的最后阶段时，期待更多的突破性变化发生。

6. 结论(Conclusion )

拥有一个现代 API 来与web上的 GPU 对话将会非常有趣。在投入时间克服最初的学习曲线之后，我真的觉得自己能够使用 JavaScript 在 GPU 上运行大规模并行工作负载。还有 wgpu，它在 Rust 中实现了 WebGPU API，允许你在浏览器之外使用 API。 wgpu 还支持将 WebAssembly 作为编译目标，因此您可以在浏览器外部和浏览器内部通过 WebAssembly 本机运行 WebGPU 程序。有趣的事实：Deno 是第一个开箱即用的也支持 WebGPU 的运行时（感谢 wgpu）。

如果您有疑问或遇到问题，可以使用 Matrix 频道，其中有许多 WebGPU 用户、浏览器工程师和标准人员，他们对我提供了极大的帮助。去把你的脚弄湿！激动人心的时刻。

感谢 Brandon Jones 校对了本文，感谢 WebGPU Matrix 频道回答了我所有的问题。

原文地址