Halaman ini diterjemahkan oleh Cloud Translation API.

Mulai menggunakan Komputasi GPU di web

Postingan ini mengeksplorasi WebGPU API eksperimental melalui contoh dan membantu Anda memulai komputasi paralel data menggunakan GPU.

François Beaufort

Latar belakang

Seperti yang mungkin sudah Anda ketahui, Unit Pemrosesan Grafis (GPU) adalah subsistem elektronik dalam komputer yang awalnya dikhususkan untuk memproses grafis. Namun, dalam 10 tahun terakhir, GPU telah berkembang menjadi arsitektur yang lebih fleksibel, sehingga developer dapat menerapkan banyak jenis algoritma, bukan hanya merender grafik 3D, sekaligus memanfaatkan arsitektur unik GPU. Kemampuan ini disebut GPU Compute, dan penggunaan GPU sebagai koprosesor untuk komputasi ilmiah tujuan umum disebut pemrograman GPU umum (GPGPU).

GPU Compute telah berkontribusi secara signifikan pada booming machine learning baru-baru ini, karena jaringan saraf konvolusi dan model lainnya dapat memanfaatkan arsitektur untuk berjalan lebih efisien di GPU. Dengan Platform Web saat ini yang tidak memiliki kemampuan GPU Compute, Grup Komunitas "GPU for the Web" W3C mendesain API untuk mengekspos API GPU modern yang tersedia di sebagian besar perangkat saat ini. API ini disebut WebGPU.

WebGPU adalah API level rendah, seperti WebGL. Alat ini sangat canggih dan cukup panjang, seperti yang akan Anda lihat. Tapi tidak apa-apa. Yang kita cari adalah performa.

Dalam artikel ini, saya akan berfokus pada bagian GPU Compute dari WebGPU dan, sejujurnya, saya hanya membahas permukaannya, sehingga Anda dapat mulai bermain sendiri. Saya akan membahas lebih dalam dan membahas rendering WebGPU (kanvas, tekstur, dll.) dalam artikel mendatang.

Mengakses GPU

Mengakses GPU sangat mudah di WebGPU. Memanggil navigator.gpu.requestAdapter() akan menampilkan promise JavaScript yang akan diselesaikan secara asinkron dengan adaptor GPU. Anggap adaptor ini sebagai kartu grafis. GPU dapat terintegrasi (di chip yang sama dengan CPU) atau terpisah (biasanya kartu PCIe yang lebih berperforma tetapi menggunakan lebih banyak daya).

Setelah Anda memiliki adaptor GPU, panggil adapter.requestDevice() untuk mendapatkan promise yang akan di-resolve dengan perangkat GPU yang akan Anda gunakan untuk melakukan beberapa komputasi GPU.

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

Kedua fungsi tersebut menggunakan opsi yang memungkinkan Anda menentukan jenis adaptor (preferensi daya) dan perangkat (ekstensi, batas) yang Anda inginkan. Untuk memudahkan, kita akan menggunakan opsi default dalam artikel ini.

Menulis memori buffer

Mari kita lihat cara menggunakan JavaScript untuk menulis data ke memori untuk GPU. Proses ini tidak mudah karena model sandboxing yang digunakan di browser web modern.

Contoh di bawah menunjukkan cara menulis empat byte ke buffering memori yang dapat diakses dari GPU. Fungsi ini memanggil device.createBuffer() yang mengambil ukuran buffer dan penggunaannya. Meskipun flag penggunaan GPUBufferUsage.MAP_WRITE tidak diperlukan untuk panggilan khusus ini, mari kita nyatakan secara eksplisit bahwa kita ingin menulis ke buffering ini. Hal ini menghasilkan objek buffering GPU yang dipetakan saat pembuatan berkat mappedAtCreation yang ditetapkan ke true. Kemudian, buffer data biner mentah terkait dapat diambil dengan memanggil metode buffer GPU getMappedRange().

