GPU 计算使用入门 (Web)

本博文通过示例探索了实验性 WebGPU API， 您将开始使用 GPU 执行数据并行计算。

François Beaufort

背景

您可能已经知道，图形处理器 (GPU) 是一种电子 子系统，这些子系统最初专用于处理 图形。然而，在过去的 10 年里，它逐渐发展成 该架构允许开发者实现许多类型的算法，而不仅仅是 渲染 3D 图形，同时充分利用 Google Cloud 平台的独特架构 GPU。这些功能称为 GPU 计算，将 GPU 用作 用于通用科学计算的协处理器称为通用协处理器， GPU (GPGPU) 编程。

GPU 计算对最近的机器学习热潮起到了重大贡献， 因为卷积神经网络和其他模型可以利用 更高效地在 GPU 上运行。使用当前的网络平台 由于缺乏 GPU 计算功能，W3C 的“适用于 Web 的 GPU”社区小组 正在设计一个 API，以公开大多数 现有设备此 API 称为 WebGPU。

WebGPU 是一种低级 API，例如 WebGL。这项功能功能强大，而且非常冗长 但没关系。我们注重的是效果。

在本文中，我将重点介绍 WebGPU 的“GPU 计算”部分。 老实说，我只是触及了皮毛，所以您可以开始在自己的设备上玩游戏了。 。我会更深入地介绍 WebGPU 渲染（画布、纹理、 等）。

访问 GPU

在 WebGPU 中可以轻松访问 GPU。正在呼叫navigator.gpu.requestAdapter() 返回一个使用 GPU 异步解析的 JavaScript promise 适配器。请将此适配器视为显卡。它可以 （在与 CPU 相同的芯片上）或离散卡（通常是 PCIe 卡， 但会消耗更多电量）。

有了 GPU 适配器后，请调用 adapter.requestDevice() 以获取 promise 将使用 GPU 设备解析，您将使用 GPU 设备执行一些 GPU 计算。

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

这两个函数都提供选项，可让您指定 所需的适配器（电源偏好设置）和设备（扩展程序、限制）。对于 为简单起见，我们将使用本文中的默认选项。

写入缓冲区内存

我们来看看如何使用 JavaScript 将数据写入 GPU 的内存。这个 这个过程并不简单，因为现代网络采用的沙盒模型 。

以下示例展示了如何将 4 个字节写入可访问的缓冲区内存 。它会调用 device.createBuffer()，后者会接受 缓冲区及其用途。虽然使用标志 GPUBufferUsage.MAP_WRITE 是 我们再明确说明一下，我们希望编写 至此缓冲区。它会导致在创建 GPU 缓冲区对象时进行映射，这要归功于 mappedAtCreation 设置为 true。然后，关联的原始二进制数据缓冲区 可通过调用 GPU 缓冲区方法 getMappedRange() 来检索。

如果您已使用 ArrayBuffer 完成过操作，那么写入字节就很熟悉；使用 TypedArray 并将值复制到其中。

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE
});
const arrayBuffer = gpuBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

此时，系统已映射 GPU 缓冲区，这意味着它归 CPU 所有，并且 可以从 JavaScript 以读写模式访问。这样 GPU 就可以访问它 必须取消映射，这与调用 gpuBuffer.unmap() 一样简单。

需要映射/未映射的概念，以防止 GPU 出现竞态条件 和 CPU 访问内存

读取缓冲区内存

现在，我们来看看如何将 GPU 缓冲区复制到另一个 GPU 缓冲区并读回该缓冲区。

由于我们是在第一个 GPU 缓冲区中写入数据，因此想要将其复制到第二个 GPU 缓冲区。 GPU 缓冲区，则需要新的用法标志 GPUBufferUsage.COPY_SRC。第二个 这次，创建 GPU 缓冲区时处于未映射状态， device.createBuffer()。其用法标志为 GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ，因为它将用作第一个 GPU 的目标 执行 GPU 复制命令后，在 JavaScript 中读取和缓冲数据。

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuWriteBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC
});
const arrayBuffer = gpuWriteBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

