`matmul` ve `matmul_residual` — Tiled Matrix Multiply (Forward)

Dosya: 03_linear.wgsl (forward kernel'ları) Pipeline adımı: Hemen her layer'da. Attention Q/K/V/O projection, FFN gate/up/down, lm_head — modelin en pahalı işlemleri.

Forward kernel'ları:

matmul — Y = X @ W (Çift arabelleğe alma ve vec4 yükleme ile)
matmul_residual — Y = X @ W + R (residual fused)
matmul_residual_swiglu_a — Y = (silu(GATE) * UP) @ W + R (FFN SwiGLU ve matmul forward fusion)
matmul_w16 — f16 weight versiyonu (mixed precision)
matmul_residual_w16
matmul_residual_swiglu_a_w16 — f16 weight ile Fused FFN forward

Backward kernel'ları (matmul_t, matmul_t_acc, matmul_at, matmul_at_acc, matmul_at_swiglu_a, matmul_at_acc_swiglu_a) 11_backward_linear.md'da incelenir.

Nedir bu ya?

Bir restoran mutfağı düşün. Her tabak (output satırı) için, elindeki malzemeleri (girdi vektörü X'in bir satırı) bir sürü tarifle (ağırlık matrisi W'nin sütunları) eşleştirip karıştırıyorsun. "Bu malzemeden şu kadar, ondan bu kadar" — her tarif için malzemeleri tartıp topluyorsun. Matris çarpımı tam olarak bu: binlerce küçük "tart-topla" işlemi (FMA, yani çarp-ve-ekle) aynı anda. Transformer'ın zamanının yarısından fazlasını burada geçirmesinin sebebi de bu; modeldeki en kalabalık mutfak burası.

Olayın bütün derdi de zaten malzeme taşımak. Bir aşçının asıl işi pişirmek ama mutfak kalabalıksa vaktinin çoğu kilere gidip malzeme getirmekle geçer. GPU'da da aynı: çarpma-toplama hesabı bedava sayılır, asıl pahalı olan veriyi bellekten çekmek. O yüzden malzemeyi tek tek değil, kasayla taşıyoruz — vec4 yüklemesi tam bu, dört sayıyı bir hamlede çek. Ve hesabı 64×64'lük "tezgah parçalarına" (tile) bölüp, bir tezgaha sığacak kadarını hızlı erişilen ortak rafa (shared memory) koyuyoruz.

İşin asıl şık kısmı şu: bir aşçı tezgahtaki malzemeyle pişirirken, yanındaki yardımcı bir sonraki partinin malzemesini çoktan kilerden getirip boş tezgaha diziyor. Pişirme bitince beklemek yok, sıradaki parti hazır. "Çift arabellek" (double-buffering) dediğimiz şey bu — iki tezgah, biri pişerken öteki hazırlanır, bellek beklemesi hesabın arkasına gizlenir.

Bir de bonus: FFN katmanında iki ayrı işi (SwiGLU karışımı + sonraki çarpım) tek mutfakta birleştiriyoruz. Normalde ara karışımı bir kaba boşaltıp tekrar kaptan alman gerekir; biz onu tezgaha çekerken anında karıştırıp doğrudan çarpıma sokuyoruz. Ara kap hiç kirlenmiyor, devasa bellek trafiği yok oluyor. Aşağıda bütün bunların offset'leriyle, bariyerleriyle nasıl kurulduğunu satır satır göreceksin.

Ne Yapar?

Klasik dense matrix multiply:

Y[M, N] = X[M, K] @ W[K, N]

Forward'da:

M = seq_len (token sayısı, 512)
K = d_model (768)
N = output dim — Attention Q/K/V/O için 768, FFN gate/up/down için 3072, lm_head için 16384.

matmul_residual ek olarak Y += R[M, N] epilogue'da residual ekler. Layer connection için fused.

matmul_residual_swiglu_a ise FFN katmanının en büyük iki işlemini birleştirir: element-wise silu(GATE) * UP birleşimi (SwiGLU) ve ardından gelen W_down matris çarpımı (hidden @ W_down + R). Bu sayede hidden ara matrisi belleğe hiç yazılmaz, hem 1 dispatch hem de yoğun bellek bant genişliği tasarrufu sağlanır.

Bu kernel grubu pipeline'ın en hot path'idir. Performance budget'ın %50+'ı burada harcanır.

Algoritma — 64×64 Output Tile, 4×4 Sub-tile per Thread

Tile Hiyerarşisi

Output Y[M, N]
  ↓ 64×64 tiles
  Each workgroup writes ONE 64×64 tile
    ↓ 16×16 thread workgroup
    Each thread writes 4×4 sub-tile (16 register accumulators)
      ↓ K-dim tiling: TK=16 elements per inner loop iteration
      Cooperative tile-load 64×16 (A) and 16×64 (B) into shared memory

Sayılar

Boyut	Değer	Açıklama
`TM`	64	Output tile rows
`TN`	64	Output tile cols
`TK`	16	K-dim block per inner loop
`TK_PAD`	17	tileA stride (bank conflict avoidance)
Workgroup	16×16=256 thread
Normal Mod (Tek-Arabellek)
Tile A	64 × 17 × 4B = 4.4 KB	f32
Tile B	16 × 64 × 4B = 4 KB	f32
Çift Arabellekli Mod (Double-Buffered matmul)
Tile A_db	2 × 64 × 17 × 4B = 8.5 KB	f32 (double-buffered)
Tile B_db	2 × 16 × 64 × 4B = 8 KB	f32 (double-buffered)
Total wg memory (db)	~16.5 KB	Apple GPU max 32 KB limitine çok rahat sığar
Per-thread accumulators	16 register f32	(4×4 sub-tile)

Önemli Optimizasyonlar

1. Vec4 Tile Loads (16-Byte Vectorized Access)

X ve W girdileri belleğe array<vec4<f32>> (veya mixed-precision için vec4<f16>) olarak bağlanır. Her bir thread, 4 ayrı skaler load yerine tek bir 16-byte'lık vektörize load işlemi ile veriyi çeker.

Kısıtlar: Host tarafının matris boyutlarında K % 4 == 0 ve N % 4 == 0 koşullarını sağlaması zorunludur.

2. Çift Arabelleğe Alma (Double-Buffering)

matmul ve matmul_w16 kernel'larında, tileA_db ve tileB_db adlı iki katı büyüklükte workgroup tamponu tanımlanmıştır.

Çalışma Prensibi: t döngüsünde mevcut tile üzerinde 4×4 FMA hesaplamaları GPU ALU'larında sürerken, bir sonraki t + 1 adımının verisi asenkron olarak bellekten boş workgroup tamponuna çekilir (prefetch).
Kazanç: Bellek gecikmesi (memory latency) hesaplama arkasına gizlenir ve döngü içi bariyer sıklığı yarıya iner.

3. SwiGLU Forward Fusion

matmul_residual_swiglu_a ve matmul_residual_swiglu_a_w16 kernel'larında, A tile'ı bellekten yüklenirken silu(gate) * up işlemi on-the-fly (yükleme anında) hesaplanarak tileA_db içine yazılır.

Kazanç: FFN katmanındaki ara aktivasyon matrisi ([seq_len, d_ff]) belleğe yazılıp tekrar okunmaktan tamamen kurtulur. Katman başına devasa bir bellek bant genişliği tasarrufu sağlanır.

Bind Group ABI

`matmul` (4 binding)

Binding	Tür	Detay
0	`storage, read`	`X: array<vec4<f32>>` — `[M × K/4]` row-major
1	`storage, read`	`W: array<vec4<f32>>` — `[K × N/4]` row-major
2	`storage, read_write`	`Y: array<f32>` — `[M × N]`
3	`uniform`	`dims: vec4<u32>` — `(M, N, K, _)`

`matmul_residual` (5 binding)

Aynısı + R: array<f32> [M × N] residual input.

`matmul_residual_swiglu_a` (6 binding)

Binding	Tür	Detay
0	`storage, read`	`GATE_mrs: array<vec4<f32>>` — FFN gate projection çıktısı
1	`storage, read`	`UP_mrs: array<vec4<f32>>` — FFN up projection çıktısı
2	`storage, read`	`W_mrs: array<vec4<f32>>` — `W_down` ağırlık matrisi
3	`storage, read`	`R_mrs: array<f32>` — layer input residual
4	`storage, read_write`	`Y_mrs: array<f32>` — final layer output
5	`uniform`	`dims_mrs: vec4<u32>` — `(M, N, K, _)`

Dispatch Şekli

workgroup_size: (16, 16, 1) → 256 threads
grid: (ceil(N/64), ceil(M/64), 1) workgroups

Satır Satır — Double-Buffered `matmul`

Double-buffered ve vec4 yüklemeli matmul kernel'ının kritik kısımları:

1) Giriş ve Yükleme Yapısı

