#llm.istanbul — WGSL Kernel Etüdleri

WebGPU üzerinde sıfırdan LLM eğitimi: WGSL ile yazılmış her kernel'ın detaylı analizi.

Bir forward pass, bir backward pass, bir optimizer step — tüm training loop, satır satır.

#LLM nedir?

En sade hâliyle bir LLM (Large Language Model), bir sonraki kelimeyi tahmin eden bir makine. Telefon klavyendeki otomatik tamamlamayı düşün — ama steroid almış hâli. Ona "Türkiye'nin başkenti" yazarsın, o "Ankara" demeyi öğrenmiştir; çünkü milyonlarca cümle görüp "şu kelimelerden sonra büyük ihtimalle şu gelir" istatistiğini içine sindirmiştir. Sohbet, çeviri, kod yazma — hepsi aslında bu tek yeteneğin üstüne kurulu: sırada ne var?

Peki bunu nasıl yapıyor? Metni önce küçük parçalara (token) bölüyoruz, sonra her token'ı 01 Embedding'de gördüğümüz gibi bir sayı vektörüne çeviriyoruz. Bu vektörler bir yığın transformer katmanından geçiyor: her katmanda kelimeler birbirine bakıp bağlam topluyor (05 Attention), sonra her kelime tek başına biraz "düşünüyor" (FFN — 03 Linear, 06 Activation). En sonunda model her pozisyon için "sıradaki token şu olabilir" diye bir olasılık dağılımı üretiyor.

Öğrenme kısmı şöyle: modelin tahminiyle gerçek bir sonraki kelimeyi kıyaslıyoruz, ne kadar yanıldığını tek bir sayıyla ölçüyoruz (loss — 07 Cross-Entropy). Sonra bu hatayı tüm ağın içinden geriye doğru yayıyoruz (backward pass) ve her ağırlık için "seni hangi yöne kıpırdatsam hata azalır?" bilgisini çıkarıyoruz. Optimizer (12 AdamW) da her ağırlığı o yönde kıl payı oynatıyor. Bunu milyonlarca kez tekrarla — model yavaş yavaş dili "kapıyor".

İşin sırrı şu: bütün bu görkemli iş, aslında devasa bir matris çarpımı yığını artı birkaç yardımcı işlemden ibaret. Sihir yok; sadece çok sayıda sayının GPU'da paralel çarpılıp toplanması. Bu etütlerde tam olarak o işlemleri — embedding lookup, normalizasyon, matris çarpımı, attention, aktivasyon, loss, gradyanlar, optimizer — birer birer, kernel kernel açıyoruz. Yani "LLM nasıl çalışır?" sorusunun cevabını teoride değil, çalışan kodun her satırında göreceksin.

Aşağıdaki sıralama da modelin doğal akışını izliyor: bir token'ı tahmin etme (forward), ne kadar yanıldığını ölçme, ve düzeltme (backward + optimizer).

Bu seri, docs/learning/ altındaki MetalFoundry etüdlerinin WebGPU/WGSL eşdeğeridir. Algoritmalar aynı; dil ve runtime farklı:

#Nasıl Okunmalı?

Her etüd aynı yapıyı takip eder:

1. Ne Yapar? — Kernel'ın amacı, tek cümleyle
2. Matematiksel Tanım — Formüller
3. Bind Group ABI — Hangi binding slot'ta ne var (WGSL'de @group(0) @binding(N))
4. Dispatch Şekli — Workgroup boyutu ve grid
5. Satır Satır Açıklama — Kod parça parça
6. WGSL-Spesifik Notlar — Metal'den farkı, dil özellikleri
7. Code Review — Risk Analizi — Bulgu tablosu (varsa)
8. Hızlı Kontrol Listesi — Test coverage / sanity

Sıralama, modelin bir token'ı "tahmin etme" ve "öğrenme" akışını takip eder.

#0. Altyapı — Preamble, Yardımcılar, Sabitler

Her şeyden önce: tüm kernel'ların paylaştığı sözleşme ve yardımcı kernel'lar.

#I. Forward Pass — Token'dan Tahmine

Giriş: bir dizi token ID. Çıkış: "bir sonraki token ne olabilir?" sorusunun cevabı.

#II. Backward Pass — Tahmin Hatasından Düzeltmeye

Giriş: loss'un gradient'ı. Çıkış: her weight'ın ne kadar değişmesi gerektiği.

#III. Optimizer — Ağırlık Güncellemesi

#IV. BPE Tokenizer — Kelime Parçalama ve Sıkıştırma

Dil modeline giren ham metin verisini token ID'lerine dönüştüren ve çıkarım sırasında tokenları çözüp birleştiren bağımsız GPU tokenizer pipeline'ı.

#Pipeline Akış Diyagramı

                               ┌──── allocate (bir kere) ────┐
                               ▼                              │