Menulis byte sudah biasa jika Anda sudah bermain dengan ArrayBuffer; gunakan TypedArray dan salin nilai ke dalamnya.

// 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]);

Pada tahap ini, buffer GPU dipetakan, yang berarti buffer tersebut dimiliki oleh CPU, dan dapat diakses dalam operasi baca/tulis dari JavaScript. Agar GPU dapat mengaksesnya, GPU harus dihapus pemetaan yang semudah memanggil gpuBuffer.unmap().

Konsep dipetakan/tidak dipetakan diperlukan untuk mencegah kondisi perlombaan saat GPU dan CPU mengakses memori secara bersamaan.

Membaca memori buffer

Sekarang, mari kita lihat cara menyalin buffer GPU ke buffer GPU lain dan membacanya kembali.

Karena kita menulis di buffer GPU pertama dan ingin menyalinnya ke buffer GPU kedua, tanda penggunaan baru GPUBufferUsage.COPY_SRC diperlukan. Buffer GPU kedua dibuat dalam status yang tidak dipetakan kali ini dengan device.createBuffer(). Flag penggunaannya adalah GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ karena akan digunakan sebagai tujuan buffering GPU pertama dan dibaca dalam JavaScript setelah perintah salin GPU dieksekusi.

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

Karena GPU adalah koprosesor independen, semua perintah GPU dijalankan secara asinkron. Inilah sebabnya ada daftar perintah GPU yang dibuat dan dikirim dalam batch jika diperlukan. Di WebGPU, encoder perintah GPU yang ditampilkan oleh device.createCommandEncoder() adalah objek JavaScript yang membuat batch perintah "dibuffer" yang akan dikirim ke GPU pada suatu waktu. Di sisi lain, metode di GPUBuffer bersifat "tanpa buffering", yang berarti metode tersebut dieksekusi secara atomik saat dipanggil.

Setelah Anda memiliki encoder perintah GPU, panggil copyEncoder.copyBufferToBuffer() seperti yang ditunjukkan di bawah untuk menambahkan perintah ini ke antrean perintah untuk dieksekusi nanti. Terakhir, selesaikan perintah encoding dengan memanggil copyEncoder.finish() dan kirimkan perintah tersebut ke antrean perintah perangkat GPU. Antrean bertanggung jawab untuk menangani pengiriman yang dilakukan melalui device.queue.submit() dengan perintah GPU sebagai argumen. Tindakan ini akan mengeksekusi semua perintah yang disimpan dalam array secara berurutan.

// 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]);

Pada tahap ini, perintah antrean GPU telah dikirim, tetapi tidak harus dieksekusi. Untuk membaca buffer GPU kedua, panggil gpuReadBuffer.mapAsync() dengan GPUMapMode.READ. Fungsi ini menampilkan promise yang akan diselesaikan saat buffer GPU dipetakan. Kemudian, dapatkan rentang yang dipetakan dengan gpuReadBuffer.getMappedRange() yang berisi nilai yang sama dengan buffering GPU pertama setelah semua perintah GPU dalam antrean telah dieksekusi.

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

Anda dapat mencoba contoh ini.

Singkatnya, berikut hal yang perlu diingat terkait operasi memori buffering:

Buffer GPU harus dibatalkan pemetaan agar dapat digunakan dalam pengiriman antrean perangkat.
Saat dipetakan, buffer GPU dapat dibaca dan ditulis dalam JavaScript.
Buffer GPU dipetakan saat mapAsync() dan createBuffer() dengan mappedAtCreation ditetapkan ke true dipanggil.

Pemrograman shader

Program yang berjalan di GPU yang hanya melakukan komputasi (dan tidak menggambar segitiga) disebut compute shader. Tugas ini dijalankan secara paralel oleh ratusan core GPU (yang lebih kecil dari core CPU) yang beroperasi bersama untuk memproses data. Input dan outputnya adalah buffering di WebGPU.

Untuk mengilustrasikan penggunaan compute shader di WebGPU, kita akan bermain dengan perkalian matriks, algoritma umum dalam machine learning yang diilustrasikan di bawah.

