llm.istanbul·Etüt
TR EN
Workbench →

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. doublefloat'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)

BindingTürDetay
0storage, readsrc: array<f32>
1storage, read_writedst: array<f16>
2uniformn: u32

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

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

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

İpucu

Tamamlanan Optimizasyon — Fused F16 Mirroring: Bu dispatch yükünü sıfırlamak için cast_f32_to_f16 adımları, optimizer aşamasındaki adamw_update_f16 (veya adamw_8bit_update_f16) kernel'larının içine fuse edilmiştir. adamw_update_f16 güncelleme kernel'ı, fp32 master weights güncellenirken aynı dispatch/pass içinde forward fp16 aynasını da (dst_w16 tamponuna) doğrudan yazar. Böylece standalone cast_f32_to_f16 döngüsü ve 110 ayrı dispatch overhead'i tamamen ekarte edilmiştir.


WGSL-Spesifik Notlar

1. enable f16; zorunlu

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

javascript
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?

RiskAçıklama
🟡 minorSağlıklı training'de hiçbir clamp tetiklenmez. Yine de her element 2 comparison + 2 select yapar. Marjinal CPU overhead. Pratikte memory-bound olduğu için ALU üstü yüke aldırış edilmiyor.

Bulgu 2: Standalone vs fused

RiskAçıklama
🟢 ÇözüldüStandalone cast adımı 110 dispatch/step gerektiriyordu. Bu işlem adamw_update_f16 içerisine fuse edilerek ekstra dispatch yükü sıfırlanmış ve ~%1.5-2.0 adım süresi iyileşmesi sağlanmıştır.

Bulgu 3: cast_f16_to_f32 kullanılıyor mu?

RiskAçıklama
🟢 yokŞu an çağrılmıyor — forward kernel'lar tile-load'da inline f32(weight_w16) yapıyor. Standalone kernel future use için dururor. Bandwidth maliyeti yok.

Hızlı Kontrol Listesi

Test SenaryosuDurum
clamp(weight = 1000) doğru clamp ediyor mu?✅ formula ±65504
f16(weight) → f32 lossless mi?✅ widening cast
f16(weight = 1.0) → f32 = 1.0 exact?✅ powers of 2 exact
f16 mirror master ile periodically synced mi?✅ adamw_update'tan sonra her step
Safari'de yüklemiyor mu?shader-f16 feature gate

Sonraki

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

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