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

Bu yayında, deneysel WebGPU API, örneklerle anlatılmış ve Böylece GPU'yu kullanarak verilere paralel hesaplamalar yapmaya başlayabilirsiniz.

François Beaufort
François Beaufort
GitHub

Arka plan

Gördüğünüz gibi Grafik İşlem Birimi (GPU) elektronik işleme almak üzere uzmanlaşmış bir bilgisayardaki alt sistemdir gösterir. Ancak son 10 yılda daha esnek bir çalışmaya dönüşerek geliştiricilerin yalnızca 3D grafikler oluşturabilir, aynı zamanda uygulamanın benzersiz mimarisinden GPU'ya dokunun. Bu özellikler GPU Compute olarak adlandırılır ve genel amaçlı bilimsel bilişimin yardımcı işlemcisine genel amaçlı GPU (GPU) programlaması.

GPU Compute, son zamanlarda ortaya çıkan makine öğrenimindeki artışa önemli katkı sağladı. konvolüsyon nöral ağları ve diğer modeller daha verimli çalışmasını sağlayacak bir mimariye sahip. Mevcut Web Platformu ile W3C'nin "GPU for the Web"i olarak kullandığında, GPU Compute Topluluk Grubu yaygın olarak kullanılan mevcut cihazlarda kullanılabilir. Bu API'ye WebGPU adı verilir.

WebGPU, WebGL gibi alt seviye bir API'dir. Hem güçlü hem de çok ayrıntılı olduğundan görürsünüz. Ama sorun değil. Aradığımız şey performans.

Bu makalede, WebGPU'nun GPU bilgi işlem kısmına odaklanacağım ve sadece küçük resmi çiziyorum ki siz de cihazınızda oynamaya başlayabilirsiniz sahip olmalıdır. Daha ayrıntılı bir inceleme yapacağım ve WebGPU oluşturma (kanvas, doku, vb.) ekleyebilirsiniz.

GPU'ya erişme

WebGPU'da GPU'ya erişmek kolaydır. navigator.gpu.requestAdapter() aranıyor bir GPU ile eşzamansız olarak çözümlenecek bir JavaScript taahhüdü döndürür adaptörü kullanabilirsiniz. Bu adaptörü bir grafik kartı olarak düşünebilirsiniz. Hem Google (CPU ile aynı çip üzerinde) veya ayrı (genellikle daha büyük bir PCIe kart) performans gösterir, ancak daha fazla güç kullanır).

GPU adaptörünü edindikten sonra söz almak için adapter.requestDevice() adlı cihazı arayın yapmak için kullanacağınız bir GPU cihazıyla çözülecektir.

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

Her iki işlev de projenizin türü hakkında spesifik olmanızı sağlayan istediğiniz adaptörü (güç tercihi) ve cihazı (uzantılar, sınırlar) seçin. basit olması açısından, bu makaledeki varsayılan seçenekleri kullanacağız.

Arabellek belleğine yaz

Şimdi de JavaScript'in, GPU için belleğe veri yazmak amacıyla nasıl kullanılacağına bakalım. Bu modern web sürümünde kullanılan korumalı alan modeli nedeniyle işlem kolay değildir izin verir.

Aşağıdaki örnekte erişilebilir arabellek için nasıl dört bayt yazılacağı gösterilmektedir tüm verimi GPU'dan alıyor. Şu boyutu alır: device.createBuffer() ve kullanımını ele alacağız. GPUBufferUsage.MAP_WRITE kullanım bayrağı Bu özel görüşme için gerekli değil. Yazmak istediğimizi açıkça belirtelim: geri yükleyebilirsiniz. Sonuç olarak, oluşturma sırasında eşlenen bir GPU arabellek nesnesi bulunur: mappedAtCreation doğru değerine ayarlandı. Daha sonra ilişkilendirilen ham ikili veri arabelleği getMappedRange() GPU arabellek yöntemini çağırarak alınır.