tokens ──► 01 Embed ──► 02 Norm ──► 03 Linear ──► 04 RoPE ──┤
                                                              │
                                                   05 Attention
                                                              │
                                                   06 Activation (FFN)
                                                              │
                               ┌────── 02 Norm ◄── residual ◄┘
                               ▼
                      03 Linear (lm_head) ──► 07 Cross-Entropy
                                                   │  dLogits (fp32)
            ╔═══════════════════════════════════ │ ════════════════╗
            ║                BACKWARD PASS                          ║
            ║  11 Linear ← 09 FFN ← 10 Attn (split) ← 02 Norm bwd ║
            ║                                       ← 01 Embed bwd ║
            ╚═══════════════════════════════════════════════════════╝
                                  │
                       reduce_norm_sq → finalize_grad_stats
                                  │
                           12 AdamW (fused weight mirroring) ──► LOOP ↑

#Toplam Sayılar

#WGSL'nin Metal'den Birkaç Farkı

WGSL alıştırılması gereken şeyler:

#1. enable direktifleri preamble'da

WGSL "enable" direktifleri dosya başında olmak zorunda. Engine 00_shared.wgsl'in en başına enable f16;\nenable subgroups; injekte eder, sonra preamble'ın geri kalanını ve kernel kodunu append eder.

#2. Workgroup boyutu kernel attribute'unda

@compute @workgroup_size(256, 1, 1)
fn my_kernel(...) { ... }

Metal'de [[max_total_threads_per_threadgroup(N)]] sadece optimization hint'idir; WGSL'de kernel imzasında zorunludur.

#3. Bind group + binding (slot yerine)

@group(0) @binding(0) var<storage, read> X: array<f32>;
@group(0) @binding(1) var<storage, read_write> Y: array<f32>;

Metal'de [[buffer(N)]]. WGSL'de iki seviye: @group (descriptor set benzeri) + @binding (slot). Bizim kodda sadece @group(0) kullanıyoruz.

#4. Atomik tip wrap

var<storage, read_write> counts: array<atomic<u32>>;
atomicAdd(&counts[i], 1u);

Atomic değer tip seviyesinde etiketlenmek zorunda. Metal'de atomic_* operation'ları doğrudan int üzerinde çalışıyor.

#5. f16 kullanımı sınırlı

enable f16; yapsan bile f16 operasyonları daha sınırlı. Bizim kodda f16 sadece storage type olarak kullanılıyor, hesaplama hâlâ f32 (cast-load → compute → cast-store).

#6. Kernel bölme (// --- KERNEL: name ---)

WGSL bir dosyada birden fazla @compute kernel barındırabiliyor ama bizim engine her dosyayı kernel marker'lara göre split edip her kernel için ayrı pipeline state oluşturuyor. Sebep: enable direktifleri ve var<workgroup> bildirimleri başka kernel'a sızmıyor.

#7. Subgroup operations (lar)

enable subgroups; ile subgroupAdd(), subgroupBroadcast() gibi cross-thread reduction'lar kullanılabiliyor — ama Apple GPU'larda Chrome flag arkasında ve wave size = 32 varsayılıyor. Bu runtime feature flag ve adapter desteğine bağlı.

#Notlar

• Bind group layout dinamik değil — her kernel'ın bind group layout'u shader compile sırasında belirleniyor; dispatch sırasında sadece buffer binding'leri değiştirilebilir.
• Multi-tensor AdamW farkı — Metal versiyonu da multi-tensor yapıyor ama WGSL'in entries[].resource.{offset, size} ile sub-range binding desteği gerek; bu standart bir WebGPU özelliği.
• F16 backward yok — backward pass tamamen fp32. Mixed precision sadece forward'da _w16 kernel'larıyla aktif.
• Batch=1 garantili — tarayıcıda WebGPU üzerinde gerçek batch desteği zor olduğu için tüm kernel'lar tek sequence için tasarlandı; gradient accumulation host-side yapılıyor.

#Kaynak Kod Konumu

bpe/src/llm/
├── engine.js              ← WebGPU device + pipeline compilation + bind helpers
├── model.js               ← Model class orchestrator (Object.assign wiring)
├── model/
│   ├── layout.js          ← paramLayout, sub-range, bufOf/offsetOf, FNV
│   ├── alloc.js           ← allocate (mega_w), allocActivations, allocBackward
│   ├── forward.js         ← _encodeLayerForward, forwardLoss
│   ├── backward.js        ← backward (split attn dispatch lives here)
│   ├── optim.js           ← multi-tensor AdamW step
│   ├── decode.js          ← allocKVCache, prefillCache, decodeOne, predictLogitsAt
│   └── checkpoint.js      ← export/import + setStepCount + destroy
└── shaders/               ← Bu etüdlerin konusu

İlk durağa: 00_shared.md →