wgsl

@compute @workgroup_size(16, 16, 1)
fn matmul(@builtin(workgroup_id) wgid: vec3<u32>,
          @builtin(local_invocation_id) lid: vec3<u32>) {
    let M = dims.x; let N = dims.y; let K = dims.z;
    let K4 = K / 4u;
    let N4 = N / 4u;
    let tx = lid.x; let ty = lid.y;
    let tid = ty * 16u + tx;

    let block_row = wgid.y * TM;
    let block_col = wgid.x * TN;

2) Vektörize A / B Tile İlk Yüklemesi (Prologue)

Döngüye girmeden önce, t = 0 için ilk tile verisi 16-byte vec4 okumalarıyla çekilir:

wgsl

    {
        let aI0 = tid * 4u;
        let aIm = aI0 / TK; let aIk = aI0 % TK;
        let axr = block_row + aIm; let axc = aIk;
        let row_in = axr < M;
        // Tek bir vec4 okuması ile 4 float birden çekilir
        let xv = X[axr * K4 + axc / 4u];
        tileA_db[aIm * TK_PAD + aIk + 0u] = select(0.0, xv.x, row_in && (axc + 0u) < K);
        tileA_db[aIm * TK_PAD + aIk + 1u] = select(0.0, xv.y, row_in && (axc + 1u) < K);
        tileA_db[aIm * TK_PAD + aIk + 2u] = select(0.0, xv.z, row_in && (axc + 2u) < K);
        tileA_db[aIm * TK_PAD + aIk + 3u] = select(0.0, xv.w, row_in && (axc + 3u) < K);

        let bI0 = tid * 4u;
        let bIk = bI0 / TN; let bIn = bI0 % TN;
        let bwr = bIk; let bwc = block_col + bIn;
        let bwr_in = bwr < K;
        let wv = W[bwr * N4 + bwc / 4u];
        tileB_db[bIk * TN + bIn + 0u] = select(0.0, wv.x, bwr_in && (bwc + 0u) < N);
        tileB_db[bIk * TN + bIn + 1u] = select(0.0, wv.y, bwr_in && (bwc + 1u) < N);
        tileB_db[bIk * TN + bIn + 2u] = select(0.0, wv.z, bwr_in && (bwc + 2u) < N);
        tileB_db[bIk * TN + bIn + 3u] = select(0.0, wv.w, bwr_in && (bwc + 3u) < N);
    }
    workgroupBarrier();

