Tuvalden WebGPU'ya: IR merdiveni

Studio'da bir iş akışı çizersiniz — düğümler, bağlantılar, parametreler. Çalışan şey ise çizim değildir: Makineniz için üretilmiş ve tam orada, tarayıcı sekmesinin içinde yürütülen GPU kodudur. İkisinin arasında bir ara temsiller merdiveni durur ve güvenceler bu merdivenden inerken kurulur — uygunluk hiçbir şey çalışmadan netleşir, determinizm varsayılmak yerine beyan edilir, her çekirdek güvenini bir referans karşısında kazanır. Bu merdivenin sahibi olan motor, bir önceki makaledeki mikroçekirdektir; burada onu iş başında izliyoruz.

Merdivenden aşağı, basamak basamak#

İndirgeme, adlandırılmış ve incelenebilir aşamalarla ilerler. Workflow dokümanı — çizdiğiniz graf — birimi blok olan tipli bir programa derlenir (BlockIR): Yapı burada açık hale gelir — hangi bloklar, hangi sırayla, hangi kontrol akışı altında çalışır. Sayısal iş, otomatik türevin gerçekleştiği basamak olan tensör düzeyinde bir işlem grafına iner (WebIR), oradan arka uca özgü çekirdeklere — WebGPU için WGSL, yanında CPU referansı (KernelIR) — ve en sonunda somut dispatch'lerden, tamponlardan ve bir bellek planından oluşan bir yürütme planına (ExecPlan) iner. Bir çalıştırma cihazlara dağıldığında, bir basamak daha — WorkPlan — onu doğrulanabilir parçalara böler.

Her basamak, açıp okuyabileceğiniz birinci sınıf bir yapıttır. Resim ile sonuç arasında kara kutu yoktur; makinenin birazdan ne yapacağını merak ettiğinizde, sorunuza cevap veren basamağı incelersiniz.

graph (mid-ladder, simplified)
  %x    = input "dataset"             shape [N, D]  dtype f32
  %w    = param "weights"             shape [D, K]  dtype f32
  %mean = reduce_mean %x axis=0       shape [D]
  %xc   = sub %x, broadcast %mean     shape [N, D]
  %y    = matmul %xc, %w              shape [N, K]

plan (computed before any dispatch)
  determinism_tier : prod             declared, never defaulted
  tolerance_class  : default          resolves to explicit thresholds per tier
  peak_memory      : parameters + activations + gradients
                     + optimizer_state + working
  verdict          : green            fits the declared budgets

Bellek çalıştırmadan önce netleşir#

Plan, uygunluk ilk dispatch'ten önce karara bağlansın diye vardır. Bellek, tahsis sınıflarına göre açık bir döküm olarak planlanır — parametreler, geri yayılım için tutulan aktivasyonlar, gradyanlar, optimizatör durumu ve çekirdek geçicileri için çalışma belleği — böylece “gizemli GPU belleği”nin saklanacak yeri kalmaz. Karma hassasiyet, f32 ana ağırlıklarını da tabloya ekler. Büyük girdiler beyan edilmiş akış sınıfları olarak planlanır; parça boyutları çökerek değil bütçeye göre hesaplanır ve karo düzeni, WebGPU'nun 128 MB'lık bağlama sınırını otomatik ele alır. Bu beş sınıfı ve sığmadıklarında neyi değiştirmek gerektiğini tarayıcıda model eğitme sayfası adım adım anlatır.

Planlama üç karardan biriyle biter. Yeşil bütçe içinde çalışır. Sarı uyarılarla çalışır — örneğin bellek içi bir geçişten daha yavaş bir akışla. Kırmızı reddeder ve bu red eyleme dönüktür: ne başarısız oldu, neden ve neyi değiştirmeli — daha küçük bir batch, akışı açmak, farklı bir hassasiyet. Çalıştırma ortasında bellek tükenmesi sürprizi, şanssızlık değil mimari kusur sayılır.

Üç determinizm kademesi — asla sessizce varsayılmaz#

Her çalıştırma, ne kadar deterministik olması gerektiğini beyan eder. Tam üç kademe vardır:

• debug — aynı cihazda, aynı tohumla, çalıştırmalar arasında bit düzeyinde özdeş sonuçlar. GPU üzerinde bu, deterministik çekirdekleri zorunlu kılar: sabit indirgeme sırası, atomik float toplama yok — hız maliyetine kesinlik satın alır.
• prod — beyan edilmiş sayısal tolerans sınıfı içinde deterministik; gerçek işin günlük duruşu.
• fast — önce iş çıkarma hızı; nondeterminizme yalnızca güvenli olduğu beyan edilen yerlerde izin verilir.

Kademe işle birlikte yolculuk eder: spesifikasyonlarda beyan edilir, yürütmede seçilir, çalıştırmayla kaydedilir. Asla sessizce varsayılmaz — geçerli kademesi olmayan bir iş akışı, sessizce bir kademe varsaymak yerine adlandırılmış bir tanıyla derlenmeyi reddeder. Birleşen parçalar farklı kademeler beyan ettiğinde en katı olan kazanır — siz açıkça geçersiz kılmadıkça; geçersiz kılma, gerçekten daha katı bir kademeyi düşürdüğü her durumda bir tanıyla kendini duyurur.