Singkatnya, berikut yang akan kita lakukan:

Buat tiga buffering GPU (dua untuk matriks yang akan dikalikan dan satu untuk matriks hasil)
Menjelaskan input dan output untuk shader komputasi
Mengompilasi kode compute shader
Menyiapkan pipeline komputasi
Kirim perintah yang dienkode ke GPU dalam batch
Membaca buffer GPU matriks hasil

Pembuatan Buffer GPU

Untuk mempermudah, matriks akan direpresentasikan sebagai daftar bilangan floating point. Elemen pertama adalah jumlah baris, elemen kedua adalah jumlah kolom, dan sisanya adalah angka sebenarnya dari matriks.

Representasi sederhana matriks dalam JavaScript dan yang setara dalam notasi matematika

Tiga buffer GPU adalah buffer penyimpanan karena kita perlu menyimpan dan mengambil data di shader komputasi. Hal ini menjelaskan mengapa flag penggunaan buffering GPU menyertakan GPUBufferUsage.STORAGE untuk semuanya. Flag penggunaan matriks hasil juga memiliki GPUBufferUsage.COPY_SRC karena akan disalin ke buffering lain untuk dibaca setelah semua perintah antrean GPU dieksekusi.

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

Tata letak grup pengikatan dan grup pengikatan

Konsep tata letak grup pengikatan dan grup pengikatan khusus untuk WebGPU. Tata letak grup binding menentukan antarmuka input/output yang diharapkan oleh shader, sedangkan grup binding mewakili data input/output sebenarnya untuk shader.

Pada contoh di bawah, tata letak grup pengikatan mengharapkan dua buffer penyimpanan hanya baca di binding entri bernomor 0, 1, dan buffer penyimpanan di 2 untuk shader komputasi. Di sisi lain, grup pengikatan, yang ditentukan untuk tata letak grup pengikatan ini, mengaitkan buffer GPU ke entri: gpuBufferFirstMatrix ke pengikatan 0, gpuBufferSecondMatrix ke pengikatan 1, dan resultMatrixBuffer ke pengikatan 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
      }
    }
  ]
});

Kode shader komputasi

Kode shader komputasi untuk mengalikan matriks ditulis dalam WGSL, WebGPU Shader Language, yang dapat diterjemahkan secara mudah ke SPIR-V. Tanpa membahas detailnya, Anda akan menemukan tiga buffering penyimpanan di bawah yang diidentifikasi dengan var<storage>. Program ini akan menggunakan firstMatrix dan secondMatrix sebagai input dan resultMatrix sebagai output-nya.

Perhatikan bahwa setiap buffering penyimpanan memiliki dekorasi binding yang digunakan yang sesuai dengan indeks yang sama yang ditentukan dalam tata letak grup pengikatan dan grup pengikatan yang dideklarasikan di atas.

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

Penyiapan pipeline

Pipeline komputasi adalah objek yang benar-benar menjelaskan operasi komputasi yang akan kita lakukan. Buat file tersebut dengan memanggil device.createComputePipeline(). Fungsi ini memerlukan dua argumen: tata letak grup pengikatan yang kita buat sebelumnya, dan tahap compute yang menentukan titik entri compute shader kita (fungsi WGSL main) dan modul compute shader sebenarnya yang dibuat dengan device.createShaderModule().

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

Pengiriman perintah

Setelah membuat instance grup pengikatan dengan tiga buffering GPU dan pipeline komputasi dengan tata letak grup pengikatan, saatnya menggunakannya.

Mari kita mulai encoder kartu komputasi yang dapat diprogram dengan commandEncoder.beginComputePass(). Kita akan menggunakannya untuk mengenkode perintah GPU yang akan melakukan perkalian matriks. Tetapkan pipeline-nya dengan passEncoder.setPipeline(computePipeline) dan grup pengikatannya pada indeks 0 dengan passEncoder.setBindGroup(0, bindGroup). Indeks 0 sesuai dengan dekorasi group(0) dalam kode WGSL.