Daha önce ArrayBuffer oynadıysanız bayt yazmak normaldir; bir TypedArray'ı tıklayın ve değerleri buna 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 ve yani CPU'ya ait olur. JavaScript'ten okuma/yazma yöntemiyle erişilebilir. Böylece GPU, işleminin kaldırılması gerekir. Bu işlem, gpuBuffer.unmap() yöntemini çağırmak kadar basit bir işlemdir.

GPU'nun eşlendiği yarış koşullarını önlemek için "Eşlenmiş/eşlenmemiş" kavramı gereklidir. ve CPU erişim belleği aynı anda kullanılabilir.

Arabellek belleğini oku

Şimdi bir GPU arabelleğini başka bir GPU arabelleğine nasıl kopyalayacağınızı ve sonra tekrar nasıl okuyacağınızı görelim.

İlk GPU arabelleğine yazdığımız ve bunu bir ikinci GPU arabelleği, yeni bir kullanım işareti GPUBufferUsage.COPY_SRC gerekli. İkinci GPU arabelleği, bu kez device.createBuffer() İlk GPU'nun hedefi olarak kullanılacağı için kullanım işareti GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ arabelleğe alma ve GPU kopyalama komutları yürütüldükten sonra JavaScript'te okuma.

// 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ı yürütülür eşzamansız olarak ayarlayabilirsiniz. Bu nedenle derlenen ve gönderilen GPU komutlarının bir listesi vardır. gruplar. WebGPU'da, device.createCommandEncoder(), "arabelleğe alınmış" komutlarını kullanabilirsiniz. Bu yöntemler Öte yandan GPUBuffer, "arabelleğe alınmamış" olduğundan atomik olarak uygulanırlar. çağrıldığında da gösterilir.

GPU komutu kodlayıcıyı edindikten sonra copyEncoder.copyBufferToBuffer() komutunu çağırın komut kuyruğuna eklemek için aşağıdaki talimatları uygulayın. Son olarak, copyEncoder.finish() çağrısı yaparak kodlama komutlarını tamamlayın ve gönderin. bunları GPU cihaz komut sırasına ekler. Sıra, projenin device.queue.submit() aracılığıyla GPU komutlarının bağımsız değişken olarak kullanıldığı gönderimler. 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 yürütülmeyebilir. İkinci GPU arabelleğini okumak için gpuReadBuffer.mapAsync() öğesini şu işlemle çağırın: GPUMapMode.READ. GPU arabelleği şu anda mevcut olduğunda çözümlenecek bir söz ayrıntılarına inceleyebilirsiniz. Ardından, gpuReadBuffer.getMappedRange() içeren, eşlenen aralığı ve sıraya alınan tüm GPU komutlarında ilk GPU arabelleğiyle aynı değerleri içerir yürütülmüş olmalıdır.

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

Bu örneği deneyebilirsiniz.

Kısaca, arabellek belleği işlemleriyle ilgili olarak hatırlamanız gerekenler şunlardır:

  • Cihaz sırası gönderiminde kullanılmak üzere GPU arabelleklerinin eşlenmemiş olması gerekir.
  • Eşlenen GPU arabellekleri, JavaScript'te okunabilir ve yazılabilir.
  • GPU arabellekleri, aşağıdaki durumlarda mapAsync() ve createBuffer() olduğunda eşlenir: True olarak ayarlanmış mappedAtCreation öğeleri çağrılır.

Gölgelendirici programlama

GPU'da çalışan, yalnızca hesaplamalar gerçekleştiren (ve çizim yapmayan) programlar üçgenler) hesaplama gölgelendirici olarak adlandırılır. Bu toplantılar yüzlerce GPU çekirdeğinden daha küçük olan (CPU çekirdeklerinden daha küçüktür) dışı verilerdir. Giriş ve çıkışları WebGPU'da arabelleklerdir.

