#gelu_inplace ve swiglu_combine — FFN Activations

Dosya: 06_activation.wgsl Pipeline adımı: FFN orta hesaplaması — gate output'u activation'dan geçirip up output'la birleştirir.

İki kernel:

• gelu_inplace — GeLU activation, in-place
• swiglu_combine — Fused silu(gate) * up

#Nedir bu ya?

Bir matmul'den çıkan ham sayılar var elinde — kimi pozitif, kimi negatif, hepsi düz bir doğru üzerinde. Eğer bunları olduğu gibi bir sonraki katmana geçirirsen, üst üste koyduğun bütün katmanlar tek bir büyük çarpan gibi davranır; ne kadar derin olursa olsun model hâlâ düz bir çizgi çizer. Aktivasyon fonksiyonu tam bu noktada devreye girip o düzlüğü kıran şey: kıvrım ekliyor, böylece model eğri büğrü şeyleri öğrenebiliyor.

En sezgili düşünme şekli: aktivasyon bir dimmer (kısık ışık) düğmesi gibi. Sert bir aç/kapa anahtarı değil — "şu sinyali ne kadar geçireyim?" diye yumuşakça karar veriyor. GeLU bunu yapıyor: küçük negatif değerleri tamamen kesmek yerine azıcık geçiriyor, büyük pozitifleri neredeyse olduğu gibi bırakıyor. Eski ReLU'nun ("0'ın altındaysa direkt kes") yumuşatılmış, daha kibar hâli diyebilirsin.

SwiGLU ise işi bir adım öteye taşıyor. Tek bir sinyal yerine iki paralel sinyal üretiyorsun: biri up (asıl içerik), öbürü gate (kapı bekçisi). Gate'i bir dimmer'dan geçirip (SiLU) sonra up ile çarpıyorsun — yani bir sinyal, diğerinin ne kadarının geçeceğini belirliyor. Tıpkı bir ses mikserindeki kanal: birinci fader sesin kendisi, ikinci fader o sesi kısıp açan kol. Bu öğrenilebilir kapı, modele "şu bağlamda bu özelliği öne çıkar, şunu bastır" esnekliği veriyor; LLaMA, Mistral, Gemma hep bunu kullanıyor.

Kernel'lerin yaptığı iş aslında basit: her eleman için ufak bir matematik formülü uygula. Asıl ince nokta burada değil — exp() gibi fonksiyonlar büyük girdilerde sonsuza fırlayıp NaN üretebiliyor. O yüzden saf formül yetmiyor; aşağıda göreceğin gibi işaret bazlı dallanma ile sayısal olarak stabil bir SiLU kuruluyor.

#Ne Yapar?

#GeLU

Gaussian Error Linear Unit. Smooth, non-monotone (negatiflerde minik dip vardır):

gelu(x) ≈ 0.5 · x · (1 + tanh(√(2/π) · (x + 0.044715 · x³)))

Bu tanh approximation. Exact GeLU x · Φ(x) where Φ standart normal CDF — pahalı. Tanh form hem hızlı hem ~%99 doğruluk.

#SwiGLU

LLaMA-style FFN'in kalbi. İki ayrı projection (gate, up) → SiLU(gate) ile up'ı element-wise çarp:

silu(x) = x · sigmoid(x)
swiglu_combine(gate, up): out[i] = silu(gate[i]) · up[i]

SwiGLU FFN tam olarak:

hidden = swiglu_combine(W_gate(x), W_up(x))     ← 2 matmul + bu kernel
out = W_down(hidden)                             ← 1 matmul

vs. GeLU FFN:

hidden = gelu(W_up(x))                           ← 1 matmul + gelu
out = W_down(hidden)                             ← 1 matmul

SwiGLU 1 matmul daha pahalı ama daha iyi quality. LLaMA, Mistral, Falcon, Gemma hepsi SwiGLU kullanır. ~%5-10 daha düşük loss.

#Matematik

#GeLU

gelu(x) = 0.5 · x · (1 + tanh(√(2/π) · (x + 0.044715 · x³)))

with constants:
  √(2/π) ≈ 0.7978845608
  0.044715 (paper-derived coefficient)

#SiLU (a.k.a. Swish-1)

silu(x) = x · sigmoid(x) = x / (1 + exp(-x))

#SwiGLU combine

swiglu_combine(g, u)[i] = silu(g[i]) · u[i]
                       = (g[i] / (1 + exp(-g[i]))) · u[i]

gate (g) ve up (u) matmul'lerden geliyor — typical FFN intermediate dim 4× d_model. Yani d_ff=3072 boyutunda her ikisi.

#Bind Group ABI

#gelu_inplace (2 binding)

#swiglu_combine (4 binding)

#Dispatch Şekli

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

n = seq × d_ff (örneğin 512 × 3072 = 1.57M). 6144 WG.

Bir thread = bir element. Tamamen parallel, no shared memory, no barriers.

#Satır Satır

#gelu_inplace