Farklı GPU'larda “aynı sonuç” ne demektir#

Açıkça söyleyelim: Farklı GPU'larda hesaplanan sonuçlar bit düzeyinde özdeş değildir. Kayan nokta toplaması birleşme özelliği taşımaz ve farklı cihazlar işi farklı sıralar. Cihazlar arasında geçerli olan, beyan edilmiş tolerans sınıfı içinde uyuşmadır — exact, tight, default, loose sınıflarından biri; her biri kademe başına açık eşiklere çözümlenir. Yani yeniden üretilebilirlik dürüstçe iki kademelidir: oracle yolunda replayable (bit düzeyinde kesin) ve GPU yolunda convergent (toleransla sınırlı). Bit düzeyinde tekrar tam iki yerde vardır — aynı cihazdaki debug kademesi ve her yerde deterministik olan CPU referans yolu. İstediğinizi karşılayamayan bir plan — örneğin GPU'larda cihazlar arası bit kesinliği — bir tanıyla reddedilir; ekranda aynı görünen daha zayıf bir güvenceye asla düşürülmez.

Güven bir referans karşısında kazanılır#

Yerleşik her GPU işlemi, bir CPU referans implementasyonuna karşı — platformun doğruluk çapası olan oracle'a karşı — beyan edilmiş tolerans sınıfı içinde doğrulanır. Oracle bir çalışma zamanı yedeği değil, doğruluk çapasıdır: Yerleşik bir işlem için eksik GPU çekirdeği düzeltilecek bir hatadır, asla daha yavaş koda sessiz bir sapma değil. Merdivenin dibinde, sınanan yapıt sıradan WGSL'dir — okunabilir ve denetlenebilir:

@group(0) @binding(0) var<storage, read>       x    : array<f32>;
@group(0) @binding(1) var<storage, read>       mean : array<f32>;
@group(0) @binding(2) var<storage, read_write> out  : array<f32>;
@group(0) @binding(3) var<uniform>             dims : vec2<u32>;  // rows, cols

@compute @workgroup_size(64)
fn center_columns(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= dims.x * dims.y) { return; }
  let col = i % dims.y;
  // one element per invocation; conformance compares this kernel
  // to its CPU reference within the declared tolerance class
  out[i] = x[i] - mean[col];
}

Kullanıcı hesaplaması da aynı şekilde kazanır#

Kullanıcı yazımı hesaplama da aynı disipline tabidir, çünkü hiçbir zaman kendi kendini onaylamaz. Bir çekirdek yazarsanız bir referans da verirsiniz; uygunluk, beyan edilen imzadan ve toleranstan durumlar üretir ve çekirdeğin eşleştiğini kanıtlar. Özel bir gradyan, kendi ileri yönünüzün sonlu farklarına karşı gradcheck'ten geçer. Geçmek bir ConformanceReceipt üretir — profile bağlıdır; farklı bir GPU zarfına sahip bir cihaz, hiçbir GPU'nun karşılayamayacağı bir iddiayı miras almak yerine denetimi yerelde yeniden koşar. Receipt yürütmeyi değil güveni sınırlar: Kendi makinenizde uygunluğu kanıtlanmamış bir çekirdeği her zaman geliştirebilirsiniz. Bütün doğruluk rejimi budur ve dispatch anında motor, bir topluluk çekirdeğini yerleşik olandan ayırt etmez.

GPU yok mu? Yüksek sesli bir CPU yolu, asla boş bir tuval değil#

Her cihazda WebGPU yoktur — ayrıca WebGPU render'ı garanti bir taban değil, bir yetenek bayrağıdır; garanti bağlamlar Canvas2D ve WebGL2'dir. Scellis açılışta yetenekleri yoklar ve her sonuca tanımlı, görünür bir durum verir. GPU yok ama çalışan bir CPU motoru var: İş akışlarının tamamı, yüksek sesli ve kalıcı bir uyarı şeridinin arkasında CPU referans arka ucunda çalışır — CPU üzerinde, daha yavaş, ağır eğitim sınırlı. İkisi de yoksa: Studio yine görüntüleme ve yazarlık için açılır — kompoze edin, inceleyin, paylaşın — ve neden hiçbir şeyin yürütülmeyeceğini tam olarak söyler. Asla olmayan şey, boş bir tuval ya da sessiz bir cevapsızlıktır.

Bir çalıştırmanın bütün yolu budur: Bir çizim basamak basamak WGSL'e iner; bellek ve uygunluk ilk dispatch'ten önce netleşir; determinizm, cihazlar arası dürüst anlamlarıyla beyan edilmiş bir sözleşmedir; ve her çekirdek, inceleyebileceğiniz bir referans karşısında görünür bir güven kazanır. Bir çalıştırmayı güvenilir kılan aynı içerik adresli yapıtlar, onu URL ile yeniden üretilebilir de kılar.