#00_shared.wgsl — Shared Preamble (Tüm Kernel'ların Tabanı)

Dosya: 00_shared.wgsl Pipeline adımı: yok — Bu bir preamble dosyasıdır, kernel değildir. Engine her kernel modülünün başına injekte eder.

#Nedir bu ya?

Bir projende 30 ayrı dosya olduğunu düşün ve her birinin en üstünde aynı #define'lar, aynı yardımcı fonksiyonlar, aynı sabitler tekrar ediyor. Elle her dosyaya kopyala-yapıştır yapmak istemezsin — bir prelude.h yazıp tepeye #include edersin, olur biter. İşte bu dosya tam olarak o "prelude", sadece bir sorun var: WGSL'de #include diye bir şey yok.

O yüzden engine kaba ama işe yarayan bir numara çeviriyor: her kernel'ı derlemeden hemen önce, bu dosyanın içeriğini düz string olarak kernel kodunun başına yapıştırıyor. Yani sen wg_reduce_sum'ı çağırdığında onu import etmiş ya da link'lemiş olmuyorsun; o fonksiyon zaten senin dosyanın içine fiziksel olarak eklenmiş durumda. Dilin built-in'iymiş gibi davranıyor ama aslında her kernel'a ayrı ayrı kopyalanan ortak metin.

İçinde ne var? Birkaç sabit (WG=256, NEG_INF gibi), bir-iki minik yardımcı (flat_id 2D dispatch'i tek bir lineer index'e indirger, nan_guard patlamış sayıları sıfırlar), ve bütün kernel'ların paylaştığı asıl ağır iş: subgroup hızlandırmalı workgroup reduction'ları (wg_reduce_sum, wg_reduce_max). Bunlar "256 thread'in değerini tek bir toplama/maksimuma indir" işini GPU'nun SIMD lane'lerini kullanarak klasik tree-reduction'dan ~%50 daha hızlı yapıyor.

İlginç kısım şurada: bu "kopyala-yapıştır" yaklaşımı tek geçerli yol olduğu için bir sürü ince kural doğuruyor — dosyada kernel marker'ı olmamalı, enable direktifleri en tepeye gelmeli, subgroup fonksiyonları "uniform control flow" altında çağrılmalı. Aşağıda bunların hepsi tek tek açılıyor.

#Ne Yapar?

Tüm WGSL kernel'larının paylaştığı ortak preamble'ı tanımlar. Her kernel compile edilmeden önce engine 00_shared.wgsl'in içeriğini en başa ekler. Yani burada tanımlı her şey:

• Sabit (WG, NEG_INF, SUBGROUP_SIZE, NUM_SUBGROUPS)
• Helper fonksiyon (flat_id, is_finite, nan_guard, wg_reduce_sum, wg_reduce_max)
• Shared workgroup memory (sh_red)

…her kernel'ın içinde görünür.

Önemli kural: Bu dosyada kernel marker (// --- KERNEL: name ---) olmamalı. Eğer olursa engine bu dosyayı kernel split'e tabi tutar ve preamble özelliği bozulur.

#Engine Tarafı — Nasıl Injekte Edilir?

engine.js'da init() fonksiyonu:

let sharedPreamble = await loadShader(SHARED_PREAMBLE_PATH, import.meta.url);
if (/^\/\/ --- KERNEL:/m.test(sharedPreamble)) {
    throw new Error(`${SHARED_PREAMBLE_PATH} must not contain kernel markers`);
}
sharedPreamble = 'enable f16;\nenable subgroups;\n' + sharedPreamble;

Yani her kernel'ın final source'u şöyle olur:

enable f16;
enable subgroups;
// ─── 00_shared.wgsl içeriği ───
const WG: u32 = 256u;
fn wg_reduce_sum(tid: u32, val: f32) -> f32 { ... }
// ... (rest of preamble)
// ─── per-file preamble (e.g., constants from 03_linear.wgsl) ───
const TM: u32 = 64u;
var<workgroup> tileA: array<f32, 1088>;
// ─── kernel marker'dan sonraki kod ───
@compute @workgroup_size(...)
fn matmul(...) { ... }

Bu, "include" sistemi olmayan WGSL'de tek geçerli yöntem — string concatenation.

#İçerik

#A. Sabitler

const WG: u32 = 256u;
const MAX_WG_DIM: u32 = 65535u;
const NEG_INF: f32 = -3.4028234e38;
const F32_MAX: f32 = 3.4028234e38;

const SUBGROUP_SIZE: u32 = 32u;
const NUM_SUBGROUPS: u32 = 8u;   // WG / SUBGROUP_SIZE

Niye SUBGROUP_SIZE sabit? WebGPU'nun subgroupSize runtime built-in'i olsa da bizim kodumuz derleme zamanında 32 varsayıyor. Apple/NVIDIA için doğru. AMD RDNA Chrome'da da 32 raporluyor (CDNA 64 ama Chrome'da CDNA yok). Eğer bir gün 64-lane bir GPU'da çalıştırsan reduction count yanlış olur — bu bilinçli bir hardware kısıtı.

#B. flat_id — 2D dispatch fallback için lineerize

