#embed_lookup ve embed_backward — Token ID → Vektör

Dosya: 01_embedding.wgsl Pipeline adımı: 0 — Forward'ın ilk adımı, token'ları sayısal vektörlere çevirir.

3 kernel: embed_lookup, embed_backward, embed_lookup_w16 (mixed precision varyantı).

#Nedir bu ya?

Diyelim elinde "kedi" kelimesi var. Bilgisayar "kedi"yi anlamaz, sadece sayı anlar. Tokenizer "kedi"ye bir numara verdi: 4213. Ama 4213 tek başına aptal bir numara — 4214'ten ("köpek" olsun) büyük olması "kedi > köpek" gibi saçma bir şey ifade etmiyor; numaralar gelişigüzel.

Embedding tam burada giriyor: her token numarasını, anlam taşıyan bir sayı vektörüne çeviriyor. 4213 → [0.2, -1.1, 0.7, …] (mesela 768 sayı). Bu vektör, modelin "kedi" hakkında bildiği her şeyin sıkıştırılmış hâli. Eğitim ilerledikçe "kedi" ile "köpek"in vektörleri birbirine yaklaşır, "kedi" ile "kamyon"unki uzaklaşır.

Kernel ne mi yapıyor? Aslında dev bir lookup table. [vocab × 768]'lik bir tablo düşün, her satır bir token'ın vektörü. embed_lookup resmen table[token_id] çekiyor — array indexleme, hash map'ten değer alma gibi. Tek fark: GPU'da binlerce token için aynı anda.

Geri tarafta (backward) iş tersine: "bu token'ın vektörünü şu yöne itersek loss düşer" bilgisini tabloya geri yazıyoruz. Ufak incelik: aynı token cümlede 10 kez geçtiyse 10 thread aynı satıra aynı anda yazmaya çalışır → "atomik" yazma gerekir (aşağıda).

#Ne Yapar?

#Forward: embed_lookup

Modelin ilk forward adımı. Token ID'ler için embedding tablosundan satır okur:

out[s, d] = table[tokens[s], d]

Eğer tokens[s] >= vocab_size (bozuk girdi) → out[s, d] = 0.

#Backward: embed_backward

Scatter-add gradient. Her thread bir output position için, kaynak token'ın grad satırına ekler. Aynı token birden fazla pozisyonda görünebileceği için atomik CAS-add gerekir:

grad_table[tokens[s], d] += grad_out[s, d]   ∀ (s, d)

#Mixed-precision: embed_lookup_w16

Tablo f16 storage'dan okur, output f32. Backward'da f16 yok — gradient her zaman fp32 (numerical stability).

#Matematiksel Tanım

#Forward

out[s, d] = table[tokens[s], d]    if tokens[s] < vocab_size
            0                        otherwise

#Backward

∂L/∂table[t, d] = Σ_s [tokens[s] == t] · ∂L/∂out[s, d]

Bu literal "sparse scatter-add" — eğer token t toplam k farklı pozisyonda geçiyorsa o satıra k farklı katkı yazılır. Aynı slot'a yazan thread'ler race condition üretir → atomik gerek.

#Bind Group ABI

#embed_lookup (4 binding)

#embed_backward (4 binding)

Dikkat: grad_table deklarasyonu array<atomic<u32>>. Aslında veri f32 ama WGSL atomic<f32> yok — bit-cast trick (aşağıda detay).

#embed_lookup_w16 (4 binding, _w16 suffixli)

#Dispatch Şekli

workgroup_size: 256
total threads:  ceil(seq_len × d_model / 256) workgroups × 256

Bir thread = bir (s, d) çifti. Yani embedding output'un her elemanı için bir thread.

Örnek (seq=512, d=768): 393K thread, 1536 WG.

#Satır Satır — embed_lookup

#1) Bindings + entry

