Web'de GPU Compute'u kullanmaya başlayın

Bu gönderide, deneysel WebGPU API, örneklerle incelenecek ve GPU'yu kullanarak veriye paralel hesaplamalar yapmaya başlamanıza yardımcı olacaktır.

François Beaufort

Arka plan

Grafik İşlem Birimi'nin (GPU) bilgisayardaki bir elektronik alt sistem olduğunu ve başlangıçta grafik işleme konusunda uzmanlaştığını biliyorsunuzdur. Ancak son 10 yılda, GPU'nun benzersiz mimarisinden yararlanırken geliştiricilerin yalnızca 3D grafikleri oluşturmakla kalmayıp birçok algoritma türünü de uygulamasını sağlayan daha esnek bir mimariye doğru gelişti. Bu özelliklere GPU Compute denir ve GPU'yu genel amaçlı bilimsel hesaplama için yardımcı işlemci olarak kullanmak genel amaçlı GPU (GPGPU) programlama olarak adlandırılır.

GPU Compute, son zamanlardaki makine öğrenimi patlamasına önemli ölçüde katkıda bulundu. Çünkü konvolüsyonel sinir ağları ve diğer modeller, GPU'larda daha verimli çalışabilmek için mimariden yararlanabilir. Mevcut Web Platformu'nda GPU bilgi işlem özellikleri bulunmadığından W3C'nin "Web için GPU" Topluluk Grubu, mevcut cihazların çoğunda bulunan modern GPU API'lerini kullanıma sunmak için bir API tasarlıyor. Bu API'nin adı WebGPU'dur.

WebGPU, WebGL gibi düşük düzey bir API'dir. Göreceğiniz gibi, çok güçlü ve oldukça ayrıntılı bir yöntemdir. Ama sorun değil. Amacımız performanstır.

Bu makalede, WebGPU'nun GPU Compute kısmına odaklanacağım. Dürüst olmak gerekirse, kendi başınıza oynamaya başlayabilmeniz için konuyu sadece yüzeysel olarak ele alacağım. Gelecekteki makalelerde WebGPU oluşturma (tuval, doku vb.) hakkında daha ayrıntılı bilgi vereceğim.

GPU'ya erişme

WebGPU'da GPU'ya erişmek kolaydır. navigator.gpu.requestAdapter() çağrısı, bir GPU bağdaştırıcısı ile eşzamansız olarak çözülecek bir JavaScript promise döndürür. Bu adaptörü grafik kartı olarak düşünebilirsiniz. Entegre (CPU ile aynı çip üzerinde) veya ayrı (genellikle daha iyi performans gösteren ancak daha fazla güç kullanan bir PCIe kart) olabilir.

GPU adaptörünü aldıktan sonra, GPU hesaplaması yapmak için kullanacağınız bir GPU cihazıyla çözülecek bir promise almak üzere adapter.requestDevice() işlevini çağırı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) hakkında net olmanıza olanak tanıyan seçenekler sunar. Basitlik açısından bu makalede varsayılan seçenekleri kullanacağız.

Yazma arabellek belleği

GPU'nun belleğine veri yazmak için JavaScript'in nasıl kullanılacağını görelim. Modern web tarayıcılarında kullanılan korumalı alan modeli nedeniyle bu işlem basit değildir.

Aşağıdaki örnekte, GPU'dan erişilebilen belleğe dört baytın nasıl yazılacağını görebilirsiniz. Arabelleğin boyutunu ve kullanımını alan device.createBuffer() öğesini ç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 belirtelim. mappedAtCreation öğesinin doğru değerine ayarlanması nedeniyle oluşturma sırasında eşlenen bir GPU arabellek nesnesi oluşturulur. Ardından, GPU arabellek yöntemi getMappedRange() çağrılarak ilişkili ham ikili veri arabelleği alınabilir.

ArrayBuffer ile daha önce oynadıysanız bayt yazmak normaldir. TypedArray kullanın ve 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 arabellek, CPU'ya aittir ve JavaScript'ten okuma/yazma olarak erişilebilir. GPU'nun erişebilmesi için haritanın kaldırılması gerekir. Bu işlem, gpuBuffer.unmap() çağrısı yapmak kadar basittir.

GPU ve CPU erişim belleğinin aynı anda gösterildiği yarış koşullarını önlemek için "eşlenmiş/eşlenmemiş" kavramı gereklidir.

Arabellek belleğini okuma

Şimdi bir GPU arabelleğinin başka bir GPU arabelleğine nasıl kopyalanacağını ve nasıl geri okunacağını görelim.

