llm.istanbul·Etüt
TR EN
Workbench →

bpe — GPU BPE Tokenizer Pipeline (Eğitim + Çıkarım)

Dosya: bpe.wgsl Pipeline adımı: GPU üzerinde pre-tokenization, eğitim adımları (pair counting, reduction, scan, compaction) ve Trie-tabanlı çıkarım (inference). Özellik: Subgroup kooperatif tarama (subgroup scan), O(1) Root LUT önbellekleme ve atomik çakışmaları sıfırlayan lokal paylaşımlı hash tabloları.

Tek bir dosyada toplanmış iki büyük pipeline:

  • Eğitim Pipeline (Training - 9 Dispatch): clear_tablepair_countfind_max4find_max_final_detsetup_mergemerge_reducescan_pass1scan_pass2scan_pass3finalize_compact
  • Çıkarım Pipeline (Inference - 3 Dispatch): trie_tokenizer_chunkedtrie_prefix_sumtrie_tokenizer_compact

Nedir bu ya?

Diyelim bir metni küçültmek istiyorsun ama tek bir karakter bile kaybetmeden. Bir editörde sürekli aynı kelime grubunu yazıp duruyorsun: "ve böylece", "ve böylece"... Bir noktada diyorsun ki "ya ben buna bir kısaltma uydurayım, vb yazıp geçeyim". İşte BPE'nin (Byte Pair Encoding) yaptığı tam olarak bu: metindeki en sık yan yana gelen ikiliyi bulup ona tek bir yeni sembol veriyor. Sonra tekrar bakıyor, yeni en sık ikiliyi buluyor, ona da bir sembol veriyor. Bu döngü binlerce kez dönüyor ve elinde bir "sözlük" (vocabulary) oluşuyor — sık geçen şeyler tek parça, nadir şeyler hâlâ harf harf.

Eğitim tarafı işte bu döngünün GPU üzerinde dönen hâli. Ama bir sorun var: "en sık ikiliyi bul" demek, milyonlarca harf çiftini saymak demek. Tek bir thread tek tek sayarsa ömür yetmez. Onun yerine milyonlarca thread aynı anda sayar — ama hepsi aynı sayaç kutusuna (global hash tablosuna) yazmaya çalışınca birbirine giriyorlar (contention). Çözüm, herkesin önce kendi masasında küçük bir hesap tutması (workgroup-lokal hash tablosu), sonra sadece toplamları tek seferde ortak deftere geçirmesi.

Çıkarım (inference) tarafıysa tam tersi: artık sözlüğün hazır, eline yeni bir metin geliyor ve onu en az sayıda parçaya bölmek istiyorsun. Burada hile şu — sözlüğü bir Trie (önek ağacı) olarak diziyorsun, tıpkı bir telefon rehberi ya da otomatik tamamlamanın çalıştığı yapı gibi. Her karakterde ağaçta bir adım iniyorsun ve "en uzun eşleşen kelimeyi" arıyorsun (greedy longest-match): "k", "ke", "ked", "kedi"... en uzun nereye kadar gidebiliyorsan o token.

Geri kalan her şey bu iki fikrin etrafındaki donanım dansı: bariyerleri 17'den 2'ye indiren subgroup taramaları, ağacın kökünü shared memory'de önbellekleyip global bellek beklemesini sıfırlayan numaralar, thread'leri aynı hizada tutup dallanmayı (divergence) engelleyen branchless aramalar. Yani "ne sayılıyor / ne aranıyor" basit; asıl iş bunu GPU'da boğulmadan yapmak.


Ne Yapar?

1. GPU Pre-tokenization (bpe_word_boundary)

BPE eğitiminde kelime sınırları boyunca birleştirmelerin (örneğin "yakınlık▁ve" → tek bir token) oluşmasını engellemek için, Türkçe-uyumlu GPT-4 tarzı pre-tokenization yapılır. Her sembolün karakter sınıfı (Letter, Digit, Space, Punct, Newline) belirlenir. Sınıf geçişlerinde o sembolün 16. biti WORD_START_BIT (0x10000) ile etiketlenir. BPE çift toplama kernel'ı bu biti görünce o sınır üzerinden çift saymayı durdurur.