@compute @workgroup_size(256, 1, 1)
fn embed_lookup(@builtin(global_invocation_id) gid: vec3<u32>,
                @builtin(num_workgroups) nwg: vec3<u32>) {
    let i = flat_id(gid, nwg);
    let seq_len = dims.x;
    let d_model = dims.y;
    let vocab_size = dims.z;
    let total = seq_len * d_model;
    if (i >= total) { return; }

• i global flat index, [0, seq×d_model)
• dims.x/y/z deconstruct — vec4<u32> ile aldık çünkü 16-byte alignment'a düşer
• Bounds check — son WG'da boş thread'ler çıkar

#2) (s, d) decode

let s = i / d_model;
let d = i % d_model;

Flat index → 2D (sequence position, dimension within embedding).

Performance note: % d_model divide-and-modulo — derleyici bir kerede hesaplar (udiv instruction). Eğer d_model = 2^k ise bit-shift'e optimize olur (768 değil, ama 1024 olsa).

#3) Token lookup + bounds

let t = tokens[s];
if (t >= vocab_size) {
    out[i] = 0.0;
    return;
}

• Eğer token bozuksa (örn padding token bizim vocab'a girmiyor) → çıktıyı sıfırla.
• Bu defansif — sağlıklı corpus'ta hiç tetiklenmemeli, ama vocab fingerprint mismatch'inde silent corruption olmaz.

#4) Tablo lookup

out[i] = table[t * d_model + d];

Row-major: table[t * d_model + d] token t'nin d'inci dim'i.

Memory pattern: Adjacent thread'ler out[i] ve out[i+1]'e yazıyor → coalesced. Ama table[t*d + d] pattern'i t'ye bağlı → adjacent thread'ler eğer aynı s'de ise (yani d farklılaşıyorsa) coalesced okuma; ama s farklı thread'lerde aynı d varsa scattered okuma. Pratikte WG=256 thread'leri i = base..base+255 aralığında ardışık i alır → genelde aynı s, farklı d → coalesced.

#Satır Satır — embed_backward (atomik CAS-add)

#1) İlk kısım embed_lookup ile aynı

let i = flat_id(gid, nwg);
// ... bounds check, decode, token lookup
let val = grad_out[i];
if (val == 0.0 || !is_finite(val)) { return; }
let dst_idx = t * d_model + d;

is_finite ile NaN/Inf gradient'ı atla — overflow propagation engellemesi.

#2) CAS-loop f32 atomic add

var old_bits = atomicLoad(&grad_table[dst_idx]);
loop {
    let new_bits = bitcast<u32>(bitcast<f32>(old_bits) + val);
    let res = atomicCompareExchangeWeak(&grad_table[dst_idx], old_bits, new_bits);
    if (res.exchanged) { break; }
    old_bits = res.old_value;
}

Niye böyle karmaşık? WGSL'in atomik fonksiyonları sadece i32 ve u32 üzerinde çalışıyor. Float atomik add yok. Çözüm:

1. Bit-cast f32'yi u32 olarak yorumla (aynı 32-bit pattern, farklı interpretation)
2. atomicLoad → old_bits (mevcut değer u32 olarak)
3. Hesapla new_bits = bits(old_as_f32 + val) — yani u32 → f32 cast → ekle → f32 → u32 cast
4. atomicCompareExchangeWeak:
   • eğer slot hâlâ old_bits ise → new_bits yaz, exchanged = true
   • eğer başka thread araya girip değiştirdiyse → exchanged = false, gerçek değeri res.old_value'da döndür
5. exchanged ise loop'tan çık; aksi halde yeni old_bits'le tekrar dene

Bu lock-free retry pattern. Race olmadığı sürece tek iter'da biter; race varsa k thread aynı slot'a yazıyorsa O(k²) toplam denemeye düşer.

Pratikte ne kadar contention?

• Türkçe corpus, en sık token "▁ve" ~%2 frekansta. Yani 512 seq'de ~10 kez geçer.
• Her geçiş bir thread → 10 thread aynı row'a yazıyor.
• 10² = 100 başarısız CAS olabilir worst-case ama paralel olarak bekleştikleri için pratik latency ~10 cycle/atomik.
• Önemli darboğaz değil.

#3) Niye bitcast<f32>(0) aslında 0.0

WGSL'de bitcast<f32>(0u) → 0.0f (zero-init grad_table sıfır f32 demek). Init için fill_zero kullanılıyor — sıfır f32 = sıfır u32, bit-pattern aynı.

#embed_lookup_w16 Farkı

out_w16[i] = f32(table_w16[t * d_model + d]);

Tek fark: table_w16 array<f16> ve okuma sırasında f32 cast. Tablo yarı boyutta (12.6MB → 6.3MB for 16K×384), bandwidth tasarrufu.

Niye _w16 versiyonun backward'ı yok?

Backward gradient hep f32 (mixed precision standardı). f32 grad'ı f16 weight'e ekleyip yazsan precision kaybı. Onun yerine:

1. f32 master copy korur
2. Forward için f16 mirror yapar (cast_f32_to_f16 kernel)
3. Backward f32 master'a yazar

Bu yüzden embed_lookup ve embed_lookup_w16 ayrı; embed_backward tek.

#WGSL-Spesifik Notlar

#1. atomic<f32> yok

Spec'te yok. CAS-bit-cast pattern bilinen workaround. Vulkan ve D3D12'de native atomicAdd var ama WebGPU portability için yok. Apple Metal'de native atomic_fetch_add (f32) var ama WebGPU'nun en küçük ortak paydası bu.

#2. bitcast<T>(x) — non-mutating

bitcast<u32>(1.0f) → 0x3f800000. Aynı 32 bit, farklı interpretation. Maliyet: yok (no-op CPU/GPU'da).

#3. atomicLoad ve atomicCompareExchangeWeak — pointer şart

atomicLoad(&grad_table[dst_idx])

Pointer-to-storage burada legal çünkü ilk argüman özel. Genelde WGSL pointer'ı fonksiyona geçirmek yasak (uniform analysis). Atomik fonksiyonlar exception.

#4. array<atomic<u32>> — type level wrap

Storage buffer deklarasyon seviyesinde atomic olarak işaretlenmek zorunda:

@group(0) @binding(2) var<storage, read_write> grad_table: array<atomic<u32>>;

Bunu array<u32> yazarsan atomicLoad/CompareExchangeWeak derleme hatası verir.

#5. loop { ... break; } — for-yapısı yerine

WGSL'de loop {} infinite loop, break ile çıkılır. for (;;) veya while(true) yok. CAS retry'da idiomatic kullanım.

#Code Review

#Bulgu 1: Forward'da out[i] = 0 OOB token için silent

Karar: Defensive zeroing kalacak. Vocab mismatch'i preflight check yakalıyor.

#Bulgu 2: is_finite filter sadece backward'da

#Bulgu 3: Atomik CAS başarısızlık metric'i yok

#Hızlı Kontrol Listesi

#Sonraki

02_norm.md — RMSNorm forward + backward. Embedding çıktısını L2-normalize eder.