ffn_gelu_backward ve ffn_swiglu_backward — FFN Activation Backward
Dosya: 09_backward_ffn.wgsl Pipeline adımı: Backward FFN'in element-wise kısmı. GEMM kısmı
matmul_t/_atkernel'ları yapar.
İki kernel:
ffn_gelu_backward— GeLU activation backwardffn_swiglu_backward— SwiGLU combine backward
Nedir bu ya?
Forward'da bir aktivasyon fonksiyonu (GeLU, SiLU) sayıyı eğip büküyor: bazı değerleri kısıyor, bazılarını geçiriyor. Sen sadece "girdi şuydu, çıktı buydu" diyorsun. Ama backward'da soru başka: "girdiyi azıcık oynatsaydım, çıktı ne kadar değişirdi?" İşte o "ne kadar"ın cevabı türev — yani fonksiyonun o noktadaki eğimi.
Şöyle düşün: bir gaz pedalına basıyorsun. Pedalı 1 mm itince araba ne kadar hızlanır? Bu, pedalın o anki konumuna bağlı. Boştayken 1 mm hiçbir şey yapmaz (eğim ~0), ortadayken ciddi bir ivme verir (eğim yüksek). Türev tam bunu söylüyor: "şu anki girdide, küçük bir dokunuş çıktıya kaç katı yansır." Backward sadece yukarıdan gelen gradient'i bu çarpanla ölçekliyor — grad_giriş = grad_çıkış × eğim. Zincir kuralı dediğimiz şey bu kadar.
GeLU'da tek bir eğim var, doğrudan çarpıyorsun. SwiGLU biraz daha kurnaz: çıktı iki ayrı girdinin (gate ve up) çarpımından geliyor, o yüzden iki ayrı eğim çıkarıp gradient'i ikiye dağıtıyorsun. İkisi de aynı kernel'da, eleman eleman, komşusundan habersiz — utanç verici derecede paralel.
Tek ince nokta eğimi nasıl hesapladığın. Sigmoid içinde e^x var ve x büyükse bu sonsuza fırlar (overflow → Inf, sonra her şey NaN). Çözüm: işaretine bakıp exp()'ye her zaman negatif argüman verecek şekilde formülü çevirmek. Aşağıda "stabil sigmoid" derken kastettiğimiz bu küçük ama hayati hile.
Ne Yapar?
GeLU backward
Forward: hidden = gelu(pre_gelu). Backward:
grad_pre[i] = grad_hidden[i] · gelu'(pre[i])SwiGLU backward
Forward: hidden = silu(gate) · up. Backward (chain rule):
grad_gate[i] = grad_hidden[i] · up[i] · silu'(gate[i])
grad_up[i] = grad_hidden[i] · silu(gate[i])İki ayrı upstream tensor için iki ayrı gradient hesabı — aynı kernel'da paralel.
Türevler
GeLU' (tanh approximation)
Forward:
gelu(x) = 0.5x · (1 + tanh(inner))
inner = √(2/π) · (x + 0.044715·x³)Türev (chain + product rule):
gelu'(x) = 0.5 · (1 + tanh(inner)) + 0.5·x · sech²(inner) · d(inner)/dx
sech²(z) = 1 - tanh²(z)
d(inner)/dx = √(2/π) · (1 + 3·0.044715·x²)Kod:
fn gelu_derivative(x_in: f32) -> f32 {
// d/dx [0.5 * x * (1 + tanh(inner))] inner = sqrt(2/pi)*(x + 0.044715*x^3)
let x = clamp(x_in, -100.0, 100.0);
let sqrt_2_over_pi = 0.7978845608;
let coeff = 0.044715;
let x2 = x * x;
let x3 = x2 * x;
let inner = sqrt_2_over_pi * (x + coeff * x3);
let t = tanh(inner);
let sech2 = 1.0 - t * t;
let d_inner = sqrt_2_over_pi * (1.0 + 3.0 * coeff * x2);
return 0.5 * (1.0 + t) + 0.5 * x * sech2 * d_inner;
}SiLU'
Forward: silu(x) = x · σ(x) where σ(x) = 1/(1+e^-x) sigmoid.
Türev:
silu'(x) = σ(x) + x · σ(x) · (1 - σ(x))
= σ(x) · (1 + x · (1 - σ(x)))Kod:
fn sigmoid(x: f32) -> f32 {
// Stable sigmoid — branch on sign so exp() only ever sees non-positive argument
if (x >= 0.0) {
return 1.0 / (1.0 + exp(-x));
}
let e = exp(x);
return e / (1.0 + e);
}
fn silu_derivative(x: f32) -> f32 {
let s = sigmoid(x);
return s * (1.0 + x * (1.0 - s));
}Bind Group ABI
ffn_gelu_backward (4 binding)
| Binding | Tür | Detay |
|---|---|---|
| 0 | storage, read | grad_hidden: array<f32> — upstream gradient (FFN'in son matmul'ünden) |
| 1 | storage, read | pre_gelu: array<f32> — forward'da save edilen pre-activation |
| 2 | storage, read_write | grad_pre: array<f32> — output, sonraki matmul backward'a gider |
| 3 | uniform | n: u32 |
ffn_swiglu_backward (6 binding)
| Binding | Tür | Detay |
|---|---|---|
| 0 | storage, read | grad_hidden |
| 1 | storage, read | gate — forward gate output |
| 2 | storage, read | up — forward up output |
| 3 | storage, read_write | grad_gate |
| 4 | storage, read_write | grad_up |
| 5 | uniform | n: u32 |
Dispatch Şekli
workgroup_size: 256
threads: ceil(n / 256) workgroups
n = seq × d_ff (örn 512 × 3072 = 1.57M)Element-wise, no shared memory, no barriers.
Satır Satır
ffn_gelu_backward
let i = flat_id(gid, nwg);
if (i >= n) { return; }
let g = grad_hidden[i];
let pg = pre_gelu[i];
let r = g * gelu_derivative(pg);
grad_pre[i] = nan_guard(r);Trivial chain rule. nan_guard finite-check — pre-activation NaN olabilir (rare), sonsuza propagation engelle.
ffn_swiglu_backward
let i = flat_id(gid, nwg);
if (i >= n) { return; }
let g = grad_hidden[i];
let gt = gate[i];
let u = up[i];
let s = silu(gt);
let dg = g * u * silu_derivative(gt);
let du = g * s;
grad_gate[i] = nan_guard(dg);
grad_up[i] = nan_guard(du);İki gradient paralel:
grad_gate = g · u · silu'(gate)grad_up = g · silu(gate)
silu(gate) ve silu'(gate) aynı gate[i] değeri için iki kez hesaplanıyor (sigmoid içinde). Ufak waste ama fp32 ALU bedava.
Niye pre_gelu save ediliyor?
GeLU non-linear. Backward'da derivative gelu'(x) hesaplamak için forward'daki x değerini bilmek gerek. Ama forward'da gelu_inplace kernel'ı x'i overwrite ediyor (in-place).
Çözüm: forward sırasında pre_gelu buffer'ına bir kopya sakla. Backward okur.
Alternatif: backward'da gelu_inverse(hidden) ile x'i geri hesapla. Ama gelu non-injective — invert imkansız.
Memory cost: seq × d_ff × 4B ek activation. seq=512, d_ff=3072 → 6 MB per layer. 12 layer → 72 MB. Anlamlı ama gerekli.
SwiGLU'da bu sorun yok — gate ve up zaten ayrı buffer'larda korunmuş (matmul output'larıydı, post-activation yazılmadı).
WGSL-Spesifik Notlar
1. tanh() and exp() — built-in
Hardware-accelerated. Apple GPU fast_tanh instruction ~3 cycles.
2. clamp(x, -100, 100) ve Dallanmalı Stabil Sigmoid
GeLU girdisi için clamp(x, -100, 100) koruması sürdürülmektedir (x³ büyümesini kontrol etmek için). Ancak sigmoid/silu hesaplamasında eski clamp(x, -50, 50) tamamen kaldırılmıştır. Bunun yerine, x >= 0.0 kontrolüyle exp(-x) veya exp(x) hesaplanır. exp() argümanının her zaman negatif veya sıfır olması garanti edildiğinden, üstel taşma (overflow → Inf) donanımsal olarak engellenir ve clamp ihtiyacı kalmaz. Geriye kalan underflow durumları ise doğrudan kararlı bir şekilde 0.0 olarak çözülür.
3. Re-using silu() and sigmoid() helpers
Forward kernel ile backward kernel aynı helper'ları kullanır. WGSL preamble injection sayesinde (each kernel module compiled separately, helpers injected via 00_shared.wgsl içine değil, file-local preamble). Bizim kodumda silu ve silu_derivative aynı dosyanın file preamble'ında.
Code Review
Bulgu 1: silu(gt) ve silu_derivative(gt) aynı sigmoid(gt) hesaplıyor
| Risk | Açıklama |
|---|---|
| 🟢 minor | Optimization fırsatı: let s = sigmoid(gt); let silu_v = gt * s; let silu_d = s * (1 + gt*(1-s));. Sigmoid bir kez. ~%5 bandwidth'siz savings. Pratikte fp32 ALU bedava, fark imperceptible. |
Bulgu 2: nan_guard her iki kernel'da
| Risk | Açıklama |
|---|---|
| 🟢 yok | Defensive. gelu_derivative veya silu_derivative extreme input'ta NaN üretebilir; downstream gradient'i kirletmek yerine 0'a çek. |
Bulgu 3: Mixed precision yok bu kernel'larda
| Risk | Açıklama |
|---|---|
| 🟢 yok (architectural) | Backward fp32 standardı. pre_gelu, gate, up forward'da fp32 aktivasyonlardan geliyor. Mixed precision sadece weight tarafında, activation tarafında değil. |
Hızlı Kontrol Listesi
| Test Senaryosu | Durum |
|---|---|
gelu_derivative(0) ≈ 0.5 mi? | ✅ formula |
silu_derivative(0) ≈ 0.5 mi? | ✅ formula |
| Negatif input için derivative pozitif mi? | ✅ |
pre_gelu forward'da save ediliyor mu? | ✅ host-side |
NaN injection'da nan_guard çalışıyor mu? | ✅ self-check |
| GeLU forward+backward → identity gradient (small steps)? | ⚠ formal test yok |
Sonraki
10_backward_attention.md — pipeline'ın en karmaşık kernel'i. 3 varyant: streaming, split-short, split-dKdV.