開始在網路上使用 GPU Compute

本文透過範例探索實驗性 WebGPU API，並協助 一開始先使用 GPU 進行資料平行運算

François Beaufort
François Beaufort
GitHub

背景

如您所知，圖形處理器 (GPU) 是電子版處理器， 電腦中最初需要處理的子系統 圖像。然而，在過去 10 年來已進化為更靈活的選擇 這個架構可讓開發人員執行多種類型的演算法，而不只是 會轉譯 3D 圖形，並利用 例如 GPU這些功能稱為「GPU Compute」 適用於一般用途科學運算的輔助處理器稱為「一般用途」 GPU (GPGPU) 程式設計。

GPU 運算為最近的機器學習蓬勃發展帶來重大影響 而卷積類神經網路和其他模型則可利用 以便在 GPU 上更有效率地執行架構使用目前的網路平台 缺少 GPU 運算功能，W3C 的「GPU for the Web」社群群組 設計 API 來公開 。這個 API 稱為 WebGPU。

WebGPU 是低階 API，例如 WebGL。這個 API 的功能相當強大而且詳細 但沒關係。我們要注意的是效能。

在本文中，我將著重介紹 WebGPU 的 GPU 運算部分， 其實我還在刮掉表面，讓你可以開始在 我會進一步深入探討 WebGPU 轉譯 (畫布、紋理、 等)。

存取 GPU

使用 WebGPU 可輕鬆存取 GPU。正在撥打 navigator.gpu.requestAdapter() 會傳回將使用 GPU 以非同步方式解析的 JavaScript 承諾 轉接頭。您可以將這個轉接器視為圖形卡。可以整合 (採用與 CPU 相同的晶片) 或離散 (通常是 PCIe 卡， 但會提高效能)。

取得 GPU 轉接器後，請呼叫 adapter.requestDevice() 來取得承諾 ，根據您用於執行部分 GPU 運算的 GPU 裝置解析，

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

這兩個函式都會提供選項，讓您可以更具體地指定 轉接器 (電源偏好設定) 以及要使用的裝置 (擴充功能、限制)。對於 為了方便起見，我們會使用本文中的預設選項。

寫入緩衝區記憶體

讓我們瞭解如何使用 JavaScript 將資料寫入 GPU 的記憶體。這個 新式網頁採用的沙箱模型，因此程序並不簡單 。

以下範例說明如何編寫四個位元組，來緩衝可存取的記憶體 從 GPU 擷取出來這會呼叫 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 指令 以非同步方式載入物件因此，我們建立並傳送許多 GPU 指令清單 執行批次作業在 WebGPU 中，GPU 命令編碼器 device.createCommandEncoder() 是 JavaScript 物件，用來建構批次 「緩衝處理」會在某個時間點向 GPU 發送的指令 另一方面，GPUBuffer 屬於「未緩衝處理」，表示它們能以不可分割的形式執行 系統在呼叫時

取得 GPU 指令編碼器後，請呼叫 copyEncoder.copyBufferToBuffer() 如下所示，將此指令新增至指令佇列，以供日後執行。 最後，呼叫 copyEncoder.finish() 來完成編碼指令並提交 新增至 GPU 裝置指令佇列佇列會負責處理 透過 device.queue.submit() 完成提交，並使用 GPU 指令做為引數。 這會依序執行陣列中儲存的所有指令。

// 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。它會傳回會在 GPU 緩衝區發生時解析的承諾 對應。然後使用 gpuReadBuffer.getMappedRange() 取得對應的範圍 包含所有已排入佇列的 GPU 指令後，與第一個 GPU 緩衝區相同的值 是否已執行。

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

您可以試用此範例。

簡單來說，以下是緩衝區記憶體作業的重要須知：

  • GPU 緩衝區必須未對應，才能用於提交裝置佇列。
  • 對應時，可讀取及寫入 GPU 緩衝區。
  • 套用 mapAsync() 和 createBuffer() 時，GPU 緩衝區會對應於 會呼叫設為 true 的 mappedAtCreation。

著色器程式設計

在 GPU 上執行的程式只會執行運算 (且不繪製) 稱為運算著色器。這些模型會由數百個 雙層運作才能處理的 GPU 核心數 (小於 CPU 核心) 資料。其輸入和輸出內容是 WebGPU 中的緩衝區。

為了說明在 WebGPU 中使用運算著色器的方式，我們將操作矩陣 乘法是機器學習的通用演算法，如下圖所示。

矩陣乘法圖表
Matrix 乘法圖表

簡單來說，我們會執行以下作業：

  1. 建立三個 GPU 緩衝區 (兩個用來相乘的矩陣，另一個則用於 結果矩陣)
  2. 說明運算著色器的輸入和輸出內容
  3. 編譯運算著色器程式碼
  4. 設定運算管道
  5. 將已編碼的指令批次提交至 GPU
  6. 讀取結果矩陣 GPU 緩衝區

建立 GPU 緩衝區

為求簡單，矩陣會以浮動清單表示 重點。第一個元素是列數，第二個元素 而其餘部分則是矩陣的實際數量。

以 JavaScript 形式呈現矩陣，其數學標記法中相等的矩陣
以 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。繫結 群組版面配置會定義著色器預期的輸入/輸出介面，而 bind group 代表著色器的實際輸入/輸出資料。

在下方範例中，繫結群組版面配置預期在下列位置有兩個唯讀儲存空間緩衝區： 針對運算著色器，有編號項目繫結 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 Shader Language 可適當翻譯為 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 上執行。我們的 目標是在結果矩陣的每個儲存格上平行執行此程式， 。例如，對大小為 16 x 32 的結果矩陣進行編碼時 而在 @workgroup_size(8, 8) 上呼叫 passEncoder.dispatchWorkgroups(2, 4) 或 passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)。 第一個引數「x」是第一個維度，第二個為「y」是第二個維度 和最新的「z」是預設為 1 的第三個維度，因為此處用不到。 在 GPU 運算環境中，對一組用於執行核心函式的指令編碼稱為「調度」。

同時對每個結果矩陣儲存格平行執行
針對每個結果矩陣儲存格並行執行

運算著色器的工作群組格線大小在 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()，然後呼叫 將 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，並等待傳回的承諾解決問題， 就能對應 GPU 緩衝區此時， 範圍為 gpuReadBuffer.getMappedRange()。

矩陣乘法結果
Matrix 乘法結果

在我們的程式碼中，開發人員工具 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 時

GPU 與 CPU 基準的比較
GPU 與 CPU 基準的比較

本文只是我開始探索 WebGPU 的第一步。預期會增加 文章即將深入探討 GPU 運算和轉譯方式 (畫布、紋理、取樣器) 支援 WebGPU。