fn flat_id(gid: vec3<u32>, nwg: vec3<u32>) -> u32 {
    return gid.x + gid.y * nwg.x * WG;
}

Neden var? WebGPU dispatchWorkgroups(x, y, z) 1D'de max 65535 workgroup verir. 200M token'lık corpus'ta (200M / 256 = 781K WG) bu sınır aşılır. Çözüm: 2D grid kullanmak.

engine.dispatch1D(...) host-side'da bu kararı veriyor:

• WG count ≤ 65535 → (N, 1, 1) 1D
• Aksi → (65535, ceil(N/65535), 1) 2D

Kernel iki durumu fark etmeden çalışmalı. flat_id her thread'i tek bir global ID'ye redüklüyor:

flat_id = gid.x + gid.y * (nwg.x * WG)
        = thread_in_x_row + (row_index * threads_per_row)

gid.x ve nwg.x thread/workgroup'un WG-cinsinden koordinatı, ama lineer ID thread-cinsinden isteniyor — × WG çarpımı bu çevirimi yapıyor.

Önemli: flat_id çağrılan kernel local_invocation_index ile concat edilmiş thread ID üretmek zorunda; sadece workgroup ID değil. Detay her kernel kendi içinde halleder.

#C. NaN/Inf Guard

fn is_finite(x: f32) -> bool {
    return (x == x) && (abs(x) < F32_MAX);
}

fn nan_guard(x: f32) -> f32 {
    return select(0.0, x, is_finite(x));
}

WGSL'de isInf/isNaN built-in yok. Self-defined:

• (x == x) → NaN için false, diğer her şey için true. NaN'ın tek tanımlayıcı özelliği: kendisine eşit değil.
• abs(x) < F32_MAX → ±Inf exclude edilir, finite floats kabul.

nan_guard(x):

• x finite ise x döner
• NaN/Inf ise 0 döner

Kullanım yeri: Backward pass kenarlarında — örneğin attention_backward çıktısında, eğer Q×K^T sırasında bir patlama olursa downstream gradient'i sıfırlamak yerine pipeline'ı NaN'le kontamine etmemek için.

#D. Subgroup-accelerated workgroup reduction

#Shared scratch

var<workgroup> sh_red: array<f32, 256>;

256 entry workgroup-shared bellek. Sadece ilk NUM_SUBGROUPS = 8 slotu anlamlı veri taşır, gerisi yastık. Niye 256? Çünkü tüm kernel'lar bu preamble'ı paylaşıyor; bazıları sh_red[tid] indeksleyip 256'ya kadar yazabilir (örn wg_reduce_* haricindeki bir helper). Defansif boyut.

#wg_reduce_sum

fn wg_reduce_sum(tid: u32, val: f32) -> f32 {
    let sg_sum = subgroupAdd(val);
    let sg_id  = tid / SUBGROUP_SIZE;
    let lane   = tid % SUBGROUP_SIZE;

    if (lane == 0u) { sh_red[sg_id] = sg_sum; }
    workgroupBarrier();

    var v: f32 = 0.0;
    if (lane < NUM_SUBGROUPS) { v = sh_red[lane]; }
    let result = subgroupAdd(v);
    workgroupBarrier();
    return result;
}

Algoritma — 2-level subgroup reduction:

1. Phase 1: Her subgroup içinde subgroupAdd → 8 ayrı subgroup-toplamı (her subgroup'un her thread'inde aynı değer)
2. Lane 0 her subgroup'un toplamını shared memory'ye yazar (8 yazma)
3. Barrier — shared memory tutarlı olsun
4. Phase 2: Her subgroup ilk 8 değeri yükleyip kendi içinde subgroupAdd yapar
5. Trailing Barrier (Kuyruk Bariyeri): sh_red tüm wg_reduce_* çağrıları tarafından paylaşılan ortak bir geçici alandır. Son okumadan hemen sonra ikinci bir workgroupBarrier(); olmazsa, hızlı bir subgroup fonksiyondan erken dönüp bir sonraki reduction adımına başlayabilir ve yavaş kalan subgroup henüz shared bellekten v = sh_red[lane] okumasını bitiremeden sh_red[sg_id] hücresini ezebilir. Bu bariyer ardışık reduction'ların (örn. attention ve cross_entropy içindeki wg_reduce_max -> wg_reduce_sum geçişleri) birbiriyle yarışmasını (race condition) önler.
6. Sonuç: Her thread aynı toplamı görür.

Niye phase 2'de tüm subgroup'lar redundant iş yapıyor?

WGSL'in uniform control flow kuralı: subgroupAdd ve diğer subgroup operasyonları uniform CF altında çağrılmak zorunda. Yani:

if (sg_id == 0u) {        // ← non-uniform!
    let total = subgroupAdd(...);  // ← compile error
}

Bunu yapmak için sg_id == 0u koşulu, tid / SUBGROUP_SIZE'a bağlı, ki o local_invocation_index'e bağlı, ki o tüm thread'ler için farklı → non-uniform, fail.

Çözüm: tüm subgroup'lar phase 2'yi çalıştırsın, hepsi aynı 8 değeri yükleyip aynı subgroupAdd'i çağırsın. Redundant computation ama uniform CF, kabul edilir.

