llm.istanbul·Etüt
TR EN
Workbench →

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 (loss07 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ı:

BoyutMetalWGSL
RuntimeApple GPU (native)tarayıcı (Chrome/Edge)
DilMetal Shading LanguageWebGPU Shading Language
Memory modelunified (CPU+GPU)mostly-uniform
Threadgroup syntaxthreadgroupvar<workgroup>
Barrierthreadgroup_barrierworkgroupBarrier()
Atomikatomic_* (built-in)atomic*<u32>, atomicAdd, vb.
f16 desteğivarsayılanenable f16; extension
SubgroupSIMD-group nativeenable subgroups; extension

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.

#EtüdDosyaÖzet
00aShared Preamble00_shared.wgslEngine her kernel'ın başına injekte eder: WG boyutu, reduction, NaN guard, char_class.
00bInfrastructure Kernels00_infrastructure.wgslTutkal kernel'lar: fill_zero, fill_const, scale, axpy, copy, clamp_inplace.

I. Forward Pass — Token'dan Tahmine

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

#EtüdShaderÖzet
01Embedding Lookup01_embedding.wgslToken ID → vektör. f32 ve f16 (_w16) varyantları.
02RMS Norm02_norm.wgslVektörü L2-normuna göre normalize et. Subgroup reduction.
03Linear Forward03_linear.wgsl (forward kernel'ları)Y = X @ W — 64×64 double-buffered tile, vec4 loads, fused SwiGLU forward.
04RoPE04_rope.wgslPozisyona göre Q/K rotate. Forward + backward (involution).
05Attention Forward05_attention.wgslOnline softmax + GQA + KV cache + decode + seg çapraz döküman maskeleme.
06Activation06_activation.wgslGeLU, SwiGLU combine (stabil dallanmalı SiLU).
07Cross-Entropy Loss07_loss.wgslLog-softmax + NLL fused. Backward'ın başlangıcı.
08F16 Cast08_cast.wgslf32 ↔ f16. Standalone mixed precision dönüşüm (cast ekarte edilmiştir).

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

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

#EtüdShaderÖzet
09Backward FFN09_backward_ffn.wgslGeLU ve SwiGLU activation backward (stabil sigmoid türevi).
10Backward Attention10_backward_attention.wgsl3 varyant: streaming, split-short, split-dKdV + seg çapraz döküman maskeleme.
11Backward Linear03_linear.wgsl (backward kernel'ları)dW = X^T @ dY, dX = dY @ W^T. matmul_t, _at, fused SwiGLU backward.

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

#EtüdShaderÖzet
12AdamW Optimizer12_optimizer.wgslMulti-tensor AdamW + fused fp16 mirroring (adamw_update_f16). fp32 + 8-bit varyantı.

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'ı.

#EtüdShaderÖzet
13BPE Tokenizer Pipelinebpe.wgslKelime sınırı tespiti, 2-seviyeli local hash count, subgroup-kooperatif scan ve Trie-tabanlı greedy longest-match.

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

MetrikDeğer
Shader dosyası14 (00–13, w16/8bit/BPE varyantlarıyla birleşik)
Etüd dosyası15 + index
Pipeline coverage%100
Forward kernel11 unique
Backward kernel13 unique
Optimizer kernel5 unique (3 fp32, 1 8bit, 1 grad-stat)
Tokenizer kernel14 unique (eğitim ve çıkarım)

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

wgsl
@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)

wgsl
@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

wgsl
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

WGSL kernel etüdleri · WebGPU üzerinde sıfırdan LLMİstanbul’da Uğur Toprakdeviren tarafından hazırlandı.