Sekarang, mari kita bahas cara shader komputasi ini akan berjalan di GPU. Sasaran kami adalah mengeksekusi program ini secara paralel untuk setiap sel matriks hasil, langkah demi langkah. Misalnya, untuk matriks hasil berukuran 16x32, untuk mengenkode perintah eksekusi, pada @workgroup_size(8, 8), kita akan memanggil passEncoder.dispatchWorkgroups(2, 4) atau passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Argumen pertama "x" adalah dimensi pertama, argumen kedua "y" adalah dimensi kedua, dan argumen terakhir "z" adalah dimensi ketiga yang ditetapkan secara default ke 1 karena kita tidak memerlukannya di sini. Dalam dunia komputasi GPU, mengenkode perintah untuk mengeksekusi fungsi kernel pada kumpulan data disebut pengiriman.

Eksekusi secara paralel untuk setiap sel matriks hasil

Ukuran petak grup kerja untuk shader komputasi adalah (8, 8) dalam kode WGSL kami. Oleh karena itu, "x" dan "y" yang masing-masing merupakan jumlah baris matriks pertama dan jumlah kolom matriks kedua akan dibagi dengan 8. Dengan demikian, kita sekarang dapat mengirim panggilan komputasi dengan passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). Jumlah petak grup kerja yang akan dijalankan adalah argumen dispatchWorkgroups().

Seperti yang terlihat pada gambar di atas, setiap shader akan memiliki akses ke objek builtin(global_invocation_id) unik yang akan digunakan untuk mengetahui sel matriks hasil yang akan dihitung.

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();

Untuk mengakhiri encoder kartu komputasi, panggil passEncoder.end(). Kemudian, buat buffer GPU untuk digunakan sebagai tujuan untuk menyalin buffer matriks hasil dengan copyBufferToBuffer. Terakhir, selesaikan perintah encoding dengan copyEncoder.finish() dan kirimkan ke antrean perangkat GPU dengan memanggil device.queue.submit() dengan perintah 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]);

Membaca matriks hasil

Membaca matriks hasil semudah memanggil gpuReadBuffer.mapAsync() dengan GPUMapMode.READ dan menunggu promise yang ditampilkan untuk diselesaikan yang menunjukkan buffer GPU kini dipetakan. Pada tahap ini, Anda bisa mendapatkan rentang yang dipetakan dengan gpuReadBuffer.getMappedRange().

Dalam kode kita, hasil yang dicatat di konsol JavaScript DevTools adalah "2, 2, 50, 60, 114, 140".

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

Selamat! Anda berhasil. Anda dapat bermain dengan sampel.

Trik terakhir

Salah satu cara untuk membuat kode Anda lebih mudah dibaca adalah dengan menggunakan metode getBindGroupLayout yang praktis dari pipeline komputasi untuk menyimpulkan tata letak grup pengikatan dari modul shader. Trik ini menghilangkan kebutuhan untuk membuat tata letak grup pengikatan kustom dan menentukan tata letak pipeline di compute pipeline seperti yang dapat Anda lihat di bawah.

Ilustrasi getBindGroupLayout untuk contoh sebelumnya tersedia.

 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: [

Temuan performa

Jadi, bagaimana perbandingan antara menjalankan perkalian matriks di GPU dengan menjalankannya di CPU? Untuk mengetahuinya, saya menulis program yang baru saja dijelaskan untuk CPU. Dan seperti yang dapat Anda lihat pada grafik di bawah, menggunakan kemampuan penuh GPU tampaknya merupakan pilihan yang jelas saat ukuran matriks lebih besar dari 256x256.

Artikel ini hanyalah awal dari perjalanan saya dalam menjelajahi WebGPU. Nantikan artikel lainnya yang akan segera menampilkan lebih banyak pembahasan mendalam tentang GPU Compute dan cara kerja rendering (kanvas, tekstur, sampler) di WebGPU.