3) Prefetch ve Hesaplama Döngüsü

Döngü içinde t + 1 verisi bir sonraki arabelleğe (nxt_a_off, nxt_b_off offsetleri ile) yüklenirken, mevcut cur tamponlarındaki veri çarpılır:

wgsl

    for (var t: u32 = 0u; t < nTiles; t = t + 1u) {
        let parity = t & 1u;
        let cur_a_off = parity * TA_DB_HALF;
        let cur_b_off = parity * TB_DB_HALF;

        if (t + 1u < nTiles) {
            let nxt_a_off = (1u - parity) * TA_DB_HALF;
            let nxt_b_off = (1u - parity) * TB_DB_HALF;
            let nxt_kBase = (t + 1u) * TK;

            let aI0 = tid * 4u;
            let aIm = aI0 / TK; let aIk = aI0 % TK;
            let axr = block_row + aIm; let axc = nxt_kBase + aIk;
            let row_in = axr < M;
            let xv = X[axr * K4 + axc / 4u];
            tileA_db[nxt_a_off + aIm * TK_PAD + aIk + 0u] = select(0.0, xv.x, row_in && (axc + 0u) < K);
            tileA_db[nxt_a_off + aIm * TK_PAD + aIk + 1u] = select(0.0, xv.y, row_in && (axc + 1u) < K);
            tileA_db[nxt_a_off + aIm * TK_PAD + aIk + 2u] = select(0.0, xv.z, row_in && (axc + 2u) < K);
            tileA_db[nxt_a_off + aIm * TK_PAD + aIk + 3u] = select(0.0, xv.w, row_in && (axc + 3u) < K);

            let bI0 = tid * 4u;
            let bIk = bI0 / TN; let bIn = bI0 % TN;
            let bwr = nxt_kBase + bIk; let bwc = block_col + bIn;
            let bwr_in = bwr < K;
            let wv = W[bwr * N4 + bwc / 4u];
            tileB_db[nxt_b_off + bIk * TN + bIn + 0u] = select(0.0, wv.x, bwr_in && (bwc + 0u) < N);
            tileB_db[nxt_b_off + bIk * TN + bIn + 1u] = select(0.0, wv.y, bwr_in && (bwc + 1u) < N);
            tileB_db[nxt_b_off + bIk * TN + bIn + 2u] = select(0.0, wv.z, bwr_in && (bwc + 2u) < N);
            tileB_db[nxt_b_off + bIk * TN + bIn + 3u] = select(0.0, wv.w, bwr_in && (bwc + 3u) < N);
        }

        // Mevcut arabellek ile 16 FMA hesaplaması
        for (var k: u32 = 0u; k < TK; k = k + 1u) {
            let a0 = tileA_db[cur_a_off + (4u * ty + 0u) * TK_PAD + k];
            let a1 = tileA_db[cur_a_off + (4u * ty + 1u) * TK_PAD + k];
            let a2 = tileA_db[cur_a_off + (4u * ty + 2u) * TK_PAD + k];
            let a3 = tileA_db[cur_a_off + (4u * ty + 3u) * TK_PAD + k];
            let b0 = tileB_db[cur_b_off + k * TN + (4u * tx + 0u)];
            let b1 = tileB_db[cur_b_off + k * TN + (4u * tx + 1u)];
            let b2 = tileB_db[cur_b_off + k * TN + (4u * tx + 2u)];
            let b3 = tileB_db[cur_b_off + k * TN + (4u * tx + 3u)];
            acc00 = fma(a0, b0, acc00); acc01 = fma(a0, b1, acc01); acc02 = fma(a0, b2, acc02); acc03 = fma(a0, b3, acc03);
            acc10 = fma(a1, b0, acc10); acc11 = fma(a1, b1, acc11); acc12 = fma(a1, b2, acc12); acc13 = fma(a1, b3, acc13);
            acc20 = fma(a2, b0, acc20); acc21 = fma(a2, b1, acc21); acc22 = fma(a2, b2, acc22); acc23 = fma(a2, b3, acc23);
            acc30 = fma(a3, b0, acc30); acc31 = fma(a3, b1, acc31); acc32 = fma(a3, b2, acc32); acc33 = fma(a3, b3, acc33);
        }
        workgroupBarrier();
    }