Maliyeti: 7× redundant subgroupAdd. Ama her biri çok ucuz (8 lane × 1 op), barrier'lara kıyasla faydalı.

Klasik alternatif (Apple Metal'da yapılan):

8 iteration tree reduction with 8 barriers + log2(256) = 8 ALU ops

Subgroup versiyonu: 2 subgroup ops + 2 barriers = ~%50 daha hızlı reduction.

#wg_reduce_max

Aynı yapı ve kuyruk bariyeri mantığı, subgroupAdd → subgroupMax, identity 0.0 → NEG_INF.

fn wg_reduce_max(tid: u32, val: f32) -> f32 {
    let sg_max = subgroupMax(val);
    let sg_id  = tid / SUBGROUP_SIZE;
    let lane   = tid % SUBGROUP_SIZE;

    if (lane == 0u) { sh_red[sg_id] = sg_max; }
    workgroupBarrier();

    var v: f32 = NEG_INF;
    if (lane < NUM_SUBGROUPS) { v = sh_red[lane]; }
    let result = subgroupMax(v);
    workgroupBarrier();
    return result;
}

NEG_INF identity önemli — phase 2'de lane >= NUM_SUBGROUPS olan thread'lerin v = NEG_INF olması gerekiyor; aksi halde v = 0.0 olsa max-toplamı yanlışlıkla 0'a çekebilir (eğer tüm sums negatifse).

#Engine'in 00_shared.wgsl'i Nereye Eklediği

engine.js'in splitKernels fonksiyonu:

function splitKernels(source, fileLabel, sharedPreamble = '') {
    const marker = /^\/\/ --- KERNEL: (\S+) ---$/gm;
    const matches = [...source.matchAll(marker)];
    if (matches.length === 0) {
        throw new Error(`No kernel markers in ${fileLabel}`);
    }
    const filePreamble = source.slice(0, matches[0].index);
    const preamble = sharedPreamble + '\n' + filePreamble;
    const kernels = {};
    for (let i = 0; i < matches.length; i++) {
        const name = matches[i][1];
        const start = matches[i].index + matches[i][0].length;
        const end = matches[i + 1]?.index ?? source.length;
        kernels[name] = preamble + source.slice(start, end);
    }
    return kernels;
}

Yani bir kernel için final source = enable f16;\nenable subgroups;\n + 00_shared.wgsl içeriği + \n + dosyanın kernel-marker'dan önceki kısmı (file preamble) + kernel kodu.

Sonuç: wg_reduce_sum, flat_id, nan_guard, vs. her kernel'da çağrılabilir sanki dilin built-in'i gibi.

#WGSL-Spesifik Notlar

#1. enable direktifleri dosya başında olmak zorunda

WGSL spec: enable directive'leri ilk non-whitespace token olmalı. Bu yüzden engine enable f16; ve enable subgroups;'u şared preamble'ın da ÖNÜNE injekte ediyor:

enable f16;            ← ilk satır olmak zorunda
enable subgroups;
// ─── 00_shared.wgsl içeriği ───
const WG: u32 = 256u;
...

Eğer 00_shared.wgsl içine enable f16; koysaydık ve engine bu dosyanın üstüne başka bir şey eklerdi, derleme hata verirdi.

#2. var<workgroup> sınırlamaları

sh_red: array<f32, 256> her kernel için ayrı bir kopya alır (workgroup başına). Apple GPU'da workgroup memory limit ~32 KB; sh_red = 1 KB küçük. Ama eğer bir kernel kendi içinde 30 KB workgroup memory zaten kullanıyorsa (örn matmul'un tileA + tileB), sh_red'in ek 1 KB'i gözle görülür occupancy düşürebilir. WGSL bu durumu compile-time'da uyarmaz; bilmek lazım.

#3. Subgroup-uniform CF — silent compile error

Bizim kodun wg_reduce_sum içindeki "redundant phase 2" gerekçesi tam buradan. Yeni başlayan biri "sg_id == 0u içinde subgroupAdd" yazsa, Tint compiler "uniformity analysis failed" hatası verir. Standart Metal/CUDA pattern bu — WGSL daha katı.

#4. select(false_val, true_val, cond) — ternary

WGSL'de ternary operator cond ? a : b yok. select(b, a, cond) kullanılıyor. Sıralama select(false_branch, true_branch, condition) — Metal'in select(false, true, cond)'una benziyor ama acaba C/C++'tan farklı.

#Code Review

#Bulgu 1: SUBGROUP_SIZE hardcoded — 64-lane'lerde bozar

Mitigasyon: Eğer 64-lane GPU eklenirse wg_reduce_*'da subgroupSize runtime built-in'i kullanmak gerekir. Kodun şu an çalıştığı target'larda (Apple, NVIDIA desktop, Intel) sorun yok.

#Bulgu 2: sh_red 256 yerine 8 olabilir

#Bulgu 3: F32_MAX literal — derleyici toleransı

#Hızlı Kontrol Listesi

#Sonraki

01_embedding.md — token ID → vektör dönüşümü, modelin ilk forward adımı.