İlk GPU arabelleğine yazdığımız ve bunu ikinci bir GPU arabelleğine kopyalamak istediğimiz için yeni bir kullanım işareti GPUBufferUsage.COPY_SRC gerekir. İkinci GPU arabelleği bu kez device.createBuffer() ile eşlenmemiş durumda oluşturulur. İlk GPU arabelleğinin hedefi olarak kullanılacağı ve GPU kopyalama komutları yürütüldüğünde JavaScript'te okunacağı için kullanım işareti 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 ek işlemci olduğundan, tüm GPU komutları eşzamansız olarak yürütülür. Bu nedenle, GPU komutlarının bir listesi oluşturulur ve gerektiğinde toplu olarak gönderilir. WebGPU'da, device.createCommandEncoder() tarafından döndürülen GPU komutu kodlayıcı, bir noktada GPU'ya gönderilecek bir "arabelleğe alınmış" komut grubu oluşturan JavaScript nesnesidir. Öte yandan GPUBuffer'teki yöntemler "önbelleğe alınmaz". Yani çağrıldıkları anda atomik olarak yürütülürler.

GPU komut kodlayıcıyı aldıktan sonra, bu komutu daha sonra yürütülmek üzere komut sırasına eklemek için aşağıda gösterildiği gibi copyEncoder.copyBufferToBuffer() işlevini çağırın. Son olarak, copyEncoder.finish() çağrısı yaparak kodlama komutlarını tamamlayın ve bu komutları GPU cihaz komut sırasına gönderin. GPU komutlarını bağımsız değişken olarak kullanarak device.queue.submit() aracılığıyla yapılan gönderimlerin işlenmesi sıranın sorumluluğundadır. Bu işlem, dizgede depolanan tüm komutları sırayla atomik olarak 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ırası komutları gönderilmiş olsa da yürütülmeyebilir. İkinci GPU arabelleğini okumak için GPUMapMode.READ ile gpuReadBuffer.mapAsync() çağrısı yapın. GPU arabelleği eşlendiğinde çözülecek bir promise döndürür. Ardından, tüm sıralı GPU komutları yürütüldükten sonra ilk GPU arabelleğiyle 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));

Bu örneği deneyebilirsiniz.

Kısacası, arabellek bellek işlemleriyle ilgili olarak hatırlamanız gerekenler şunlardır:

• Cihaz sırası gönderiminde kullanılmak üzere GPU arabelleklerinin eşlenmemiş olması gerekir.
• Haritalandığında GPU arabellekleri JavaScript'te okunabilir ve yazılabilir.
• mappedAtCreation doğru olarak ayarlanmışken mapAsync() ve createBuffer() çağrıldığında GPU arabellekleri eşlenir.

Gölgelendirici programlama

GPU'da çalışan ve yalnızca hesaplama yapan (üçgen çizmeyen) programlara hesaplama gölgelendiricileri denir. Bunlar, verileri analiz etmek için birlikte çalışan yüzlerce GPU çekirdeği (CPU çekirdeklerinden daha küçük) tarafından paralel olarak yürütülür. Giriş ve çıkışları WebGPU'da tamponlardır.

WebGPU'da hesaplama gölgelendiricilerinin kullanımını açıklamak için makine öğrenimindeki yaygın bir algoritma olan matris çarpımıyla oynayacağız.

Özetlemek gerekirse şunları yapacağız:

1. Üç GPU arabelleği oluşturun (çarpılacak matrisler için iki tane ve sonuç matrisi için bir tane)
2. Bilgi işleme gölgelendiricinin girişini ve çıkışını açıklama
3. Compute gölgelendirici kodunu derleyin
4. İşlem ardışık düzeni oluşturma
5. Kodlanmış komutları GPU'ya toplu olarak gönderme
6. Sonuç matrisi GPU arabelleğini okuma

GPU arabelleklerinin oluşturulması

Kolaylık sağlamak için matrisler, kayan nokta sayılarının listesi olarak gösterilir. İlk öğe satır sayısı, ikinci öğe sütun sayısı, geri kalanı ise matrisin gerçek sayılarıdır.

Bilgi işlem gölgelendiricisinde veri depolayıp almamız gerektiğinden, üç GPU tamponu depolama tamponudur. Bu, GPU arabellek kullanım işaretlerinin neden hepsi için GPUBufferUsage.STORAGE içerdiğini açıklar. Sonuç matrisi kullanım işaretinde de GPUBufferUsage.COPY_SRC vardır. Bunun nedeni, tüm GPU sıra komutları yürütüldükten sonra okunması için başka bir arabelleğe kopyalanmasıdır.

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 bağlama ve grup bağlama

Bağlama grubu düzeni ve bağlama grubu kavramları WebGPU'ya özgüdür. Bağlama grubu düzeni, bir gölgelendirici tarafından beklenen giriş/çıkış arayüzünü tanımlar. Bağlama grubu ise bir 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 hesaplama gölgelendirici için 2 konumunda bir depolama arabelleği olmasını bekler. Öte yandan, bu bağlama grubu düzeni için tanımlanan bağlama grubu, GPU arabelleklerini girişlerle ilişkilendirir: gpuBufferFirstMatrix, 0 bağlamasıyla, gpuBufferSecondMatrix, 1 bağlamasıyla ve resultMatrixBuffer, 2 bağlamasıyla ilişkilendirilir.

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

Compute gölge kodu