// Unmap buffer so that it can be used later for copy.
gpuWriteBuffer.unmap();

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: 4,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

由于 GPU 是一个独立的协处理器，因此所有 GPU 命令都会在 异步执行。因此，Google 会构建和发送 GPU 命令列表， 根据需要进行批量处理。在 WebGPU 中，由 Transformer 返回的 GPU 命令编码器 device.createCommandEncoder() 是一个 JavaScript 对象，用于构建一批 “已缓冲”在某个时间点发送到 GPU 的命令上的方法 另一方面，GPUBuffer 是“无缓冲”，这意味着它们以原子方式执行 。

有了 GPU 命令编码器后，请调用 copyEncoder.copyBufferToBuffer() 如下所示，将此命令添加到命令队列中，以便稍后执行。 最后，通过调用 copyEncoder.finish() 完成编码命令并提交 发送到 GPU 设备命令队列队列负责处理 使用 GPU 命令作为参数，通过 device.queue.submit() 完成提交。 这将以原子方式按顺序执行存储在数组中的所有命令。

// Encode commands for copying buffer to buffer.
const copyEncoder = device.createCommandEncoder();
copyEncoder.copyBufferToBuffer(
  gpuWriteBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  4 /* size */
);

// Submit copy commands.
const copyCommands = copyEncoder.finish();
device.queue.submit([copyCommands]);

此时，GPU 队列命令已发送，但未必执行。 如需读取第二个 GPU 缓冲区，请使用以下命令调用 gpuReadBuffer.mapAsync()： GPUMapMode.READ。它返回一个 promise，会在 GPU 缓冲区 映射。然后使用 gpuReadBuffer.getMappedRange() 获取映射的范围， 所有已排入队列的 GPU 命令包含与第一个 GPU 缓冲区相同的值 错误。

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

您可以试用此示例。

简而言之，您需要注意有关缓冲区内存操作的内容：

• GPU 缓冲区必须取消映射，才能在设备队列提交中使用。
• 映射后，可以使用 JavaScript 读取和写入 GPU 缓冲区。
• 当使用 mapAsync() 和 createBuffer() 时，系统会映射 GPU 缓冲区， 调用设置为 true 的 mappedAtCreation。

着色器编程

在 GPU 上运行的程序，这些程序只执行计算， 称为计算着色器。它们通过数百个 的 GPU 核心（比 CPU 核心小）一起运行以处理数据量 数据。它们的输入和输出是 WebGPU 中的缓冲区。

为了说明计算着色器在 WebGPU 中的用法，我们将使用矩阵 乘法，它是机器学习中的一种常见算法，如下所示。

<ph type="x-smartling-placeholder">

</ph>

矩阵乘法图

简而言之，我们将执行以下操作：

1. 创建三个 GPU 缓冲区（两个用于相乘，一个用于 结果矩阵）
2. 描述计算着色器的输入和输出
3. 编译计算着色器代码
4. 设置计算流水线
5. 向 GPU 批量提交已编码的命令
6. 读取结果矩阵 GPU 缓冲区

GPU 缓冲区创建

为简单起见，矩阵将表示为 点号。第一个元素是行数，第二个元素是 即矩阵的实际列数。

<ph type="x-smartling-placeholder">

</ph>

使用 JavaScript 表示矩阵的简单表示法及采用数学符号的对等函数

这三个 GPU 缓冲区是存储缓冲区，因为我们需要在 计算着色器这解释了为什么 GPU 缓冲区使用标记包括 GPUBufferUsage.STORAGE。结果矩阵使用标志也具有 GPUBufferUsage.COPY_SRC，因为系统会将相应数据复制到另一个缓冲区 所有 GPU 队列命令都执行完毕后读取。

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();


// First Matrix

const firstMatrix = new Float32Array([
  2 /* rows */, 4 /* columns */,
  1, 2, 3, 4,
  5, 6, 7, 8
]);

const gpuBufferFirstMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: firstMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferFirstMatrix = gpuBufferFirstMatrix.getMappedRange();
new Float32Array(arrayBufferFirstMatrix).set(firstMatrix);
gpuBufferFirstMatrix.unmap();


// Second Matrix

const secondMatrix = new Float32Array([
  4 /* rows */, 2 /* columns */,
  1, 2,
  3, 4,
  5, 6,
  7, 8
]);

const gpuBufferSecondMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: secondMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferSecondMatrix = gpuBufferSecondMatrix.getMappedRange();
new Float32Array(arrayBufferSecondMatrix).set(secondMatrix);
gpuBufferSecondMatrix.unmap();


// Result Matrix

const resultMatrixBufferSize = Float32Array.BYTES_PER_ELEMENT * (2 + firstMatrix[0] * secondMatrix[1]);
const resultMatrixBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
});

绑定组布局和绑定组

绑定组布局和绑定组的概念特定于 WebGPU。绑定 组布局定义了着色器预期的输入/输出接口，而 绑定组表示着色器的实际输入/输出数据。

在以下示例中，绑定组布局需要两个只读存储缓冲区，位于 为计算着色器使用带编号的条目绑定 0、1 以及位于 2 的存储缓冲区。 另一方面，为此绑定组布局定义的绑定组 将 GPU 缓冲区发送到以下条目：将 gpuBufferFirstMatrix 发送到绑定 0， 将 gpuBufferSecondMatrix 绑定到绑定 1，将 resultMatrixBuffer 关联到 绑定 2。

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

const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [
    {
      binding: 0,
      resource: {
        buffer: gpuBufferFirstMatrix
      }
    },
    {
      binding: 1,
      resource: {
        buffer: gpuBufferSecondMatrix
      }
    },
    {
      binding: 2,
      resource: {
        buffer: resultMatrixBuffer
      }
    }
  ]
});

计算着色器代码

用于相乘的计算着色器代码用 WGSL（即 WebGPU 着色器语言，可以轻松转换为 SPIR-V。不包含 您会发现下列三个存储缓冲区 与 var<storage> 集成。该程序会将 firstMatrix 和 secondMatrix 用作 resultMatrix 作为其输出。

请注意，每个存储缓冲区都有一个使用的 binding 装饰，对应于 绑定组布局和上面声明的绑定组中定义的相同索引。

const shaderModule = device.createShaderModule({
  code: `
    struct Matrix {
      size : vec2f,
      numbers: array<f32>,
    }

    @group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
    @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
    @group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;

    @compute @workgroup_size(8, 8)
    fn main(@builtin(global_invocation_id) global_id : vec3u) {
      // Guard against out-of-bounds work group sizes
      if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
        return;
      }

      resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);

      let resultCell = vec2(global_id.x, global_id.y);
      var result = 0.0;
      for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
        let a = i + resultCell.x * u32(firstMatrix.size.y);
        let b = resultCell.y + i * u32(secondMatrix.size.y);
        result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
      }

      let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
      resultMatrix.numbers[index] = result;
    }
  `
});

流水线设置

计算流水线是用于描述计算操作 示例。通过调用 device.createComputePipeline() 进行创建。 它需要两个参数：我们之前创建的绑定组布局，以及 定义计算着色器（main WGSL 函数）入口点的阶段 以及使用 device.createShaderModule() 创建的实际计算着色器模块。

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

命令提交

在将绑定组实例化后，我们使用了三个 GPU 缓冲区和一个计算资源 具有绑定组布局的流水线，是时候开始使用它们了。

我们来启动一个可编程计算传递编码器， commandEncoder.beginComputePass()。我们将使用它对 GPU 命令进行编码 执行矩阵乘法的命令。使用 passEncoder.setPipeline(computePipeline) 及其在索引 0 处的绑定组， passEncoder.setBindGroup(0, bindGroup)。索引 0 对应于 WGSL 代码中的 group(0) 装饰。