2. İki Seviyeli Çift Sayımı (bpe_pair_count_b)

Milyonlarca thread doğrudan global hash tablosuna atomik ekleme (atomicAdd) yapmaya çalışırsa korkunç bir çakışma (contention) ve performans kaybı oluşur. Bunu önlemek için iki seviyeli sayım yapılır:

  1. Thread'ler kendi workgroup'ları içinde 1024 slotlu lokal paylaşımlı bir hash tablosuna (local_ids ve local_counts) Murmur3 hashing + quadratic probing ile lokal olarak veri biriktirir.
  2. Lokal reduction bittikten sonra, tek bir hamlede global hash tablosuna sadece workgroup toplamları aktarılır. Bu sayede global atomik çarpışmalar en aza indirilir.

3. Subgroup Kooperatif Kompaktlama (bpe_finalize_compact_b)

Blelloch exclusive scan algoritmalarında geleneksel olarak 8 aşamalı up-sweep, 8 aşamalı down-sweep ve 17 adet workgroupBarrier kullanılır. bpe_finalize_compact_b bu ağır bükümü subgroup-cooperative scan ile çözer:

  • Subgroup içi taramalar (subgroupExclusiveAdd ve subgroupAdd) donanımsal düzeyde (bariyersiz) çalışır.
  • Sadece subgroup toplamları shared belleğe (sh_sg_excl) yazılarak workgroup düzeyinde taranır.
  • 17 yerine sadece 2 bariyer ve 4 subgroup yönergesiyle 256 elemanlık exclusive scan tamamlanır.
  • Branchless select() ve WORD_START_BIT taşımalarıyla semboller tek dispatch'te hedefe saçılır (scatter).

4. Trie-Tabanlı Hızlı Çıkarım (trie_tokenizer_chunked)

GPU üzerinde metin çıkarımı yaparken (trie binary search) bellek gecikmesini önlemek için şu üstün optimizasyonlar uygulanmıştır:

  • Root LUT önbellekleme: Trie'nin kök düğümü (Node 0) her karakter eşleşme başlangıcında tekrar tekrar sorgulanır. Byte-level trie için kök düğümün 256 çocuğu shared memory'de root_lut olarak cache'lenir. Böylece ilk seviyedeki binary search tamamen atlanarak O(1) doğrudan erişim sağlanır.
  • Derinlik-1 Cache: Kökün hemen altındaki (derinlik 1) düğümün firstChild, numChildren ve tokenId değerleri de shared memory'de cache'lenerek token başına 3 global bellek okuması sıfırlanır.
  • Branchless Arama: Non-root düğüm aramalarında select kullanılarak warp divergence (SIMD mask split) engellenir. Tüm thread'ler tek hizada (uniform execution) çalışır.

Bind Group ABI

bpe_pair_count_b (4 binding)

Lokal paylaşımlı hash tablosunu kullanarak batched çift toplar.

BindingTürDetay
0storage, readsymbols: array<u32> — Giriş sembol dizisi
1storage, read_writepair_counts: array<atomic<u32>> — Global çift sayıları
2storage, read_writepair_ids: array<atomic<u32>> — Global çift kimlikleri (packed pair)
3storage, readstate: IterState — symbol_count, table_size vb. iteration state

bpe_merge_reduce_b (4 binding — Fused Merge Mark & Reduce)

En iyi çifti maskeler ve aynı workgroup içinde valid sembol toplamlarını reduction ile toplar.

