Mulai menggunakan Komputasi GPU di web

Postingan ini membahas WebGPU API eksperimental melalui contoh dan membantu Anda mulai melakukan komputasi paralel data menggunakan GPU.

François Beaufort
François Beaufort
GitHub

Latar belakang

Seperti yang mungkin sudah Anda ketahui, Graphic Processing Unit (GPU) adalah server subsistem dalam komputer yang awalnya khusus untuk pemrosesan grafis. Namun dalam 10 tahun terakhir, sistem ini telah berkembang menjadi lebih fleksibel arsitektur yang memungkinkan pengembang untuk mengimplementasikan banyak jenis algoritma, tidak hanya membuat grafik 3D, sambil memanfaatkan arsitektur unik oleh GPU. Kemampuan ini disebut sebagai Komputasi GPU, dan menggunakan GPU sebagai koprosesor untuk komputasi ilmiah tujuan umum disebut Pemrograman GPU (GPGPU).

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

WebGPU adalah API tingkat rendah, seperti WebGL. Cara ini sangat ampuh dan cukup panjang, karena yang akan Anda lihat. Tapi tidak apa-apa. Yang kami cari adalah performa.

Dalam artikel ini, saya akan fokus pada bagian Komputasi GPU dari WebGPU dan, jujur, saya hanya membersihkan permukaannya, agar Anda bisa mulai bermain sendiri. Saya akan membahas lebih dalam dan membahas rendering WebGPU (kanvas, tekstur, dll.) dalam artikel mendatang.

Mengakses GPU

Mengakses GPU itu mudah di WebGPU. Menelepon navigator.gpu.requestAdapter() menampilkan promise JavaScript yang akan di-resolve secara asinkron dengan GPU {i>adaptor<i}. Bayangkan adaptor ini sebagai kartu grafis. Bisa diintegrasikan (pada chip yang sama dengan CPU) atau diskrit (biasanya kartu PCIe dengan ukuran berperforma tinggi tetapi menggunakan lebih banyak daya).

Setelah Anda memiliki adaptor GPU, panggil adapter.requestDevice() untuk mendapatkan promise yang akan diselesaikan 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 {i>function<i} mengambil opsi yang memungkinkan Anda untuk spesifik tentang jenis adaptor (preferensi daya) dan perangkat (ekstensi, batas) yang diinginkan. Untuk demi kesederhanaan, kami akan menggunakan opsi {i>default<i} dalam artikel ini.

Menulis memori buffer

Mari kita lihat cara menggunakan JavaScript untuk menulis data ke memori untuk GPU. Ini prosesnya tidak mudah karena model {i>sandbox <i}yang digunakan di web modern browser.

