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

Bu gönderi, deneysel WebGPU API'yi örneklerle keşfeder ve GPU'yu kullanarak veri paralel hesaplamalar yapmaya başlamanıza yardımcı olur.

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 geliştiricilerin yalnızca 3D grafikler oluşturmakla kalmayıp GPU'nun benzersiz mimarisinden yararlanmasına olanak tanıyan birçok algoritma türünü uygulamasına olanak tanıyan daha esnek bir mimariye doğru ilerledi. 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ü bir grafik kartı olarak düşünebilirsiniz. Entegre (CPU ile aynı çipte) veya ayrı (genellikle daha yüksek performanslı 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. Kolaylık sağlaması açısından, bu makaledeki 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. Tamponun boyutunu ve kullanımını alan device.createBuffer() işlevini ç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'nun belleğe aynı anda eriştiği yarış koşullarını önlemek için eşlenmiş/eşlenmemiş kavramı gerekir.

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ı çalıştırıldıktan sonra JavaScript'te okunacağı için kullanım işareti GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ'tür.

// 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 yardımcı 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, dizide depolanan tüm komutları 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ırası komutları gönderilmiştir ancak mutlaka uygulanmamıştır. İkinci GPU arabelleğini okumak için GPUMapMode.READ ile gpuReadBuffer.mapAsync() işlevini çağırı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 kuyruğu gönderiminde kullanılabilmesi için GPU arabelleklerinin eşlemesinin kaldırılması gerekir.
• Haritalandığında GPU arabellekleri JavaScript'te okunabilir ve yazılabilir.
• mappedAtCreation doğru değerine ayarlanmış 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 işlemek için birlikte çalışan yüzlerce GPU çekirdeği (CPU çekirdeklerinden daha küçüktür) tarafından paralel olarak yürütülür. Giriş ve çıkışları WebGPU'da arabelleklerdir.

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. Hesaplama gölge kodunu derleme
4. Compute 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 tamponu ve bilgi işlem gölgelendirici için 2 adresinde bir depolama tamponu 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
      }
    }
  ]
});

Gölgelendirici kodunu hesapla

Matrisleri çarpma işlemini yapan hesaplama gölge kodu, WGSL (WebGPU Gölge Kodu Dili) ile yazılmıştır. Bu kod, SPIR-V'ye kolayca çevrilebilir. Ayrıntılara girmeden, var<storage> ile tanımlanan üç depolama arabelleğinin aşağıda gösterildiğini görebilirsiniz. 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

Hesaplama ardışık düzeni, gerçekleştireceğimiz hesaplama işlemini açıklayan nesnedir. device.createComputePipeline() numaralı telefonu arayarak 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ğimizle bir bağlama grubu ve bağlama grubu düzeniyle bir hesaplama ardışık düzeni oluşturduktan sonra bunları kullanma zamanı gelmiştir.

commandEncoder.beginComputePass() ile programlanabilir işlem kartı kodlayıcı başlatalım. Bu matrisi, matris çarpımını gerçekleştirecek GPU komutlarını kodlamak için kullanırız. passEncoder.setPipeline(computePipeline) ile ardışık düzenini ve passEncoder.setBindGroup(0, bindGroup) ile 0 dizinindeki bağlama grubunu 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, artık passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) ile bir Compute ç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();

İşlem kartı kodlayıcıyı sonlandırmak için passEncoder.end() komutunu çağırın. Ardından, sonuç matris arabelleğinin copyBufferToBuffer ile kopyalanması 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 okuma

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 ipucu

Kodunuzun okunmasını kolaylaştırmanın bir yolu, bağlantı grubu düzenini gölgelendirici modülünden tahmin etmek için hesaplama ardışık düzeninin kullanışlı getBindGroupLayout yöntemini kullanmaktır. 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? Öğrenmek için az önce CPU için anlattığım programı yazdım. Aşağıdaki grafikte görebileceğiniz gibi, matrislerin boyutu 256x256'dan büyük olduğunda GPU'nun tam gücünü kullanmak gayet doğaldır.

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.