#rope_q ve rope_k — Rotary Position Embedding (Forward)

Dosya: 04_rope.wgsl Pipeline adımı: Q ve K projeksiyon'undan hemen sonra, attention'dan hemen önce.

Forward kernel'ları:

• rope_q — Query'ye rotasyon
• rope_k — Key'e rotasyon (n_kv_heads ile, GQA için)

Backward kernel'ları (rope_q_backward, rope_k_backward) 10_backward_attention.md'da incelenir — birlikte çalıştıkları için.

#Nedir bu ya?

Diyelim elinde bir cümle var ve modelin her kelimenin cümlede nerede durduğunu bilmesi lazım. "Köpek adamı ısırdı" ile "Adamı köpek ısırdı" aynı kelimeler ama sıra her şeyi değiştiriyor. Sorun şu: attention, kelimelerin vektörlerini bir torbaya atıp birbirleriyle kıyaslar — torbada sıra yok. Pozisyonu bir şekilde vektörün içine sokman gerek.

RoPE'nin numarası şu: her vektörü bir saatin akrebi gibi düşün. Token cümlede ne kadar ileriyse, o akrebi o kadar çok döndürüyorsun. 1. kelime biraz dönmüş, 5. kelime daha çok, 50. kelime epey. Açının kendisi pozisyonu kodluyor — ayrı bir pozisyon tablosu tutmuyorsun, bilgiyi doğrudan vektörü çevirerek gömüyorsun.

