Postingan ini membahas WebGPU API eksperimental melalui contoh dan membantu Anda mulai melakukan komputasi paralel data menggunakan GPU.
Latar belakang
Seperti yang mungkin sudah Anda ketahui, Graphic Processing Unit (GPU) adalah subsistem elektronik dalam komputer yang awalnya dikhususkan untuk memproses grafis. Namun, dalam 10 tahun terakhir, GPU telah berevolusi menjadi arsitektur yang lebih fleksibel yang memungkinkan developer menerapkan berbagai jenis algoritma, tidak hanya merender grafis 3D, sekaligus memanfaatkan arsitektur unik GPU. Kemampuan ini disebut sebagai Komputasi GPU, dan menggunakan GPU sebagai koprosesor untuk komputasi ilmiah tujuan umum disebut pemrograman GPU tujuan umum (GPGPU).
Komputasi GPU telah berkontribusi signifikan terhadap boom machine learning baru-baru ini karena jaringan neural konvolusi dan model lain dapat memanfaatkan arsitektur ini untuk berjalan lebih efisien pada GPU. Karena Platform Web saat ini tidak memiliki kemampuan Komputasi GPU, Grup Komunitas "GPU untuk Web" W3C sedang 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. Sangat andal dan cukup panjang, seperti yang akan Anda lihat. Tapi itu tidak menjadi masalah. Yang kita cari adalah kinerja.
Dalam artikel ini, saya akan fokus pada bagian Komputasi GPU dari WebGPU dan, sejujurnya, saya hanya sedikit memperlihatkannya agar Anda dapat mulai bermain sendiri. Saya akan membahas lebih dalam dan membahas rendering WebGPU (kanvas, tekstur, dll.) di artikel mendatang.
Mengakses GPU
Mengakses GPU itu 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. Kartu ini dapat diintegrasikan
(pada chip yang sama dengan CPU) atau diskret (biasanya kartu PCIe yang lebih
berperforma 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 fungsi mengambil opsi yang memungkinkan Anda untuk lebih spesifik tentang jenis adaptor (preferensi daya) dan perangkat (ekstensi, batas) yang Anda inginkan. Agar mudah, kami 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 sederhana karena model sandbox yang digunakan di browser web modern.
Contoh di bawah ini menunjukkan cara menulis empat byte untuk mem-buffer memori yang dapat diakses dari GPU. Fungsi ini memanggil device.createBuffer()
yang menggunakan ukuran
buffer dan penggunaannya. Meskipun flag penggunaan GPUBufferUsage.MAP_WRITE
tidak diperlukan untuk panggilan khusus ini, mari kita jelaskan bahwa kita ingin menulis ke buffer ini. Tindakan ini menghasilkan objek buffer 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 tidak asing lagi jika Anda sudah pernah menggunakan ArrayBuffer
; gunakan
TypedArray
dan salin nilainya 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 dimiliki oleh CPU, dan
dapat diakses dalam mode baca/tulis dari JavaScript. Agar dapat mengaksesnya, GPU
harus tidak dipetakan. Sesederhana memanggil gpuBuffer.unmap()
.
Konsep pemetaan/yang tidak dipetakan diperlukan untuk mencegah kondisi race 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 dalam buffer GPU pertama dan ingin menyalinnya ke buffering GPU kedua, flag penggunaan baru GPUBufferUsage.COPY_SRC
diperlukan. Kali ini buffer GPU kedua dibuat dalam status yang tidak dipetakan dengan device.createBuffer()
. Tanda penggunaannya adalah GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ
karena akan digunakan sebagai tujuan buffer GPU
pertama dan dibaca 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. Itulah sebabnya ada daftar perintah GPU yang dibuat dan dikirim dalam
batch saat diperlukan. Dalam WebGPU, encoder perintah GPU yang ditampilkan oleh
device.createCommandEncoder()
adalah objek JavaScript yang membuat batch
perintah "di-buffer" yang akan dikirimkan ke GPU pada suatu waktu. Di sisi lain, metode pada GPUBuffer
bersifat "tidak di-buffer", yang berarti metode tersebut berjalan secara atomik pada saat 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 nanti.
Terakhir, selesaikan encoding perintah dengan memanggil copyEncoder.finish()
dan mengirimkannya
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 atomik 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 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
dieksekusi.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Anda dapat mencoba contoh ini.
Singkatnya, inilah yang perlu Anda ingat mengenai operasi memori buffer:
- Buffer GPU harus tidak dipetakan agar dapat digunakan dalam pengiriman antrean perangkat.
- Saat dipetakan, buffer GPU dapat dibaca dan ditulis dalam JavaScript.
- Buffer GPU dipetakan saat
mapAsync()
dancreateBuffer()
denganmappedAtCreation
yang ditetapkan ke benar dipanggil.
Pemrograman shader
Program yang berjalan di GPU yang hanya melakukan komputasi (dan tidak menggambar segitiga) disebut shader komputasi. Proses ini dijalankan secara paralel oleh ratusan core GPU (yang lebih kecil dari inti CPU) yang beroperasi bersama untuk mengolah data. Input dan outputnya merupakan buffer di WebGPU.
Untuk menggambarkan penggunaan shader komputasi dalam WebGPU, kita akan mempelajari perkalian matriks, sebuah algoritma umum dalam machine learning yang diilustrasikan di bawah ini.
Singkatnya, inilah yang akan kita lakukan:
- Buat tiga buffer GPU (dua untuk matriks yang akan dikalikan dan satu untuk matriks hasil)
- Menjelaskan input dan output untuk shader komputasi
- Mengompilasi kode shader komputasi
- Menyiapkan pipeline komputasi
- Mengirimkan perintah yang dienkode ke GPU dalam batch
- Membaca buffer GPU matriks hasil
Pembuatan Buffer GPU
Agar lebih praktis, matriks akan direpresentasikan sebagai daftar bilangan floating point. Elemen pertama adalah jumlah baris, elemen kedua adalah jumlah kolom, dan sisanya adalah jumlah matriks sebenarnya.
Ketiga buffer GPU tersebut merupakan buffer penyimpanan karena kita perlu menyimpan dan mengambil data dalam
shader komputasi. Ini menjelaskan mengapa tanda penggunaan buffer GPU menyertakan
GPUBufferUsage.STORAGE
untuk semuanya. Tanda penggunaan matriks hasil juga memiliki GPUBufferUsage.COPY_SRC
karena akan disalin ke buffer lain untuk dibaca setelah semua perintah antrean GPU 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 khusus untuk WebGPU. Tata letak bind group menentukan antarmuka input/output yang diharapkan oleh shader, sedangkan bind group mewakili data input/output yang sebenarnya untuk shader.
Pada contoh di bawah ini, tata letak bind group mengharapkan dua buffer penyimpanan hanya baca pada
binding entri bernomor 0
, 1
, dan buffer penyimpanan di 2
untuk shader komputasi.
Di sisi lain, bind group, yang ditentukan untuk tata letak bind group ini, mengaitkan
buffer GPU ke entri: gpuBufferFirstMatrix
ke 0
binding,
gpuBufferSecondMatrix
ke 1
binding, dan resultMatrixBuffer
ke
2
binding.
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 mudah diterjemahkan ke SPIR-V. Tanpa
mendetailnya, Anda akan menemukan di bawah tiga buffer penyimpanan yang diidentifikasi
dengan var<storage>
. Program ini akan menggunakan firstMatrix
dan secondMatrix
sebagai
input dan resultMatrix
sebagai outputnya.
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 kita buat sebelumnya, dan tahap
komputasi yang menentukan titik masuk shader komputasi (fungsi WGSL main
)
serta modul shader komputasi 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 bind group dengan tiga buffer GPU dan pipeline komputasi dengan tata letak bind group, sekarang saatnya menggunakannya.
Mari kita mulai encoder penerusan komputasi yang dapat diprogram dengan
commandEncoder.beginComputePass()
. Kita akan menggunakannya untuk mengenkode perintah GPU
yang akan melakukan perkalian matriks. Tetapkan 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. Tujuan
kita adalah untuk menjalankan 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, yang kedua "y" adalah dimensi kedua,
dan yang terbaru "z" adalah dimensi ketiga yang default-nya adalah 1 karena kita tidak membutuhkannya di sini.
Dalam dunia komputasi GPU, mengenkode perintah untuk menjalankan fungsi kernel pada sekumpulan data disebut pengiriman.
Ukuran grid workgroup untuk shader komputasi adalah (8, 8)
dalam kode
WGSL. Karena itu, "x" dan "y" yang masing-masing merupakan jumlah baris matriks pertama dan jumlah kolom matriks kedua akan dibagi 8. Dengan demikian, kini kita dapat mengirim panggilan komputasi dengan
passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)
. Jumlah
grid 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 mana 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 penerusan komputasi, panggil passEncoder.end()
. Kemudian, buat
buffer GPU untuk digunakan sebagai tujuan menyalin buffer matriks hasil dengan
copyBufferToBuffer
. Terakhir, selesaikan encoding perintah dengan
copyEncoder.finish()
dan kirimkan ke antrean perangkat GPU dengan memanggil
device.queue.submit()
menggunakan 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 dapat memperoleh 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.
Satu trik terakhir
Salah satu cara membuat kode Anda lebih mudah dibaca adalah dengan menggunakan metode
getBindGroupLayout
yang praktis dari pipeline komputasi untuk menyimpulkan tata letak
bind group dari modul shader. Trik ini menghilangkan kebutuhan dari membuat
tata letak bind group kustom dan menentukan tata letak pipeline di pipeline
komputasi 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 perbandingan antara menjalankan perkalian matriks pada 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 kekuatan penuh GPU tampak seperti pilihan yang jelas jika ukuran matriks lebih besar dari 256 x 256.
Artikel ini hanyalah awal dari perjalanan saya untuk menjelajahi WebGPU. Nantikan artikel lainnya yang akan segera menampilkan pembahasan lebih mendalam tentang Komputasi GPU dan cara kerja rendering (kanvas, tekstur, sampler) di WebGPU.