Matrisleri çarpmak için kullanılan hesaplama gölgelendirici kodu, WebGPU Gölgelendirici Dili olan WGSL ile yazılır. Bu dil, günlük olarak SPIR-V'ye çevrilebilir. Ayrıntılara girmeden, var<storage> ile tanımlanan üç depolama arabelleğinin altında olduğunu göreceksiniz. Program, giriş olarak firstMatrix ve secondMatrix'ı, çıkış olarak ise resultMatrix'yi kullanır.

Her depolama arabelleğinin, yukarıda açıklanan bağlama grubu düzenlerinde ve bağlama gruplarında tanımlanan aynı dizeye karşılık gelen bir binding süslemesi olduğunu 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

Compute ardışık düzeni, gerçekleştireceğimiz işlem işlemini tanımlayan nesnedir. device.createComputePipeline() numaralı telefonu arayarak hesap oluşturun. İki bağımsız değişken alır: Daha önce oluşturduğumuz bağlama grubu düzeni ve bilgi işlem gölgelendiricimizin giriş noktasını (main WGSL işlevi) ve device.createShaderModule() ile oluşturulan gerçek bilgi işlem gölgelendirici modülünü tanımlayan bir bilgi işlem aşaması.

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

Komut gönderme

Üç GPU arabelleğimiz ve bağlantı grubu düzenine sahip bir işlem ardışık düzeniyle bağlantı grubu örneklendirildikten sonra sıra bunları kullanmaya geldi.

commandEncoder.beginComputePass() ile programlanabilir bir hesaplama geçişi kodlayıcısı başlatalım. Bu matrisi, matris çarpımını gerçekleştirecek GPU komutlarını kodlamak için kullanırız. Ardışık düzeni passEncoder.setPipeline(computePipeline) ve bağlama grubunu passEncoder.setBindGroup(0, bindGroup) ile 0 dizininde ayarlayın. 0 dizini, WGSL kodundaki group(0) süslemesine karşılık gelir.

Şimdi bu hesaplama gölgelendiricinin GPU'da nasıl çalışacağından bahsedelim. Amacımız, bu programı sonuç matrisinin her hücresi için paralel olarak adım adım yürütmektir. Örneğin, 16x32 boyutunda bir sonuç matrisinde yürütme komutunu kodlamak için @workgroup_size(8, 8) üzerinde passEncoder.dispatchWorkgroups(2, 4) veya passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) çağrıları yaparız. İlk bağımsız değişken "x" ilk boyuttur, ikinci bağımsız değişken "y" ikinci boyuttur ve son bağımsız değişken "z", burada ihtiyaç duymadığımız için varsayılan olarak 1 olan üçüncü boyuttur. GPU bilişim dünyasında, bir veri kümesi üzerinde çekirdek işlevini yürütecek bir komutu kodlamaya gönderme denir.

WGSL kodumuzda, hesaplama gölgelendiricimiz için iş grubu ızgarasının boyutu (8, 8)'tür. Bu nedenle, sırasıyla ilk matrisin satır sayısı ve ikinci matrisin sütun sayısı olan "x" ve "y" 8'e bölünür. Böylece, passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) ile bir hesaplama çağrısı gönderebiliriz. Ç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ç matris 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();

Hesaplama geçiş kodlayıcısını sonlandırmak için passEncoder.end() değerini çağırı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()'ı çağırarak bunları GPU cihaz kuyruğuna 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, gpuReadBuffer.mapAsync() ile GPUMapMode.READ çağrısı yapmak ve döndürülen promise'ın çözülmesini beklemek kadar kolaydır. Bu, GPU arabelleğinin artık eşlendiğini gösterir. Bu noktada, gpuReadBuffer.getMappedRange() ile eşlenen aralığı almak mümkündür.

Kodumuzda, DevTools 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. Sana Özel içeriğinizi oynatabilirsiniz.

Son bir numara

Kodunuzu okumayı kolaylaştırmanın bir yolu, işlem ardışık düzeninin pratik getBindGroupLayout yöntemini kullanarak gölgelendirici modülünden bağlama grubu düzenini tahmin etmektir. Bu hile, aşağıda görebileceğiniz gibi özel bir bağlama grubu düzeni oluşturma ve hesaplama ardışık düzeninizde bir ardışık düzen belirtme ihtiyacını ortadan kaldırır.

Önceki örnek için getBindGroupLayout'ün bir resmi mevcuttur.

 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ını GPU'da çalıştırmak ile CPU'da çalıştırmak arasında nasıl bir fark var? Bunu öğrenmek için, daha önce açıklanan programı bir CPU için yazdım. Aşağıdaki grafikte de görebileceğiniz gibi, matrislerin boyutu 256x256'tan büyük olduğunda GPU'nun tam gücünü kullanmak mantıklı bir seçimdir.

Bu makale, WebGPU'yu keşfetme yolculuğumun başlangıcıydı. GPU Compute'e ve WebGPU'da oluşturma işleminin (tuval, doku, örnekleyici) işleyişine dair daha ayrıntılı bilgiler içeren daha fazla makaleyi yakında yayınlayacağız.