#00_infrastructure.wgsl — Element-wise Utility Kernel'ları

Dosya: 00_infrastructure.wgsl Pipeline adımı: yok — Bunlar tutkal kernel'lar, forward/backward/optimizer arası kullanılır.

#Nedir bu ya?

Bir dizi düşün, içinde milyonlarca sayı var. Sen de en sıkıcı şeyi yapmak istiyorsun: hepsini sıfırla. Ya da her birini 2 ile çarp. Ya da "şu listeyi öbürünün üstüne ekle". JavaScript'te bunlar tek satır: arr.fill(0), arr.map(x => x * 2), a[i] += b[i]. Hiç düşünmeden yazdığın for döngüleri.

İşte bu dosya tam olarak onlar — ama GPU'da. Modelin ağır işleri (matris çarpımı, attention) başka dosyalarda; burası ise o işleri birbirine bağlayan küçük yardımcıların çekmecesi. fill_zero diziyi sıfırlar, scale her elemanı bir sayıyla çarpar, axpy bir diziyi (ölçekleyip) öbürünün üstüne ekler, copy kopyalar, clamp taşan değerleri sınır içine iter. Hepsi "her elemana aynı basit şeyi uygula" kalıbında — yani element-wise.

Tek kelimeyle: bunlar mutfaktaki kaşık-bıçak takımı. Yıldız değiller, ama onlar olmadan hiçbir yemek tabağa gelmez. Gradient buffer'ını her adımda sıfırlamak, Adam optimizer'da moment güncellemek, gradient clipping yapmak — hepsi bu ufak kernel'lardan birini çağırıyor. Az iş yapan ama her yerde geçen tutkal.

Peki tek satırlık bir döngü GPU'da neden ayrı bir dosyayı hak ediyor? Çünkü "her elemana dokun" işini binlerce thread'e bölmek, hangi thread'in hangi elemana baktığını çözmek ve dizinin sonundaki artık thread'leri elemek gerekiyor. Basit görünen şeyin GPU tarafındaki minik düzeneği aşağıda.

#Ne Yapar?

6 küçük element-wise kernel sağlar — modelin "tutkal" işlemleri:

Hepsi tek tip pattern: 1D dispatch, WG=256, flat_id ile global index, i < n bounds check, atomik fonksiyon yok.

#Matematiksel Tanımlar

fill_zero:     dst[i] = 0                    ∀ i ∈ [0, n)
fill_const:    dst[i] = c                    ∀ i ∈ [0, n)
scale:         dst[i] ← α · dst[i]           ∀ i ∈ [0, n)
axpy:          dst[i] ← dst[i] + α · src[i]  ∀ i ∈ [0, n)   (BLAS-1 axpy)
copy:          dst[i] = src[i]               ∀ i ∈ [0, n)
clamp_inplace: dst[i] ← max(-c, min(c, dst[i]))  ∀ i ∈ [0, n)

axpy ismi BLAS-1 standardından gelir: "a · x plus y". Numerical computing'in en eski rutini.

#Bind Group ABI

Tüm kernel'larda aynı: @group(0) üzerinde 2-4 binding.

#fill_zero (2 binding)

#fill_const, scale (3 binding)

#axpy (4 binding)

#copy (3 binding)

#clamp_inplace (3 binding)

#Dispatch Şekli

workgroup_size:  256
total threads:   ceil(n / 256) workgroups × 256
threadgroup mem: 0 (hiçbiri workgroup memory kullanmıyor)

Host (engine.dispatch1D veya benzeri):

const wgCount = Math.ceil(n / 256);
const dim = wgCount <= MAX_WG_DIM
    ? [wgCount, 1, 1]
    : [MAX_WG_DIM, Math.ceil(wgCount / MAX_WG_DIM), 1];
pass.dispatchWorkgroups(...dim);

flat_id(gid, nwg) 1D ve 2D dispatch'i fark etmeden lineerize ediyor — bkz. 00_shared.md.

#Satır Satır Açıklama (axpy — temsili)

axpy en zengin örnek; diğerleri şablonu basitleştirir.

#1) Bind group declarations

@group(0) @binding(0) var<storage, read_write> dst: array<f32>;
@group(0) @binding(1) var<storage, read>       src: array<f32>;
@group(0) @binding(2) var<uniform> n: u32;
@group(0) @binding(3) var<uniform> alpha: f32;

• dst ve src ayrı bind'ler — WGSL'de storage, read ve storage, read_write aynı buffer'ı alias yapamaz (compiler gözünden non-aliasing olmalı). Eğer dst === src istenirse bu kernel kullanılamaz, başka bir kernel gerek.
• n ve alpha uniform — küçük değerler, push-constant gibi davranır. Read-only access guarantee'siyle GPU constant cache'inde durur.

#2) Kernel imzası