WebGPU'da bilgi işlem gölgelendiricilerinin kullanımını göstermek için matrisle oynayacağız. makine öğreniminde yaygın olarak kullanılan bir algoritma olan çarpma işlemi.

Matris çarpım diyagramı
Matris çarpım diyagramı

Kısaca şunları yapacağız:

  1. Üç GPU tamponu oluşturun (matrislerin çarpılması için iki, sonuç matrisi)
  2. Compute gölgelendirici için giriş ve çıkışı açıklama
  3. Compute gölgelendirici kodunu derleyin
  4. Compute ardışık düzeni oluşturma
  5. Kodlanmış komutları GPU'ya toplu olarak gönderme
  6. Sonuç matrisi GPU arabelleğini oku

GPU Arabelleği oluşturma

Basitlik açısından, matrisler hareketli bir liste . İlk öğe satır sayısı, ikinci öğe ise sütun sayısını, geri kalanlar ise matrisin gerçek sayılarıdır.

Bir matrisin JavaScript'te basit gösterimi ve matematiksel gösterimdeki eşdeğeri
Bir matrisin JavaScript'te basit gösterimi ve matematiksel gösterimdeki eşdeğeri

Verileri Google’da depolamamız ve almamız gerektiğinden üç GPU arabelleği de depolama arabelleğidir. Compute gölgelendiriciyi devre dışı bırakır. Bu, GPU arabellek kullanım işaretlerinde Tümü için GPUBufferUsage.STORAGE. Sonuç matrisi kullanım işareti de GPUBufferUsage.COPY_SRC çünkü başka bir arabelleğe kopyalanacak. tüm GPU sırası komutlarının tümü yürütüldükten sonra okuma.

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 grubu bağlama

Bağlama grubu düzeni ve bağlama grubu kavramları WebGPU'ya özgüdür. Bağlama grup düzeni bir gölgelendirici tarafından beklenen giriş/çıkış arayüzünü tanımlarken bağlama grubu, gölgelendirici için gerçek giriş/çıkış verilerini temsil eder.

Aşağıdaki örnekte, bağlama grubu düzeni şu konumda iki salt okunur depolama arabelleğinin olmasını bekler: numaralı giriş bağlamaları 0, 1 ve Compute gölgelendirici için 2 konumunda bir depolama arabelleği bulunur. Diğer yandan, bu bağlama grubu düzeni için tanımlanan bağlama grubu, Girişlere GPU arabellekleri: gpuBufferFirstMatrix bağlantıyla 0, gpuBufferSecondMatrix 1 ve resultMatrixBuffer < 2 bağlantısı.

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 çarpmak için kullanılan hesaplama gölgelendirici kodu WGSL'de yazılır. SPIR-V'ye çevrilebilen WebGPU Shader Dili. Yok: ayrıntılı olarak ele aldığımızda, belirlenen üç depolama tamponu var<storage> ile. Program, firstMatrix ve secondMatrix özelliklerini giriş ve çıkış olarak resultMatrix kullanabilirsiniz.

Her depolama arabelleğinde şuna karşılık gelen bir binding süslemesinin kullanıldığını unutmayın: yukarıda belirtilen bağlama grubu düzenlerinde ve bağlama gruplarında tanımlanan aynı dizin.

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, işlem işlemini gerçekten açıklayan nesnedir tam olarak bunu yapar. device.createComputePipeline() numaralı telefonu arayarak oluşturun. İki bağımsız değişken gerekir: Daha önce oluşturduğumuz bağlama grubu düzeni ve bir compute işlem gölgelendiricimizin giriş noktasını tanımlayan aşama (main WGSL işlevi) ve device.createShaderModule() ile oluşturulan gerçek Compute gölgelendirici modülü.

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

Komut gönderme

Üç GPU arabelleğimiz ve bir ardışık düzeniyle çalışıyor. Şimdi sıra bunları kullanmaya geldi.