@compute @workgroup_size(256, 1, 1)
fn gelu_inplace(@builtin(global_invocation_id) gid: vec3<u32>,
                @builtin(num_workgroups) nwg: vec3<u32>) {
    let i = flat_id(gid, nwg);
    if (i >= n) { return; }
    x[i] = gelu(x[i]);
}

Trivial — flat_id, bounds check, in-place transform.

gelu() helper:

fn gelu(x: f32) -> f32 {
    let xc = clamp(x, -100.0, 100.0);
    let inner = 0.7978845608 * (xc + 0.044715 * xc * xc * xc);
    return 0.5 * xc * (1.0 + tanh(inner));
}

clamp(-100, 100) neden? x³ çok büyük girdilerde overflow. x=100 → x³ = 1M, x=300 → x³ = 27M, x=1000 → x³ = 1e9, x=10000 → x³ = 1e12. f32 max ~3.4e38, hâlâ güvenli. Ama 0.044715 * x³ daha sonra tanh parametresi → tanh(very_large) = ±1 → fonksiyon zaten satured. Clamp erken çıkış, NaN propagation engelliyor.

tanh() WGSL built-in. Apple GPU hardware'ında native instruction.

#swiglu_combine

@compute @workgroup_size(256, 1, 1)
fn swiglu_combine(@builtin(global_invocation_id) gid: vec3<u32>,
                  @builtin(num_workgroups) nwg: vec3<u32>) {
    let i = flat_id(gid, nwg);
    if (i >= n) { return; }
    hidden[i] = silu(gate[i]) * up[i];
}

silu() helper:

fn silu(x: f32) -> f32 {
    var sig: f32;
    if (x >= 0.0) {
        sig = 1.0 / (1.0 + exp(-x));
    } else {
        let e = exp(x);
        sig = e / (1.0 + e);
    }
    return x * sig;
}

• Dallanmalı Stabilizasyon: Eski yaklaşımdaki clamp(x, -50, 50) yerine, x'in işaretine göre dallanma yapılır. Bu sayede exp() fonksiyonunun argümanı her zaman polarize negatif (exp(-x) veya exp(x)) kalır. exp()'in pozitif argüman görerek Inf (sonsuz) üretmesi ve ardından Inf/Inf belirsizliğiyle NaN oluşturması donanımsal düzeyde tamamen engellenir (underflow durumu ise doğrudan sıfıra yuvarlanarak kararlı çalışır). Bu yaklaşımla clamp sınırlamalarına ihtiyaç kalmaz ve büyük girdilerde silu(x) -> x asimptotiği kusursuz korunur.

Pratikte SwiGLU input'u (matmul output) ~[-3, 3] aralığında olur normal training'de. Bu yapı, olağandışı gradyan sıçramalarında bile pipeline'ın NaN ile kontamine olmasını engeller.

#Niye Fused swiglu_combine?

Klasik PyTorch:

gate_out = gate_proj(x)
up_out = up_proj(x)
silu_out = F.silu(gate_out)
hidden = silu_out * up_out

= 4 ayrı kernel call (silu in-place + multiply ayrı, veya silu + multiply fused).

Bizim swiglu_combine:

hidden[i] = silu(gate[i]) · up[i]

= 1 kernel call. Memory traffic:

• Klasik: silu in-place (1 read + 1 write of gate) + multiply (2 reads + 1 write) = 5 op
• Fused: 2 reads (gate, up) + 1 write (hidden) = 3 op → 40% bandwidth tasarrufu

Plus 1 daha az dispatch overhead.

#WGSL-Spesifik Notlar

#1. tanh() ve exp() built-in

WGSL spec garantisi. Hardware'a optimize olur (Apple GPU'da native).

#2. clamp(x, lo, hi) — saturated arithmetic

clamp(NaN, a, b) davranışı WGSL spec'inde belirsiz. Apple Metal'de "ya NaN ya bound, çoğu zaman bound". Bizim kullanım için fark etmiyor (gradient already finite check yapılıyor backward'da).

#3. In-place vs separate output

gelu_inplace adı in-place imayı yapıyor — gerçekten input'un üzerine yazıyor. Eski değer geri kazanılamaz. Backward'da gerekirse forward'da extra copy buffer tutmak zorunda kalırız (bizim kod öyle yapıyor — pre_act buffer).

swiglu_combine ayrı output buffer (hidden). Backward için gate ve up korunmuş kalıyor — ek save yok.

#Performance

Profil ölçümlerinden:

• swiglu_combine: 12 layer × ~186 µs = 2.2 ms total = 0.4% of step

Çok ufak. Element-wise, memory-bound, hot kernel değil.

#Code Review

#Bulgu 1: GeLU clamp ±100 belki gereksiz

#Bulgu 2: gelu_inplace backward için pre-act gerek

#Bulgu 3: SwiGLU pre-activation save yok

#Hızlı Kontrol Listesi

#Sonraki

07_loss.md — Cross-entropy + sum_losses. Forward'ın son adımı, backward'ın başlangıcı.