#cast_f32_to_f16 ve cast_f16_to_f32 — Mixed Precision Conversion

Dosya: 08_cast.wgsl Pipeline adımı: Optimizer step çıkışında çalışır — master fp32 → forward fp16 mirror.

İki kernel:

• cast_f32_to_f16 — saturated clamp + cast (forward weight'lerin fp16 mirror'una)
• cast_f16_to_f32 — basit cast

#Nedir bu ya?

Şöyle düşün: elinde bir RAW fotoğraf var, dosya kocaman ama her detay yerli yerinde. Üzerinde düzenleme yaparken hep o RAW'ı kullanırsın — kayıpsız, sağlam. Ama fotoğrafı web'e koyacaksan o devasa dosyayı atmazsın, ondan küçük bir JPEG çıkarırsın. JPEG biraz detay kaybeder ama yarısı kadar yer kaplar ve çok daha hızlı yüklenir. İşte bu file tam o JPEG'i üretme işini yapıyor.

Burada RAW = f32 (32-bit, 4 byte), JPEG = f16 (16-bit, 2 byte). Model eğitilirken ağırlıkların "master" kopyası hep fp32'de durur, çünkü minik gradient güncellemeleri toplamak için hassasiyet lazım. Ama her forward pass'te o devasa fp32'leri okumak bellek bandwidth'ini boğar. O yüzden yanında bir de fp16 "ayna" tutuyoruz — yarı boyutta, okuması iki kat hızlı. cast_f32_to_f16 her optimizer adımından sonra fp32 master'dan bu fp16 aynayı tazeliyor.

Cast'in kendisi aptal kadar basit — bir sayıyı al, 2 byte'a sığacak şekilde yeniden yaz. double'ı float'a düşürmekten farkı yok aslında. Tek ekstra incelik: fp16'nın taşıyabileceği bir tavan değer var (65504). Bunun üstüne çıkan bir sayıyı cast edersen sonsuza (Inf) dönüşür — o yüzden önce bir clamp ile değeri o sınıra sıkıştırıyoruz, tıpkı aşırı pozlanmış fotoğrafın bembeyaz patlamasını engellemek gibi.

Ters yön (cast_f16_to_f32) ise tamamen dertsiz — küçük kutudan büyük kutuya geçmek hiçbir şey kaybettirmez, JPEG'i tekrar açmak gibi. İşin asıl ilginç kısmı bu basit kopyalamanın 110 kez ayrı ayrı çağrılınca nasıl bir performans derdine dönüştüğü — ona aşağıda geliyoruz.

#Ne Yapar?

Mixed precision pipeline'da master weight'ler fp32 olarak tutulur (numerical stability — Adam moments + small gradient updates). Forward pass için fp16 mirror kullanılır (memory bandwidth tasarrufu).

Optimizer step:

1. AdamW master weight'leri fp32'de günceller
2. cast_f32_to_f16 her parametre tensor'u için çalışır → fp16 mirror'u günceller
3. Sonraki forward pass fp16 mirror'dan okur

cast_f16_to_f32 simetrik — fp16 → fp32. Aslında bizim hot path'te kullanılmıyor (forward kernel'lar tile-load sırasında inline cast yapıyor); standalone helper olarak duruyor.

#Niye Saturated Clamp (±65504)?

f16 max = 65504 (= 2^15 × (2 - 2^-10)). Daha büyük değer cast'te ±Inf'e dönüşür.

Sağlıklı training'de weights |w| < 5 kalır — clamp hiç tetiklenmez. Ama:

• Gradient explosion (rare ama olur)
• Adam state corruption
• Numerical instability spike

→ master fp32'da büyük değer olabilir → naive cast → fp16 Inf → forward NaN propagation.

clamp(±65504) defensive net. Maliyet: 1 instruction per element. Forward pass'in NaN'le kontamine olmasını önler — değer.

#Bind Group ABI

#cast_f32_to_f16 (3 binding)

#cast_f16_to_f32

Tip ters; geri kalan aynı.

#Dispatch Şekli

workgroup_size: 256
threads:        ceil(n / 256) workgroups × 256

1 thread = 1 element. Memory-bound, no shared memory, no barriers.

#Satır Satır

#cast_f32_to_f16

@compute @workgroup_size(256, 1, 1)
fn cast_f32_to_f16(@builtin(global_invocation_id) gid: vec3<u32>,
                   @builtin(num_workgroups) nwg: vec3<u32>) {
    let i = flat_id(gid, nwg);
    if (i >= n) { return; }
    dst[i] = f16(clamp(src[i], -65504.0, 65504.0));
}

• flat_id for 2D dispatch fallback
• clamp to f16 range
• f16(...) cast — WGSL spec built-in (only available with enable f16;)

#cast_f16_to_f32

dst[i] = f32(src[i]);

f16 → f32 her zaman safe (lossless). No clamp gerek.

#cast_f32_to_f16 Hot Spot — 110 Çağrı ve Optimizasyonu

Eski profil çıktılarına göre her parametre tensor için ayrı bir dispatch yapılıyordu:

cast_f32_to_f16 (×110)  8.585 ms total  →  ~1.7% of step

110 = 12 layer × 9 weight per layer (att Q/K/V/O + FFN gate/up/down + 2 norms) + outliers (embed, lm_head, final norm).

Her dispatch küçük (~78 µs) olsa da, biriken dispatch overhead (~110 dispatch × ~50 µs ≈ 5.5 ms) ciddi bir gecikme yapıyordu.

#WGSL-Spesifik Notlar

#1. enable f16; zorunlu

Bu file sadece shader-f16 feature varken yüklenir. engine.js'da:

if (!adapter.features.has('shader-f16')) {
    throw new Error('...');
}

enable f16; shared preamble'a injekte edilir → f16 type her kernel'da kullanılabilir. Safari'de shader-f16 yok → bu file yüklenmez, mixed precision mode disable.

#2. clamp(NaN, -c, c)

WGSL spec'de davranış belirsiz. Apple Metal'de pratikte c (en yakın bound). Pratikte sorun değil — gradient clip ve z_loss zaten NaN'ı engelliyor.

#3. f16 storage layout

array<f16> 16-bit storage — fp16 max 65504. Bizim weight'ler ~5'in altında, abundant headroom.

WebGPU spec: f16 storage packed (2 bytes per element, no padding). Apple Metal hardware half ile aynı — native f16x4 instruction'ları var. NVIDIA/Intel desktop'ta da hardware support.

#Code Review

#Bulgu 1: Clamp every element — wasted bandwidth?

#Bulgu 2: Standalone vs fused

#Bulgu 3: cast_f16_to_f32 kullanılıyor mu?

#Hızlı Kontrol Listesi

#Sonraki

09_backward_ffn.md — FFN activation backward'ları: GeLU + SwiGLU.