WGSL-Spesifik Notlar

1. `var<workgroup>` ve Çift Arabellek Maliyeti

Bellek boyutu derleme anında sabit olmalıdır. 2 * 1088 (tileA_db) + 2 * 1024 (tileB_db) floats = 16.5 KB workgroup memory tutar. WebGPU donanım limiti 32 KB olduğu için occupancy kaybı yaşanmadan mükemmel hızlanma elde edilir.

2. `fma()` ve Vektör Optimizasyonları

Döngü içindeki FMAs (acc = fma(a, b, acc)) GPU'nun Fused Multiply-Add donanım hızlandırıcısını direkt tetikler. Vektörel binding'ler sayesinde bellek veriyolu üzerindeki yük %75 oranında hafifletilir.

Code Review

Bulgu 1: `select` ile Bounds Masking Maliyeti

Risk	Açıklama
🟢 yok	`select(0.0, X[...], okay)` ifadesinde sınır dışı okumalar yapılsa dahi WebGPU'nun robust buffer erişimi sayesinde donanım kilitlenmesi yaşanmaz ve OOB veriler maskelenerek ekarte edilir. M ve N, 64'ün katı olduğundan edge-branch pratik eğitimde sıfır maliyetle çalışır.

Hızlı Kontrol Listesi

Test Senaryosu	Durum
`K` ve `N` 4'ün katı mı? (vec4 gereksinimi)	✅ Host tarafında doğrulandı
Double buffering prefetch yarışmasız çalışıyor mu?	✅ `workgroupBarrier()` doğruluğu
SwiGLU forward fusion ara bellek gereksinimini sıfırladı mı?	✅ `hidden` tamponu elendi
Bank conflict engellendi mi?	✅ `TK_PAD = 17`