İşin güzel kısmı kıyaslamada ortaya çıkıyor. İki kelimenin akreplerini kıyasladığında (attention'ın yaptığı dot product) önemli olan akreplerin mutlak yönü değil, aralarındaki açı farkı — yani kaç pozisyon uzaktalar. İki akrebi de aynı miktar daha döndürsen, aralarındaki fark yine değişmez. Böylece model "bu iki kelime 4 adım arayla" bilgisini bedavaya kapıyor, kelimeler cümlenin başında da olsa sonunda da.

Tek bir akrep yerine vektörü ikişerli çiftlere bölüp her çifti farklı hızda döndürmek (bazıları yavaş, bazıları hızlı tik tik) işin biraz dönüp duran kısmı — ama özü bu: konum bilgisi = dönme miktarı.

#Ne Yapar?

LLaMA-style rotary position embedding. Q ve K vektörlerinin her ardışık (d, d+1) çiftini token'ın pozisyonuna göre 2D rotasyonla döndürür. Position information modele böyle aktarılıyor — pozisyonel bilgi vector'ün kendisinde, ek embedding tablosu yok.

Klasik mantığa göre matematik:

Pair i (içinde head, d_pair = 2i, 2i+1):
  freq  = base^(-2i / head_dim)             ← pair index'e göre frekans
  angle = (s + pos_offset) · freq           ← position'a göre açı
  q[2i]'   = q[2i] · cos(angle) - q[2i+1] · sin(angle)
  q[2i+1]' = q[2i] · sin(angle) + q[2i+1] · cos(angle)

Bu 2D rotation matrix dot product:

[q'[2i]  ]   [ cos -sin]   [q[2i]  ]
[q'[2i+1]] = [ sin  cos] · [q[2i+1]]

Çift sayıdaki pair'lar i = 0, 1, 2, ... her biri farklı frekansta:

• i=0 (en düşük dim): freq=1 → low-frequency rotation, slow position-axis traversal
• i=head_dim/2-1 (en yüksek dim): freq=base^(-1+2/head_dim) → high-frequency

Bu spektrum sayesinde model relative position'u attention dot product içinde doğal olarak yakalar:

Q_s · K_s' = (rotated_q_s) · (rotated_k_s') = q_unrotated · k_unrotated · f(s - s')

Yani Q⋅K hesabı mutlak konumdan bağımsız, sadece relative position s−s' fonksiyonu olur. Detail: RoFormer paper.

#Bind Group ABI

#rope_q (3 binding)

#rope_k (3 binding)

Aynısı; n_heads yerine n_kv_heads. GQA için n_kv_heads < n_heads.

Why pos_offset? Decoder use-case'inde, KV cache'le devam ediyorsan pos_offset ile cache'deki tokens'a denk gelen position offset'i veriyorsun. Training'de pos_offset = 0.

#rope_base ne?

Default LLaMA-2'de 10000. Bu frekans spektrumunun "dağılımı"nı belirler — büyük base = daha geniş low-frequency aralık = daha uzun maksimum sequence length yakalar.

LLaMA-3 ve Mistral'da context window büyüdükçe base = 50000 veya 1000000 deniyor. Bu kernel base'i runtime parameter, değiştirilebilir.

#Dispatch Şekli

Bir thread = bir (s, h, p) çifti (s: seq pos, h: head, p: pair index).

threads per kernel:  seq_len × n_heads × (head_dim/2)
workgroup_size:      256
total WG:            ceil(threads / 256)

Örnek (seq=512, n_heads=12, head_dim=64):

• pairs per token = n_heads × head_dim/2 = 12 × 32 = 384
• total threads = 512 × 384 = 196,608
• WG = 768

#Satır Satır — rope_q

#1) Index decode

let i = flat_id(gid, nwg);
let half_d = head_dim / 2u;
let pairs_per_pos = n_heads * half_d;
let total = seq_len * pairs_per_pos;
if (i >= total) { return; }

let s = i / pairs_per_pos;
let rest = i % pairs_per_pos;
let h = rest / half_d;
let p = rest % half_d;

Flat index → 3D (s, h, p). p = pair index içinde head (0..head_dim/2-1).

#2) Frequency hesapla

let freq = rope_freq(p, head_dim, base);
// freq = exp(-log(base) * 2*p / head_dim) = base^(-2p/head_dim)

exp(-log(base) * 2p/d) formülü base^(-2p/d) ile aynı ama WGSL'in pow(b, x) daha stabil olur ki spec emin değil. exp(-log(base) * x) her zaman çalışır.

#3) Açı + sin/cos

let angle = f32(s + pos_offset) * freq;
let c = cos(angle);
let sn = sin(angle);

pos_offset = 0 training'de. Decoder'da KV cache offset'i.

cos/sin WGSL built-in. Apple GPU'da hardware acceleration var (fast_sin, fast_cos instruction).

#4) Pair indeksleri

let row_base = s * (n_heads * head_dim);
let head_base = row_base + h * head_dim;
let i0 = head_base + 2u * p;
let i1 = i0 + 1u;

Q[i0] = Q[s, h, 2p], Q[i1] = Q[s, h, 2p+1].

#5) Rotation in-place

let q0 = Q[i0];
let q1 = Q[i1];
Q[i0] = q0 * c - q1 * sn;
Q[i1] = q0 * sn + q1 * c;

Klasik 2D rotation. q0 ve q1'i önce register'a kopyala, sonra yaz — yoksa Q[i0]'i yazıp Q[i1]'i hesaplarken q0 artık eski değerin değil yenisinin olur, race condition içinde kendiyle.

Niye in-place? Q ve K büyük tensor'lar (seq × n_heads × d_head). Ayrı output buffer kullansak bandwidth doublerdı. RoPE element-wise rotation, mainpass içinde fused değil — ayrı kernel ama in-place sayesinde bandwidth sorun değil.

#rope_k Farkı

Sadece head sayısı:

• rope_q n_heads (örn 12)
• rope_k n_kv_heads (örn 4 — GQA)

Aksi her şey aynı. Q ve K paralel olarak rotate ediliyor — same angle (same s), same pair index, same frequency. Bu eşitlik kritik: attention'da Q·K dot product'ı hesaplandığında same-pair rotasyonlar birbirini götürür.

Q_s = R(s) · q_unrotated
K_s' = R(s') · k_unrotated
Q_s · K_s' = q · R(s)^T · R(s') · k = q · R(s' - s) · k

R(s'-s) sadece relative offset s'-s'e bağlı. Bu RoPE'nin ana gücü.

#RoPE Backward — Önemli not

WGSL kodumda backward kernel'ları 10_backward_attention.wgsl içinde:

• rope_q_backward (line 499)
• rope_k_backward (line 544)

Backward formula:

∂L/∂q[2i]   = (∂L/∂q'[2i])  · cos + (∂L/∂q'[2i+1]) · sin
∂L/∂q[2i+1] = (∂L/∂q'[2i+1]) · cos - (∂L/∂q'[2i])  · sin

Yani forward'un transpose (= rotation matrix'in inverse'i = same matrix with -angle). Detay 10_backward_attention.md bölümünde.

#WGSL-Spesifik Notlar

#1. cos(angle) precision

WGSL cos/sin IEEE round-to-nearest. Apple Metal'de fast-math açıksa precision düşebilir, ama bizim kullanım için yeterli (relative position rotation, küçük floating error attention output'a noise olarak yansır, gradient signal-to-noise oranı yine yüksek).

#2. exp ve log overhead

Her thread exp(-log(base) * 2p/d) hesaplıyor — p farklı, ama log(base) constant. Compiler log(base)'i hoisting yapabilir ama garanti değil. Alternatif: host'tan precomputed table geçirmek (örn freq[p] lookup table). Bizim kod runtime'da hesaplıyor — basitlik için.

#3. In-place write race

let q0 = Q[i0];
let q1 = Q[i1];
Q[i0] = q0 * c - q1 * sn;  // safe — q0/q1 already in registers
Q[i1] = q0 * sn + q1 * c;

Eğer önce Q[i0] = ... sonra q1 = Q[i1] yapsaydım, q0'in eski değerini kaybetmiş olurdum (kendi yazdığım yeni değeri kullanıyordum). Doğru sıralama: iki read, sonra iki write.

#Code Review

#Bulgu 1: freq lookup table optimization fırsatı

#Bulgu 2: pos_offset her zaman 0 (training)

#Hızlı Kontrol Listesi

#Sonraki

05_attention.md — modelin kavramsal kalbi: scaled dot-product attention. RoPE'd Q ve K, V ile birlikte.