Contoh di bawah ini menunjukkan cara menulis empat byte untuk buffer akses memori dari GPU. Fungsi ini memanggil device.createBuffer() yang menggunakan ukuran {i>buffer <i}dan penggunaannya. Meskipun tanda penggunaan GPUBufferUsage.MAP_WRITE adalah tidak diperlukan untuk panggilan khusus ini, mari kita masukkan ke {i>buffer<i} ini. Ini menghasilkan objek {i>buffer <i}GPU yang dipetakan pada saat pembuatan berkat mappedAtCreation ditetapkan ke benar (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, lalu salin nilainya.

// 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, artinya buffer tersebut dimiliki oleh CPU, dan dapat diakses dengan cara membaca/menulis dari JavaScript. Agar GPU dapat mengaksesnya, GPU harus tidak dipetakan, yang semudah memanggil gpuBuffer.unmap().

Konsep dipetakan/tidak dipetakan diperlukan untuk mencegah kondisi race saat GPU dan CPU akses 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 Buffering GPU, tanda penggunaan baru GPUBufferUsage.COPY_SRC diperlukan. Yang kedua Kali ini, buffer GPU dibuat dalam keadaan tidak dipetakan device.createBuffer(). Tanda penggunaannya adalah GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ karena akan digunakan sebagai tujuan GPU pertama buffer dan membaca dalam JavaScript setelah perintah penyalinan GPU telah 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 dieksekusi secara asinkron. Inilah sebabnya mengapa ada daftar perintah GPU yang dibangun dan dikirim di batch saat diperlukan. Di WebGPU, encoder perintah GPU yang ditampilkan oleh device.createCommandEncoder() adalah objek JavaScript yang membuat batch "di-buffer" perintah yang akan dikirim ke GPU. Metode pada Di sisi lain, GPUBuffer adalah "unbuffered", yang berarti mengeksekusi secara atomik pada saat mereka dipanggil.

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

// 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 belum tentu dieksekusi. Untuk membaca buffer GPU kedua, panggil gpuReadBuffer.mapAsync() dengan GPUMapMode.READ. Metode ini menampilkan promise yang akan diselesaikan saat buffer GPU dipetakan. Kemudian dapatkan rentang yang dipetakan dengan gpuReadBuffer.getMappedRange() yang berisi nilai yang sama dengan buffer GPU pertama setelah semua perintah GPU yang diantrekan 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 adalah hal yang perlu Anda ingat mengenai operasi memori buffer:

  • Buffer GPU harus tidak dipetakan agar dapat digunakan dalam pengiriman antrean perangkat.
  • Jika dipetakan, buffer GPU dapat dibaca dan ditulis dalam JavaScript.
  • Buffer GPU dipetakan saat mapAsync() dan createBuffer() dengan mappedAtCreation yang disetel ke benar (true) akan dipanggil.

Pemrograman shader

Program yang berjalan di GPU yang hanya melakukan komputasi (dan tidak menggambar segitiga) disebut shader komputasi. Operasi ini dijalankan secara paralel dengan ratusan core GPU (yang lebih kecil dari core CPU) yang beroperasi bersama untuk crunch layanan otomatis dan data skalabel. Input dan output-nya adalah buffer di WebGPU.

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

Diagram perkalian matriks
Diagram perkalian matriks

Singkatnya, inilah yang akan kita lakukan:

  1. Buat tiga buffer GPU (dua untuk matriks yang akan dikalikan dan satu untuk buffer matriks hasil)
  2. Menjelaskan input dan output untuk shader komputasi
  3. Mengompilasi kode shader komputasi
  4. Menyiapkan pipeline komputasi
  5. Kirimkan perintah yang dienkode ke GPU secara berkelompok
  6. Membaca buffer GPU matriks hasil

Pembuatan Buffer GPU

Demi kemudahan, matriks akan ditampilkan sebagai daftar angka poin. Elemen pertama adalah jumlah baris, elemen kedua adalah jumlah kolom, dan sisanya adalah bilangan aktual dari matriks.

Representasi sederhana matriks di JavaScript dan yang setara dalam notasi matematika
Representasi sederhana matriks di JavaScript dan yang setara dalam notasi matematika

Tiga buffer GPU adalah buffer penyimpanan karena kita perlu menyimpan dan mengambil shader komputasi. Ini menjelaskan mengapa flag penggunaan buffer GPU menyertakan GPUBufferUsage.STORAGE untuk semuanya. Penanda penggunaan matriks hasil juga memiliki GPUBufferUsage.COPY_SRC karena akan disalin ke buffer lain untuk membaca setelah semua perintah antrean GPU semuanya telah 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 Bind group dan bind group

Konsep tata letak bind group dan bind group bersifat khusus untuk WebGPU. Pemberkasan tata letak grup mendefinisikan antarmuka input/output yang diharapkan oleh shader, sementara bind group mewakili data input/output aktual untuk shader.

Dalam contoh di bawah ini, tata letak bind group mengharapkan dua buffer penyimpanan hanya baca binding entri bernomor 0, 1, dan buffer penyimpanan di 2 untuk shader komputasi. Di sisi lain, bind group, yang didefinisikan untuk tata letak bind group ini, mengaitkan Buffering GPU ke entri: gpuBufferFirstMatrix ke 0 binding, gpuBufferSecondMatrix ke 1 binding, dan resultMatrixBuffer ke mengikat 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, Bahasa Shader WebGPU, yang mudah diterjemahkan ke SPIR-V. Tanpa secara mendetail, Anda akan menemukan tiga buffer penyimpanan di bawah ini, dengan var<storage>. Program ini akan menggunakan firstMatrix dan secondMatrix sebagai input dan resultMatrix sebagai output-nya.

Perhatikan bahwa setiap buffer penyimpanan memiliki dekorasi binding yang digunakan sesuai dengan indeks yang sama yang ditentukan dalam tata letak bind group dan bind group 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 dengan memanggil device.createComputePipeline(). Dibutuhkan dua argumen: tata letak bind group yang telah kita buat sebelumnya, dan komputasi Tahap yang menentukan titik entri shader komputasi kita (fungsi main WGSL) dan modul shader komputasi aktual yang dibuat dengan device.createShaderModule().

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

Pengiriman perintah

Setelah membuat instance bind group dengan tiga buffer GPU dan pipeline dengan tata letak {i>bind group<i}, sekarang saatnya untuk menggunakannya.

Mari kita mulai encoder compute pass yang dapat diprogram dengan commandEncoder.beginComputePass(). Kita akan menggunakannya untuk mengenkode perintah GPU yang akan melakukan perkalian matriks. Menetapkan pipelinenya dengan passEncoder.setPipeline(computePipeline) dan bind group-nya pada indeks 0 dengan passEncoder.setBindGroup(0, bindGroup). Indeks 0 sesuai dengan Dekorasi group(0) dalam kode WGSL.

Sekarang, mari kita bahas bagaimana shader komputasi ini akan berjalan di GPU. Dengan tujuannya adalah menjalankan program ini secara paralel untuk setiap sel dari matriks hasil, langkah demi langkah. Misalnya, untuk matriks hasil berukuran 16 x 32, perintah eksekusi, di @workgroup_size(8, 8), kita akan memanggil passEncoder.dispatchWorkgroups(2, 4) atau passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Argumen pertama "x" adalah dimensi pertama, dimensi kedua adalah "y" adalah dimensi kedua, dan yang terakhir "z" adalah dimensi ketiga yang nilai defaultnya adalah 1 karena kita tidak membutuhkannya di sini. Dalam dunia komputasi GPU, pengkodean perintah untuk mengeksekusi fungsi {i>kernel<i} pada sekumpulan data disebut pengiriman.

Eksekusi secara paralel untuk setiap sel matriks hasil
Eksekusi secara paralel untuk setiap sel matriks hasil

Ukuran petak workgroup untuk shader komputasi adalah (8, 8) dalam WGSL pada kode sumber. Karena itu, "x" dan "y" yang masing-masing merupakan jumlah baris matriks pertama dan jumlah kolom matriks kedua akan dibagi sebelum 8 tahun. Dengan itu, kita sekarang bisa mengirim panggilan komputasi menggunakan passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). Tujuan jumlah petak workgroup yang akan dijalankan adalah argumen dispatchWorkgroups().