Şimdi şununla programlanabilir: Compute Pass kodlayıcı: commandEncoder.beginComputePass() Bunu, GPU komutlarını kodlamak için kullanacağız fonksiyonunu kullanmanız gerekir. Ardışık düzeni şununla değiştirin: passEncoder.setPipeline(computePipeline) ve 0 dizinindeki bağlama grubu passEncoder.setBindGroup(0, bindGroup). Dizin 0, WGSL kodunda group(0) süslemesi.

Şimdi, bu Compute gölgelendiricinin GPU'da nasıl çalışacağından bahsedelim. Bizim hedef, bu programı sonuç matrisindeki her bir hücre için paralel olarak yürütmektir. adım adım anlatacağım. Örneğin, 16'ya 32 boyutunda bir sonuç matrisi için komutun @workgroup_size(8, 8) üzerinde çalışacağız passEncoder.dispatchWorkgroups(2, 4) veya passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Birinci bağımsız değişken olan "x" birinci boyut, ikinci boyut "y"dir. ikinci boyuttur, ve sonuncusu "z" burada gerekli olmadığı için varsayılan olarak 1'e ayarlanan üçü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.

Her sonuç matrisi hücresi için paralel yürütme
Her sonuç matrisi hücresi için paralel yürütme

WGSL'mizde Compute gölgelendiricimiz için çalışma grubu ızgarasının boyutu (8, 8) girin. Bu nedenle, "x" ve "y" sırasıyla ilk matris ve ikinci matrisin sütun sayısı bölünür 8 tarihine kadar. Böylece, artık iletilerimizi tarayarak passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) İlgili içeriği oluşturmak için kullanılan ç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 bir gölgelendirici, benzersiz bir Hangi sonucun kullanılacağını öğrenmek için kullanılacak builtin(global_invocation_id) nesne matris hücresini hesaplamalısınız.

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, Hedef matrisi arabelleğini copyBufferToBuffer Son olarak, kodlama komutlarını copyEncoder.finish() ve bunları çağırarak GPU cihaz sırasına gönderin: GPU komutlarını içeren device.queue.submit().

// 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() öğesini çağırmak kadar kolay GPUMapMode.READ ve geri dönen sözün tamamlanmasını bekliyor. Bu, GPU arabelleği eşlenmiştir. Bu noktada, artık verileri analiz etmek için gpuReadBuffer.getMappedRange() ile aralığı.

Matris çarpım sonucu
Matris çarpım sonucu

Kodumuzda, Geliştirici Araçları JavaScript konsoluna kaydedilen sonuç şu şekildedir: "2, 2, 50, 60, 114, 140 inç.

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));

Tebrikler! Ba. Örnekle oynayabilirsiniz.

Son bir ipucu

Kodunuzu daha kolay okunur hale getirmenin bir yolu da Bağlama grubunu belirlemek için işlem ardışık düzeninin getBindGroupLayout yöntemi düzenini kontrol edin. Bu yöntem, aynı zamanda özel bağlantı grubu düzeni ve işleminizde ardışık düzen düzeni belirtme aşağıdaki gibi ardışık düzen oluşturun.

Önceki örnek için getBindGroupLayout görseli sunuluyor.

 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, GPU'da matris çarpım işlemi çalıştırmak ile GPU'da CPU mu? Öğrenmek için az önce CPU için anlattığım programı yazdım. Bu nedenle, aşağıdaki grafikte GPU'nun tam gücünü kullanmak gayet makul bir tercih gibi görünüyor. olduğunda matrislerin boyutu 256'ya 256'dan büyük olur.

GPU ve CPU karşılaştırması
GPU ve CPU karşılaştırması

Bu makale, WebGPU'yu keşfetme yolculuğumun sadece başlangıcıydı. Daha fazlası yakında GPU Compute hakkında daha ayrıntılı bilgi içeren ve oluşturma (tuval, doku, örnekleyici) WebGPU'da çalışır.