04_rope.md — Rotary Position Embedding. Q ve K vektörlerine pozisyon-bazlı rotasyon uygular.

#matmul ve matmul_residual — Tiled Matrix Multiply (Forward)

#Nedir bu ya?

#Ne Yapar?

#Algoritma — 64×64 Output Tile, 4×4 Sub-tile per Thread

#Tile Hiyerarşisi

#Sayılar

#Önemli Optimizasyonlar

#1. Vec4 Tile Loads (16-Byte Vectorized Access)

#2. Çift Arabelleğe Alma (Double-Buffering)

#3. SwiGLU Forward Fusion

#Bind Group ABI

#matmul (4 binding)

#matmul_residual (5 binding)

#matmul_residual_swiglu_a (6 binding)

#Dispatch Şekli

#Satır Satır — Double-Buffered matmul

#1) Giriş ve Yükleme Yapısı

#2) Vektörize A / B Tile İlk Yüklemesi (Prologue)

#3) Prefetch ve Hesaplama Döngüsü

#WGSL-Spesifik Notlar

#1. var<workgroup> ve Çift Arabellek Maliyeti

#2. fma() ve Vektör Optimizasyonları

#Code Review

#Bulgu 1: select ile Bounds Masking Maliyeti

#Hızlı Kontrol Listesi

#Sonraki

`matmul` ve `matmul_residual` — Tiled Matrix Multiply (Forward)

Nedir bu ya?

Ne Yapar?

Algoritma — 64×64 Output Tile, 4×4 Sub-tile per Thread

Tile Hiyerarşisi

Sayılar

Önemli Optimizasyonlar

1. Vec4 Tile Loads (16-Byte Vectorized Access)

2. Çift Arabelleğe Alma (Double-Buffering)

3. SwiGLU Forward Fusion

Bind Group ABI

`matmul` (4 binding)

`matmul_residual` (5 binding)

`matmul_residual_swiglu_a` (6 binding)

Dispatch Şekli

Satır Satır — Double-Buffered `matmul`

1) Giriş ve Yükleme Yapısı

2) Vektörize A / B Tile İlk Yüklemesi (Prologue)

3) Prefetch ve Hesaplama Döngüsü

WGSL-Spesifik Notlar

1. `var<workgroup>` ve Çift Arabellek Maliyeti

2. `fma()` ve Vektör Optimizasyonları

Code Review

Bulgu 1: `select` ile Bounds Masking Maliyeti

Hızlı Kontrol Listesi

Sonraki