#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/_at kernel'ları yapar.

İki kernel:

• ffn_gelu_backward — GeLU activation backward
• ffn_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)

#ffn_swiglu_backward (6 binding)

#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

#Bulgu 2: nan_guard her iki kernel'da

#Bulgu 3: Mixed precision yok bu kernel'larda

#Hızlı Kontrol Listesi

#Sonraki

10_backward_attention.md — pipeline'ın en karmaşık kernel'i. 3 varyant: streaming, split-short, split-dKdV.