BindingTürDetay
0storage, read_writesymbols: array<u32> — A-side merge işaretlerini yazar
1storage, read_writevalid_mask: array<u32> — `valid
2storage, read_writeblock_sums: array<u32> — Her bloğun geçerli sembol toplamı
3storage, readstate: IterState — symbol_a, symbol_b, new_symbol

bpe_finalize_compact_b (5 binding — Fused Scan + Merge Apply + Scatter)

Subgroup exclusive scan ile offsetleri çözer ve sembolleri sıkıştırarak hedefe saçar.

BindingTürDetay
0storage, readvalid_mask: array<u32> — bit 0 = valid, bit 1 = merge intent
1storage, readblock_sums: array<u32> — exclusive scanned block offsets
2storage, readinput_symbols: array<u32> — Kaynak sembol buffer'ı
3storage, read_writeoutput_symbols: array<u32> — Hedef sıkıştırılmış sembol buffer'ı
4storage, readstate: IterState

trie_tokenizer_chunked (6 binding — Trie-tabanlı Çıkarım)

Karakter dizisini pre-compiled binary trie ile greedy longest-match olarak tokenize eder.

BindingTürDetay
0storage, readinput: array<u32> — packed UTF-8 baytları (u32 başına 4 bayt)
1storage, readnodes: array<u32> — Trie düğüm bilgileri (firstChild, numChildren, tokenId)
2storage, readedges: array<u32> — Trie kenar bilgileri (byte, childNodeIndex)
3storage, read_writetoken_output: array<u32> — Çıktı token ID'leri
4storage, read_writechunk_counts: array<u32> — Her chunk'ın ürettiği token sayısı
5uniformparams: TrieParams — input_length, chunk_size, max_tokens_per_chunk

Satır Satır Analiz

1) bpe_word_boundary — Byte Sınıflandırma

wgsl
fn char_class(tok: u32) -> u32 {
    if (tok == 0x0Au) { return 4u; }
    if (tok == 0x20u) { return 2u; }
    if (tok - 0x30u <= 9u) { return 1u; }
    if (tok >= 0x80u) { return 0u; }
    if (tok - 0x61u <= 25u) { return 0u; }
    if (tok - 0x41u <= 25u) { return 0u; }
    return 3u;
}
  • Branch-less dostu çıkarma hilesi: (tok - base) <= range karşılaştırması tek işlemde tüm ASCII aralığını kapsar. SIMD mask dallanmasını (divergence) azaltarak donanımın tüm thread'lerini hizada tutar.
  • 0x80u üzerindeki tüm UTF-8 kontinyasyon ve öncü baytları doğrudan sınıf 0 (Letter) olarak işaretlenir. Bu sayede Türkçe karakterler (ı, ş, ğ, ç, ö, ü vb.) otomatik olarak harf sınıfına alınır.

2) bpe_finalize_compact_b — Subgroup Kooperatif Scan

wgsl
let local_excl = subgroupExclusiveAdd(v);
let sg_total   = subgroupAdd(v);

if (lane == 0u) { sh_sg_excl[sg_id] = sg_total; }
workgroupBarrier();

var t: u32 = 0u;
if (lane < NUM_SUBGROUPS) { t = sh_sg_excl[lane]; }
let cross_excl = subgroupExclusiveAdd(t);

if (sg_id == 0u && lane < NUM_SUBGROUPS) { sh_sg_excl[lane] = cross_excl; }
workgroupBarrier();

let final_excl = sh_sg_excl[sg_id] + local_excl;
  • Geleneksel 256 elemanlık Blelloch exclusive scan yerine, donanım düzeyinde subgroup intrinsics kullanılır:
    1. subgroupExclusiveAdd(v) ile subgroup içi exclusive toplam hızlıca (bariyersiz) çözülür.
    2. Her subgroup'ın toplamı (sg_total) shared belleğe (sh_sg_excl) atılır.
    3. Tüm thread'ler shared bellekteki bu 8 subgroup toplamı üzerinden bir kez daha exclusive subgroup scan çalıştırır (cross_excl).
    4. Geriye sadece kendi subgroup taban offset'ini lokal offset'e eklemek kalır.
  • Maliyet: 17 bariyer yerine sadece 2 bariyer ve donanımsal subgroup ALU birimleri. Hız artışı muazzamdır.

3) trie_tokenizer_chunked — Root LUT ve Derinlik-1 Cache

wgsl
let nc = min(cached_root_nc, MAX_CACHED_EDGES);
let fc = cached_root_fc;
if (lid.x < nc) {
    let sym     = edges[(fc + lid.x) * 2u] & 0xFFu;
    let d1_node = edges[(fc + lid.x) * 2u + 1u];
    root_lut[sym] = d1_node;
    d1_fc[sym]    = nodes[d1_node * 3u];
    d1_nc[sym]    = nodes[d1_node * 3u + 1u] & 0xFFFFu;
    d1_tid[sym]   = nodes[d1_node * 3u + 2u];
}
workgroupBarrier();
  • Her token eşleşmesi Trie kök düğümünden (Node 0) başlar.
  • Kökün tüm dalları (root_lut) ve derinlik-1 düğümlerinin meta verileri (d1_fc, d1_nc, d1_tid) shared memory'de cache'lenir.
  • Arama döngüsünde cn == 0u ise (derinlik 0'dan 1'e geçiş) arama O(1) hıza düşer:
    wgsl
    nn = root_lut[bv];
  • Derinlik 1'den 2'ye geçişte ise global bellek pointer-chasing yerine shared cache okunur:
    wgsl
    nn = find_child_global(d1_fc[rb], d1_nc[rb], bv);
    Bu iki adım, en çok tetiklenen arama seviyelerinde global bellek darboğazını tamamen yok eder.

Code Review

Bulgu 1: bpe_pair_count_b Local Table Sizing ve Occupancy İlişkisi

RiskAçıklama
🟡 perfLOCAL_TABLE_SIZE shared memory üzerinde 1024 slot (8 KB) olarak belirlenmiştir. Bu boyutun 2048 slot (16 KB) yapılması denendiğinde Apple M GPU'larında occupancy oranının yarı yarıya düştüğü (4 WG/SM → 2 WG/SM) ve bellek gizleme (latency hiding) kaybı nedeniyle adım süresinin %70 yavaşladığı saptanmıştır. 1024 slotluk mevcut denge (%25 doluluk + Murmur3 + quadratic probing) en optimum donanımsal occupancy profilini vermektedir.

Bulgu 2: trie_tokenizer_compact Coalesced Bellek Yazımı

RiskAçıklama
🟢 ÇözüldüEski tasarımdaki 1 thread = 1 chunk yaklaşımı, ardışık thread'lerin bellek üzerinde stride-512 sıçramalı (coalescing bozuk) yazmasına neden oluyordu. trie_tokenizer_compact kernel'ında ise 1 workgroup = 1 chunk eşleşmesi yapılmış; 256 thread ardışık yazma indislerini (compact_output[db + i]) kooperatif kopyalayarak coalesced memory access başarısını maksimuma çıkarmıştır.

Bulgu 3: PCIe Darboğazının Giderilmesi (trie_prefix_sum)

RiskAçıklama
🟢 ÇözüldüTokenizer çıktılarının sıkıştırılması için gerekli olan prefix sum işlemi eski yapıda CPU'ya mapAsync ile çekilip CPU'da taranıp geri yükleniyordu. trie_prefix_sum tek thread'lik (workgroup_size(1)) bir kernel olarak GPU'ya taşınmış, CPU roundtrip'i sıfırlanarak 2 PCIe transfer gecikmesi ve 1 GPU fence beklemesi tamamen ekarte edilmiştir.

Hızlı Kontrol Listesi

Test SenaryosuDurum
Kelime sınırları doğru taranıyor mu?symbols[id] | WORD_START_BIT
subgroup operasyonları donanımca desteklenmiyorsa?⚠️ Safari/Webkit üzerinde training fallback mekanizması gerektirir
Trie aramasında warp divergence engellendi mi?select() tabanlı branchless lower_bound
Local hash table modulo power-of-2 mi?LOCAL_TABLE_MASK = 1023
1B sembole kadar Blelloch scan tırmanabiliyor mu?✅ Hierarchical 3-pass scan
Kök düğüm LUT cache güvenli mi?✅ 256 byte ASCII LUT

Host Mimarisi ve Koordinasyon (trie.js & tokenizer.js)

BPE tokenizer'ın üstün hızı yalnızca GPU shader'larından değil, host (JS) tarafındaki akıllı bellek yönetimi ve asenkron yürütme mimarisinden de kaynaklanır.

1. BFS Tabanlı Önbellek-Dostu Trie Derleme (compileVocabToTrie)

Bellekteki Trie ağacı ham kelime listesinden (byte dizileri) çıkarıldıktan sonra, flat binary (v3 formatı) olarak serileştirilir.

  • BFS Sıralaması (Breadth-First Search): Ağaç düzleştirilirken BFS sırası izlenir. Bu sayede bir düğümün çocukları bellek üzerinde ardışık (consecutive) yerleşir. GPU'nun L1/L2 önbellekleri bu ardışık indislere erişirken maksimum hit oranına ulaşır.
  • Sıralı Kenarlar (Sorted Children): Çocuk kenarlar sembol bayt değerine göre küçükten büyüğe sıralanır. Bu sıralama, GPU tarafındaki non-root düğüm aramalarında find_child_global helper fonksiyonunun uniform binary search yapabilmesini sağlar.

2. Kalıcı Bellek Havuzu (Persistent Buffer Pooling)

WebGPU'da her tokenizasyon çağrısında sıfırdan GPUBuffer oluşturup yok etmek (garbage collection) ve GPU allocation overhead oluşturmak büyük bir gecikme sebebidir. TrieTokenizer bu sorunu amorti edilmiş O(1) bellek havuzu ile çözer:

  • Giriş metninin boyutuna göre kalıcı tamponlar (#inputBuf, #tokenBuf, #countsBuf, #offsetsBuf, #totalBuf, #compactBuf) bir kez oluşturulur ve sonraki tüm encode çağrılarında yeniden kullanılır (re-use).
  • Sadece yeni gelen metin mevcut havuz kapasitesinden büyük olduğunda havuz 1.5 kat büyütülerek yeniden tahsis edilir.

3. Dinamik Dilimleme (Multi-pass Slicing)

WebGPU donanımsal maxBufferSize limitine sahiptir. Çok büyük bir metin tek seferde GPU'ya yüklenmeye çalışılırsa bellek taşması oluşur.

  • TrieTokenizer gelen metin boyutunu GPU limitlerine göre analiz eder.
  • Eğer girdi limitleri aşıyorsa, metni güvenli dilimlere (sliceSize = Math.floor(maxInputPerPass / chunkSize) * chunkSize) bölerek ardışık dispatche'ler ile güvenli ve kararlı bir şekilde işler.

4. Tekil Submit ve Sıfır-CPU Gecikmesi (Single Submit Command Encoding)

En kritik optimizasyonlardan biri, CPU ile GPU arasındaki asenkron komut iletişimindedir:

  • 3 ayrı pass (trie_tokenizer_chunkedtrie_prefix_sumtrie_tokenizer_compact) tek bir CommandEncoder içinde birleştirilir ve tek bir submit (device.queue.submit) ile GPU'ya gönderilir.
  • Bu sayede ara adımlarda CPU asla GPU'yu beklemez (zero CPU-GPU roundtrip). GPU komutları kesintisiz olarak boru hattından (pipeline) geçirir.
  • Akıllı DMA Kopyası: Sıkıştırma (compaction) öncesinde çıktının tam boyutu bilinmez (çünkü her chunk'ın ürettiği token sayısı değişkendir). CPU-GPU senkronizasyonunu kilitlememek için önce 4 byte'lık total_tokens sayısı GPU üzerinde hesaplanır ve mapAsync ile host'a near-instant olarak çekilir (çünkü iş çoktan bitmiştir). Ardından tam boyutta bir DMA copyBufferToBuffer tetiklenerek sadece sıkıştırılmış gerçek token dizisi okunur. Bellek israfı sıfırlanır.

Sonraki

Tekrar dizine dönmek için: index.md.

WGSL kernel etüdleri · WebGPU üzerinde sıfırdan LLMİstanbul’da Uğur Toprakdeviren tarafından hazırlandı.