Bu yayın, örnekler üzerinden deneysel WebGPU API'sini ele almaktadır ve GPU'yu kullanarak verilere paralel hesaplamalar yapmaya başlamanıza yardımcı olmaktadır.
Arka plan
Bildiğiniz gibi, Grafik İşlem Birimi (GPU), başlangıçta grafik işleme konusunda uzmanlaşmış, bilgisayar içinde bulunan bir elektronik alt sistemdir. Ancak geçtiğimiz 10 yılda, geliştiricilerin yalnızca 3D grafikler oluşturmakla kalmayıp GPU'nun benzersiz mimarisinden de faydalanmakla kalmayıp birçok algoritma türünü uygulamasını da sağlayan daha esnek bir mimariye dönüştü. Bu özelliklere GPU Compute denir. Genel amaçlı bilimsel işlemler için ek işlemci olarak GPU kullanılmasına, genel amaçlı GPU (GPU) programlaması denir.
Konvolüsyonlu nöral ağlar ve diğer modeller GPU'lar üzerinde daha verimli bir şekilde çalışmak için bu mimariden yararlanabildiğinden GPU Compute, yakın zamanda yaşanan makine öğrenimi artışına önemli ölçüde katkı sağladı. Mevcut Web Platformu'nda GPU Compute özellikleri bulunmadığı için W3C'nin "GPU for the Web" Topluluk Grubu, en güncel cihazlarda bulunan modern GPU API'lerini açığa çıkaracak bir API tasarlıyor. Bu API'ye WebGPU adı verilir.
WebGPU, WebGL gibi alt düzey bir API'dir. Göreceğiniz gibi, bu oldukça güçlü ve oldukça ayrıntılı. Ama bu sorun değil. Aradığımız şey performans.
Bu makalede WebGPU'nun GPU İşlem kısmına odaklanacağım ve dürüst olmak üzere, kendi başınıza oynamaya başlayabilmeniz için yalnızca yüzeyleri kazıyacağım. Gelecek makalelerde WebGPU oluşturmayı (ör. tuval, doku vb.) ayrıntılı bir şekilde ele alacağım.
GPU'ya erişme
WebGPU'da GPU'ya erişmek kolaydır. navigator.gpu.requestAdapter()
çağrıldığında, bir GPU adaptörüyle eşzamansız olarak çözülecek bir JavaScript sözü döndürülür. Bu adaptörü grafik kartı olarak düşünebilirsiniz. Entegre (CPU ile aynı çip üzerinde) veya ayrı (genellikle daha yüksek performanslı ancak daha fazla güç kullanan bir PCIe kartı) olabilir.
GPU bağdaştırıcısına sahip olduğunuzda, GPU hesaplamaları yapmak üzere kullanacağınız bir GPU cihazıyla sonuçlanacak sözü almak için adapter.requestDevice()
çağrısı yapın.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
Her iki işlev de istediğiniz adaptör türü (güç tercihi) ve cihaz (uzantılar, sınırlar) konusunda spesifik olmanıza imkan veren seçenekler sunar. Basitlik sağlaması için bu makaledeki varsayılan seçenekleri kullanacağız.
Arabelleğe alma
GPU için belleğe veri yazmak üzere JavaScript'in nasıl kullanılacağını görelim. Modern web tarayıcılarında kullanılan korumalı alana alma modeli nedeniyle bu işlem kolay değildir.
Aşağıdaki örnekte, GPU'dan erişilebilen belleği arabelleğe almak için dört baytın nasıl yazılacağı gösterilmektedir. Arabelleğin boyutunu ve kullanımını alan device.createBuffer()
yöntemini çağırır. Bu çağrı için GPUBufferUsage.MAP_WRITE
kullanım işareti gerekli olmasa da bu arabelleğe yazmak istediğimizi açıkça belirtmek isteriz. mappedAtCreation
doğru değerine ayarlandığında, oluşturma sırasında bir GPU arabellek nesnesi eşlenir. Daha sonra, ilişkili ham ikili veri arabelleği getMappedRange()
GPU arabelleği yöntemi çağrılarak alınabilir.
Daha önce ArrayBuffer
ile oynadıysanız bayt yazmaya aşina olursunuz. TypedArray
kullanıp değerleri buraya kopyalayın.
// 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]);
Bu noktada GPU arabelleği eşlenir, yani CPU'nun sahibi olur ve JavaScript'ten okuma/yazma yöntemiyle erişilebilir. GPU'nun erişebilmesi için eşlemenin kaldırılması gerekir. Bu, gpuBuffer.unmap()
yöntemini çağırmak kadar basittir.
Eşlenmiş/eşlenmemiş kavramı, GPU ve CPU'nun belleğe aynı anda eriştiği yarış koşullarını önlemek için gereklidir.
Arabellek belleğini okuma
Şimdi bir GPU arabelleğinin başka bir GPU arabelleğine nasıl kopyalanacağını ve geri okunacağını görelim.
İlk GPU arabelleğine yazıyor ve ikinci bir GPU arabelleğine kopyalamak istediğimiz için yeni bir kullanım işareti GPUBufferUsage.COPY_SRC
gereklidir. İkinci GPU arabelleği bu kez device.createBuffer()
ile eşlenmemiş bir durumda oluşturulur. Kullanım işareti, ilk GPU arabelleğinin hedefi olarak kullanılacağı ve GPU kopyalama komutları yürütüldükten sonra JavaScript'te okunacağı için GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ
şeklindedir.
// 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 bağımsız bir ortak işlemci olduğundan tüm GPU komutları eşzamansız olarak yürütülür. Bu nedenle, gerektiğinde derlenip toplu olarak gönderilen GPU komutlarının bir listesi bulunur. WebGPU'da device.createCommandEncoder()
tarafından döndürülen GPU komutu kodlayıcı, bir noktada GPU'ya gönderilecek "arabelleğe alınmış" komut grubunu oluşturan JavaScript nesnesidir. Öte yandan, GPUBuffer
üzerindeki yöntemler "arabelleğe alınmamış"tır. Yani çağrıldıkları sırada atomik olarak yürütülürler.
GPU komut kodlayıcınız olduğunda, bu komutu daha sonra çalıştırmak üzere komut kuyruğuna eklemek için aşağıda gösterildiği gibi copyEncoder.copyBufferToBuffer()
çağrısı yapın.
Son olarak copyEncoder.finish()
komutunu çağırarak kodlama komutlarını tamamlayın ve bunları GPU cihazı komut sırasına gönderin. Sıra, GPU komutlarıyla bağımsız değişken olarak device.queue.submit()
aracılığıyla yapılan gönderimlerden sorumludur.
Bu işlem, dizide depolanan tüm komutları atomik olarak sırayla yürütür.
// 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]);
Bu noktada GPU sıra komutları gönderilmiş olsa da yürütülmesi gerekmez.
İkinci GPU arabelleğini okumak için GPUMapMode.READ
ile gpuReadBuffer.mapAsync()
çağrısı yapın. GPU arabelleği eşlendiğinde çözümlenecek bir söz döndürür. Ardından, sıraya alınan tüm GPU komutları yürütüldükten sonra, ilk GPU arabelleği ile aynı değerleri içeren gpuReadBuffer.getMappedRange()
ile eşlenen aralığı alın.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Arabellek belleği işlemleriyle ilgili olarak hatırlamanız gereken noktalar şunlardır:
- Cihaz sırası gönderiminde kullanılacak GPU arabelleklerinin eşlemesi kaldırılmalıdır.
- Eşlenen GPU arabellekleri, JavaScript'te okunabilir ve yazılabilir.
mappedAtCreation
değeri true olarak ayarlanmışmapAsync()
vecreateBuffer()
çağrıldığında GPU arabellekleri eşlenir.
Gölgelendirici programlama
GPU'da çalışan, yalnızca hesaplama gerçekleştiren (ve üçgen çizmeyen) programlara işlem gölgelendiriciler adı verilir. Bu uygulamalar, verileri analiz etmek için birlikte çalışan yüzlerce GPU çekirdeği (CPU çekirdeğinden daha küçük) ile paralel olarak yürütülür. Giriş ve çıkışları WebGPU'da arabellektir.
WebGPU'da hesaplama gölgelendiricilerinin kullanımını göstermek için aşağıda gösterilen ve makine öğreniminde sık kullanılan bir algoritma olan matris çarpımıyla çalışacağız.
Kısacası, şunları yapacağız:
- Üç GPU tamponu oluşturun (matrislerin çarpılması için iki, sonuç matrisi için bir tane)
- Compute gölgelendirici için giriş ve çıkışı açıklama
- İşlem gölgelendirici kodunu derleme
- Compute ardışık düzen oluşturma
- Kodlanmış komutları GPU'ya toplu olarak gönderme
- Sonuç matrisi GPU arabelleğini okuma
GPU Arabellekleri oluşturma
Basitlik açısından, matrisler kayan nokta sayıları listesi olarak gösterilir. İlk öğe satır sayısı, ikinci öğe sütun sayısıdır ve gerisi de matrisin gerçek sayılarıdır.
Üç GPU tamponu, verileri işlem gölgelendiricide depolayıp almamız gerektiğinden depolama arabellekleridir. Bu, GPU arabellek kullanım işaretlerinin neden tümü için GPUBufferUsage.STORAGE
içerdiğini açıklar. Sonuç matrisi kullanım işareti, tüm GPU sırası komutlarının tamamı yürütüldükten sonra okuma için başka bir arabelleğe kopyalanacağından bu işarette GPUBufferUsage.COPY_SRC
de bulunur.
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
});
Grup düzenini ve bağlama grubunu bağlayın
Bağlama grubu düzeni ve bağlama grubu kavramları WebGPU'ya özgüdür. Bağlama grubu düzeni, gölgelendiricinin beklediği giriş/çıkış arayüzünü tanımlarken, bağlama grubu, gölgelendiricinin gerçek giriş/çıkış verilerini temsil eder.
Aşağıdaki örnekte bağlama grubu düzeni, 0
ve 1
numaralı giriş bağlamalarında iki salt okunur depolama arabelleği ve işlem gölgelendirici için 2
konumunda bir depolama arabelleği bekler.
Diğer yandan, bu bağlama grubu düzeni için tanımlanan bağlama grubu, GPU arabelleklerini girişlerle ilişkilendirir: gpuBufferFirstMatrix
; 0
bağlama, gpuBufferSecondMatrix
, 1
ve resultMatrixBuffer
bağlama 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
}
}
]
});
Gölgelendirici kodunu hesaplayın
Matrisleri çoğaltmaya yönelik işlem gölgelendirici kodu, SPIR-V'ye trivilite çevrilebilir olan WGSL, yani WebGPU Gölgelendirici Dili ile yazılır. Ayrıntıya girmeden var<storage>
ile tanımlanan üç depolama arabelleğini aşağıda bulabilirsiniz. Program, firstMatrix
ve secondMatrix
öğelerini, çıktı olarak da resultMatrix
kullanır.
Her depolama arabelleğinin, yukarıda tanımlanan bağlama grubu düzenlerinde ve bağlama gruplarında tanımlanan aynı dizine karşılık gelen bir binding
süslemesinin kullanıldığını unutmayın.
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;
}
`
});
Ardışık düzen kurulumu
İşlem ardışık düzeni, gerçekleştireceğimiz işlem işlemini tanımlayan nesnedir. device.createComputePipeline()
yöntemini çağırarak oluşturun.
Bunun için iki bağımsız değişken gerekir: daha önce oluşturduğumuz bağlama grubu düzeni ve işlem gölgelendiricimizin giriş noktasını tanımlayan işlem aşaması (main
WGSL işlevi) ve device.createShaderModule()
ile oluşturulan asıl işlem gölgelendirici modülü.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Komut gönderme
Üç GPU arabelleğimizle bir bağlama grubu ve bağlama grubu düzenine sahip bir işlem ardışık düzeniyle örnek oluşturduktan sonra sıra bunları kullanmaya geldi.
commandEncoder.beginComputePass()
ile programlanabilir bir Compute Pass kodlayıcısı başlatalım. Bunu, matris çarpım işlemini gerçekleştirecek GPU komutlarını kodlamak için kullanacağız. Ardışık düzenini passEncoder.setPipeline(computePipeline)
ile, bağlama grubunu da passEncoder.setBindGroup(0, bindGroup)
ile 0 dizininde ayarlayın. Dizin 0, WGSL kodundaki group(0)
süslemesine karşılık gelir.
Şimdi de bu bilgi işlem gölgelendiricisinin GPU'da nasıl çalışacağından bahsedelim. Amacımız, bu programı sonuç matrisindeki her bir hücre için, adım adım, paralel olarak yürütmektir. Örneğin, 16x32 boyutundaki bir sonuç matrisi için yürütme komutunu kodlamak için @workgroup_size(8, 8)
öğesinde passEncoder.dispatchWorkgroups(2, 4)
veya passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)
adını veririz.
Birinci bağımsız değişken "x" birinci boyut, ikinci "y" ikinci boyut ve son "z" de üçüncü boyuttur. Bu boyut, burada gerekli olmadığı için varsayılan olarak 1'dir.
GPU hesaplama dünyasında, bir veri kümesi üzerinde çekirdek işlevini yürütmek için bir komut kodlamaya "gönderme" adı verilir.
İşlem gölgelendiricimiz için çalışma grubu ızgarasının boyutu, WGSL kodumuzda (8, 8)
'tır. Bu nedenle, sırasıyla birinci matrisin satır sayısı ve ikinci matrisin sütun sayısı olan "x" ve "y" 8'e bölünür. Bu sayede artık passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)
ile bir işlem çağrısı dağıtabiliyoruz. Çalıştırılacak çalışma grubu ızgaralarının sayısı dispatchWorkgroups()
bağımsız değişkenleridir.
Yukarıdaki çizimde görüldüğü gibi, her gölgelendirici, hangi sonuç matrisi hücresinin hesaplanacağını bilmek için kullanılacak benzersiz bir builtin(global_invocation_id)
nesnesine erişebilir.
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();
Compute Pass kodlayıcıyı sonlandırmak için passEncoder.end()
numaralı telefonu arayın. Ardından, sonuç matrisi arabelleğini copyBufferToBuffer
ile kopyalamak için hedef olarak kullanılacak bir GPU arabelleği oluşturun. Son olarak copyEncoder.finish()
ile kodlama komutlarını tamamlayın ve GPU komutlarıyla device.queue.submit()
yöntemini çağırarak bunları GPU cihaz sırasına gönderin.
// 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]);
Sonuç matrisini oku
Sonuç matrisini okumak, GPUMapMode.READ
ile gpuReadBuffer.mapAsync()
yöntemini çağırmak ve geri dönen sözün çözümlenmesini beklemek kadar kolaydır. Bu, GPU arabelleğinin eşlendiğini gösterir. Bu noktada, eşlenen aralığı gpuReadBuffer.getMappedRange()
ile almak mümkündür.
Kodumuzda, Geliştirici Araçları JavaScript konsoluna kaydedilen sonuç "2, 2, 50, 60, 114, 140" şeklindedir.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
Tebrikler! Ba. Örnekle oynayabilirsiniz.
Son bir ipucu
Kodunuzun okunmasını daha kolay hale getirmenin yollarından biri, bağlama grubu düzenini gölgelendirici modülünden çıkarmak için işlem ardışık düzeninin kullanışlı getBindGroupLayout
yöntemini kullanmaktır. Bu yöntem, aşağıda gösterildiği gibi özel bir bağlama grubu düzeni oluşturma ve işlem ardışık düzeninizde ardışık düzen düzeni belirtme ihtiyacını ortadan kaldırır.
Önceki örnek için getBindGroupLayout
görseli kullanılabilir.
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: [
Performans bulguları
Peki, matris çarpım işlemini GPU'da çalıştırmakla CPU üzerinde çalıştırma arasındaki fark nedir? Öğrenmek için biraz önce CPU için açıklanan programı yazdım. Ayrıca, aşağıdaki grafikte görebileceğiniz gibi, matrislerin boyutu 256x256'dan büyük olduğunda GPU'nun tam gücünü kullanmak bariz bir seçenek gibi görünmektedir.
Bu makale WebGPU'yu keşfetme yolculuğumun yalnızca başlangıcıydı. Yakında GPU Compute ve WebGPU'da oluşturma (kanvas, doku, sampler) ile ilgili daha ayrıntılı bilgiler içeren daha fazla makale alacaksınız.