Seperti yang terlihat pada gambar di atas, setiap shader akan memiliki akses ke Objek builtin(global_invocation_id) yang akan digunakan untuk mengetahui hasil mana sel matriks untuk 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 penerusan komputasi, panggil passEncoder.end(). Kemudian, buat Buffering GPU yang akan digunakan sebagai tujuan menyalin buffer matriks hasil dengan copyBufferToBuffer. Akhirnya, selesaikan perintah penyandian 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]);

Baca matriks hasil

Membaca matriks hasil semudah memanggil gpuReadBuffer.mapAsync() dengan GPUMapMode.READ dan menunggu promise yang ditampilkan untuk diselesaikan, yang menunjukkan buffer GPU sekarang sudah dipetakan. Pada tahap ini, dimungkinkan untuk memperoleh rentang dengan gpuReadBuffer.getMappedRange().

Hasil perkalian matriks
Hasil perkalian matriks

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-main dengan sampel.

Satu trik terakhir

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

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 menjalankan perkalian matriks di GPU dibandingkan dengan menjalankannya di CPU? Untuk mengetahuinya, saya menulis program yang baru saja dijelaskan untuk CPU. Seperti yang Anda bisa, lihat pada grafik di bawah ini, menggunakan kekuatan penuh GPU sepertinya pilihan yang tepat ketika ukuran matriks lebih besar dari 256 x 256.

Tolok ukur GPU vs CPU
Tolok ukur GPU vs CPU

Artikel ini hanyalah awal dari perjalanan saya mengeksplorasi WebGPU. Dapatkan lebih banyak artikel yang akan segera menampilkan pembahasan mendalam tentang Komputasi GPU dan cara rendering (kanvas, tekstur, sampler) berfungsi dalam WebGPU.