现在，我们来谈谈此计算着色器如何在 GPU 上运行。我们的 目标是对结果矩阵的每个单元并行执行此程序， 逐步学习。例如，对于大小为 16x32 的结果矩阵， 在 @workgroup_size(8, 8) 上调用执行命令， passEncoder.dispatchWorkgroups(2, 4) 或 passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)。 第一个参数“x”是第一个维度，第二个维度“y”是第二个维度， 最新的一个“z”是第三个维度，默认为 1，因为我们在这里不需要。 在 GPU 计算领域，对一组数据执行内核函数的命令称为调度。

<ph type="x-smartling-placeholder">

</ph>

针对每个结果矩阵单元并行执行

在 WGSL 中，计算着色器的工作组网格大小为 (8, 8) 代码。因此，“x”和“y”它们分别是 第一个矩阵和第二个矩阵的列数将相除 乘以 8。这样，我们现在可以使用 passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)。通过 要运行的工作组网格数是 dispatchWorkgroups() 参数。

如上图所示，每个着色器都可以访问一个唯一的 builtin(global_invocation_id) 对象，用于了解哪个结果 要计算的矩阵单元格。

const commandEncoder = device.createCommandEncoder();

const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const workgroupCountX = Math.ceil(firstMatrix[0] / 8);
const workgroupCountY = Math.ceil(secondMatrix[1] / 8);
passEncoder.dispatchWorkgroups(workgroupCountX, workgroupCountY);
passEncoder.end();

如需结束计算传递编码器，请调用 passEncoder.end()。然后，创建一个 用作复制结果矩阵缓冲区的目标的 GPU 缓冲区 copyBufferToBuffer。最后，使用以下代码 copyEncoder.finish()，并通过调用这些命令将其提交到 GPU 设备队列 device.queue.submit() 替换为 GPU 命令。

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
  resultMatrixBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  resultMatrixBufferSize /* size */
);

// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
device.queue.submit([gpuCommands]);

读取结果矩阵

读取结果矩阵就像使用以下代码调用 gpuReadBuffer.mapAsync() 一样简单： GPUMapMode.READ 并等待返回的 promise 来解析，这表明 GPU 缓冲区现在已映射此时，可获取 范围为 gpuReadBuffer.getMappedRange()。

<ph type="x-smartling-placeholder">

</ph>

矩阵乘法结果

在我们的代码中，DevTools JavaScript 控制台中记录的结果是“2, 2, 50, 60, 114, 140 英寸。

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));

恭喜！顺利到达。您可以播放相应示例。

最后一个技巧

为了使代码更易于阅读，一种方法是使用方便的 计算流水线的 getBindGroupLayout 方法，用于推断绑定组 布局。这个技巧无需创建 自定义绑定组布局，并在计算中指定流水线布局 如下图所示。

上一个示例的 getBindGroupLayout 插图可供使用。

 const computePipeline = device.createComputePipeline({
-  layout: device.createPipelineLayout({
-    bindGroupLayouts: [bindGroupLayout]
-  }),
   compute: {

-// Bind group layout and bind group
- const bindGroupLayout = device.createBindGroupLayout({
-   entries: [
-     {
-       binding: 0,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 1,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 2,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "storage"
-       }
-     }
-   ]
- });
+// Bind group
  const bindGroup = device.createBindGroup({
-  layout: bindGroupLayout,
+  layout: computePipeline.getBindGroupLayout(0 /* index */),
   entries: [

性能发现结果

那么，在 GPU 上运行矩阵乘法与在 GPU 上运行矩阵乘法相比有何不同： CPU？为了弄清楚这一点，我编写了刚刚针对 CPU 描述的程序。您可以 如下图所示，全部使用 GPU 似乎是理所当然的选择 当矩阵大小大于 256 x 256 时触发。

<ph type="x-smartling-placeholder">

</ph>

GPU 与 CPU 基准测试

本文才刚刚开始探索 WebGPU。更多精彩功能 我们将很快发布几篇文章，深入探讨 GPU 计算，以及如何渲染 （画布、纹理、采样器）适用于 WebGPU。