@compute @workgroup_size(256, 1, 1)
fn axpy(@builtin(global_invocation_id) gid: vec3<u32>,
        @builtin(num_workgroups) nwg: vec3<u32>) {

• @workgroup_size(256, 1, 1) — 256 thread per WG, kernel attribute olarak derleme-zamanı sabit. Metal'in [[max_total_threads_per_threadgroup(N)]] hint'inden farklı: WGSL'de bu garanti.
• global_invocation_id (Metal'in thread_position_in_grid'i) — bu thread'in tüm dispatch grid'indeki global koordinatı.
• num_workgroups — dispatchWorkgroups(x, y, z) çağrısındaki x, y, z değerleri.

#3) Lineerize

let i = flat_id(gid, nwg);

Detay 00_shared.md bölümünde. Sonuç: thread'in flat 1D index'i.

#4) Bounds check

if (i >= n) { return; }

WG=256 ile n 256'nın katı olmayabilir (örn n=1027 → 4×256 = 1024 thread yeter, ama 5×256 = 1280 dispatch ediliyor; son 253 thread bound dışı). Bunu yakalayan defensive check.

#5) Element-wise işlem

dst[i] = dst[i] + alpha * src[i];

Tek satır. dst[i], src[i] global memory'den iki ayrı load; * ve + ALU; tek store.

Memory pattern: Adjacent thread'ler adjacent memory okuyor — coalesced access (her warp tek 256-byte cache line'a hit). Bandwidth-optimal.

#Variants — Ne Farkı Var?

#fill_zero vs fill_const

fill_zero özel kernel çünkü 0 ile fill çok yaygın (gradient sıfırlama her step) ve value: f32 uniform binding'i kaldırarak 1 binding daha az → bind group setup'ı daha hızlı. Anlamlı pratikte mi? Marjinal — ama hot path için doğal optimizasyon.

#scale vs axpy

axpy 2 buffer ister (dst + src). Sadece dst *= alpha istiyorsan scale kullan (1 buffer). Daha az read = daha az memory traffic.

#clamp_inplace

İçinde simetrik clamp (-max_abs ile +max_abs arası). Asimetrik clamp lazımsa (örn ReLU = clamp(x, 0, +Inf)) bu kernel uyumsuz; başka bir formül lazım.

#Atomik Yok — Önemli Detay

Dosyanın başındaki yorum:

// All buffers are f32. Atomics are used for multi-writer gradient
// accumulation (CAS-add via bitcast — WGSL has no native atomic<f32>).

Ama bu dosyada hiçbiri atomik kullanmıyor — her bir thread tek bir slot'a yazıyor (dst[i]), her slot tek thread tarafından dokunulduğu için race yok.

Atomik fonksiyonlar başka yerlerde kullanılır:

• attention_backward (split öncesi version) — dK/dV'a CAS-add
• embed_backward — aynı token ID birden fazla pozisyonda → atomik
• reduce_norm_sq — atomicAdd to scalar

Bu utility dosyası "single-writer" pattern'inde — atomic gerekmez, daha hızlı.

#WGSL-Spesifik Notlar

#1. Buffer aliasing yasak

@group(0) @binding(0) var<storage, read_write> dst: array<f32>;
@group(0) @binding(1) var<storage, read>       src: array<f32>;

dst ve src farklı buffer olmalı. Aynı GPU buffer'ı iki binding'e bağlamak WebGPU runtime'da hata verir (bind group validation).

Eğer "in-place axpy" istiyorsan (x += alpha * x) bu kernel kullanılamaz; ya yeni bir kernel (scale_in_place_axpy: x *= 1+alpha) ya da scale + alpha=1+α numarası.

#2. array<f32> size'sız

var<storage, read_write> dst: array<f32>;

WGSL'de storage array'in size'i runtime determined (host buffer size'ından). array<f32, N> (sabit boy) sadece workgroup memory'de kullanılır.

#3. Uniform buffer alignment

@group(0) @binding(2) var<uniform> n: u32;
@group(0) @binding(3) var<uniform> alpha: f32;

Her bir uniform ayrı binding — host tarafında 4 ayrı buffer demek. Bunun yerine struct olarak da yapılabilir:

struct Params { n: u32, alpha: f32 }
@group(0) @binding(2) var<uniform> params: Params;

Ama bizim kodun stil tercihi: tek-değer uniform'lar ayrı. Pipeline binding count daha yüksek olur ama kod daha okunaklı.

#4. i32 vs u32 index

WGSL array indeksleme u32 ister. Metal'de int sıkıntısı yok, hepsi int. WGSL'de i: u32 olarak deklare etmek zorunda yoksa dst[i32(i)] cast ekleme gereği çıkar.

#Code Review

#Bulgu 1: Atomik yok ama isim/yorum atomik ima ediyor

#Bulgu 2: scale ve clamp_inplace neden f32 sınırlı?

Mitigasyon: Eğer bir gün f16 grad gerekirse scale_w16 variant eklenir. Şimdilik gerekli değil.

#Bulgu 3: fill_zero vs fill_const(0)

#Hızlı Kontrol Listesi

#Sonraki

01_embedding.md — token ID'sinden gerçek vektör çıkarma. Modelin ilk forward adımı.