Bellek Duvarının Fiziği: Standart Attention'dan FlashAttention'a O(N²) Çıkmazı, Online Softmax ve SRAM Mimarisi

Modern yapay zekânın en büyük darboğazı modelin zekâsı değil, verinin fiziksel hareketidir. Bir GPT mimarisinin katmanları boyunca akan hesaplamanın kalbindeki self-attention operasyonu, kâğıt üzerinde zarif bir matris çarpımı ve yumuşatma fonksiyonundan ibarettir; ancak pratikte, işlemcinin üretebildiği ham hesap gücünü tek başına tüketen bir bellek erişim yüküne dönüşür. Bu makale o yükün anatomisini çıkarıyor: O(N²) karmaşıklığının cebirsel kökeninden, bellek duvarının 1995’te adını konmuş fiziksel sebebine; softmax’ın sayısal olarak neden blok blok parçalanamaz göründüğünden, onu akışlı (streaming) hale getiren online softmax türetmesine; ve FlashAttention’ın GPU’nun on-chip SRAM’ini algoritmanın birincil tasarımıyla barıştırarak bu duvarı nasıl aştığına kadar — tüm formüller, sabitler, IO alt sınırları ve nedensel maskeleme, dropout ile decoding aşamasının ayrılmaz parçası Flash-Decoding dahil.

Bu, WebTuring Principia’nın ilk yazısıdır. Principia, haber değil kalıcı mühendislik birinci ilkeleri yayınlar: burada yazılanlar gündemle eskimez, çünkü fizik ve asimptotik karmaşıklık gündemi takip etmez. Yine de dürüstlük gereği: bu yazının donanım sayılarıyla kısmı (A100, H100, B200) zamanla eskiyecek olan katmandır; asimptotik kısım eskimez. İkisini birbirinden ayırarak okuyun.

---

Attention’ın Cebiri: Formüller ve Karmaşıklık Sınıfları

Bir transformerdaki N adet token, her katmanda d boyutlu (head başına d_k) vektörlerle temsil edilir. Girdi temsilleri X ∈ ℝ^N×d_model olmak üzere, sorgu (query), anahtar (key) ve değer (value) matrisleri öğrenilmiş projeksiyonlarla üretilir:

Ölçekli nokta-işaret attention tanımı (Vaswani ve ark., 2017):

Bu tek satır, üç ayrı maliyet nesnesi içerir. Skor matrisi S = QK^T/√d_k ∈ ℝ^N×N, satır bazında softmax P = softmax(S) ∈ ℝ^N×N ve çıktı O = PV ∈ ℝ^N×d_k.

Zaman karmaşıklığı: QK^T çarpımı head başına 2N²d_k kayan-nokta işlemi (FLOP) gerektirir; PV çarpımı bir o kadar daha. Toplam:

Softmax’ın kendisi (üs, toplam, bölme) satır başına O(N) olmak üzere toplamda O(N²) FLOP’tur. İki ifade aynı N² mertebesinde ama katsayıları farklı: matris çarpımları 4N²d_k, softmax ise ~3N² mertebesinde FLOP üretir — yani fark yaklaşık d_k katıdır, iki kat değil. d_k = 128 için softmax’ın toplam hesabı içindeki payı %1’in altındadır; bu, sonraki bölümde kritik olacak bir ayrımdır: softmax FLOP bakımından önemsiz, veri hareketi bakımından belirleyicidir.

Bellek karmaşıklığı: Asıl sorun burada. S ve P matrislerinin her biri N² eleman tutar. Bu, attention’ın depolama maliyetinin token sayısının karesiyle büyüdüğü anlamına gelir:

Neden √d_k ile ölçekleme? Varsayalım ki q ve k bileşenleri bağımsız, ortalaması 0, varyansı 1 rastgele değişkenler. O hâlde q·k = Σ_i=1^d_k q_ik_i toplamının varyansı d_k’dır (her terimin varyansı 1, kovaryanslar sıfır). Yani nokta-işaretin standart sapması √d_k ile büyür; bu değer softmax’a girince üstel fonksiyon doygunluğa sürüklenir ve gradyanlar pratik olarak ölür. 1/√d_k ölçeklemesi varyansı tam olarak 1’e geri getirir:

O(N²) Çıkmazı: Kare Matrisin Fiziksel Bedeli

Standart attention implementasyonları (naif PyTorch dahil) S ve P matrislerini tam olarak malzemeleştirir: N×N tam bir matris GPU’nun ana belleğine (HBM/DRAM) yazılır, okunur, üzerine yazılır. Bu, teorik karmaşıklık analizinde “sabit” görünen bir faktörün pratikte nasıl büyüdüğünün dersidir.

Sayısal bir ölçek tahmini yapalım. FP16’da bir eleman 2 bayt tutar. Tek bir attention head için P matrisinin bellek ihtiyacı:

Bağlam uzunluğu N	N² (eleman)	Tek head, FP16 bellek
4.096 (klasik GPT-2)	1,68 × 10⁷	33,6 MB
16.384	2,68 × 10⁸	537 MB
65.536	4,29 × 10⁹	8,6 GB
131.072 (128k bağlam)	1,72 × 10¹⁰	34,4 GB

128k token’lık bir bağlamda tek bir head’in olasılık matrisi, 80 GB’lık bir A100’un belleğinin yaklaşık %43’ünü tek başına kullanır (34,4 / 80 ≈ 0,43) — yani yarısından az, ama bir head için tek başına yetersiz bir pay. Bir modelde onlarca head ve onlarca katman olduğunu, bir de eğitimin geri yayılım için bu tensörleri saklaması gerektiğini düşünürsek, kare matris malzemeleştirmek uzun bağlamda pratik olarak olanaksız hale gelir. O(N²) burada asimptotik bir not değil, doğrudan bir bağlam uzunluğu tavanıdır.

Daha sinsi olan ikinci maliyet ise veri hareketidir (IO). Standart bir forward geçişinde, head başına yaklaşık şu HBM trafiği oluşur:

1. Q ve K’nın okunması: 2 × Nd_k eleman
2. S’in yazılması: N² eleman
3. Softmax için S’in okunması + P’nin yazılması: 2N² eleman
4. PV için P’nin okunması: N² eleman
5. O’nun yazılması: Nd_k eleman

Toplam: Θ(N² + Nd_k) ≈ Θ(N²) HBM erişimi — yani hesaplamanın kendisi kadar, hatta bant genişliği kısıtında ondan daha fazla, veri taşıma zorunluluğu.

Bellek Duvarının Fiziği: Ölüm Sebebi Hesap Değil, Taşıma

Bellek duvarı (memory wall) terimi, işlemci hızının bellek hızından sistematik olarak çok daha hızlı büyümesi ve aradaki makasın her nesilde açılması olgusunu adlandırır (Wulf & McKee, 1995). Rakamlarla: 1980’lerden 2010’lara uzanan 30+ yıllık pencerede tepe kayan-nokta hesap hızı ortalama yılda ~%52-60 büyürken, DRAM bant genişliği yalnızca yılda ~%10, DRAM gecikmesi ise neredeyse hiç iyileşmeden ilerledi (Patterson, IAS 2013). Yani bugün bir GPU’nun transistörleri, veriyi kendisine taşıyabileceğiniz hızdan kat kat fazla sayıda çarpma yapabilecek durumdadır.

Bu dengesizliği ölçen standart araç roofline modelidir (Williams, Waterman & Patterson, 2009). Bir çekirdeğin ulaşabileceği performans, iki duvarın altındaki en iyidir:

Burada π_tepe tepe FLOP/s, β bant genişliği (byte/s) ve I = FLOP/byte ise iş yükünün aritmetik yoğunluğu (arithmetic intensity). İki duvarın kesiştiği sırt noktası (ridge point) kritik bir eşiktir; altında bellek-bağlı (memory-bound), üstünde hesap-bağlı (compute-bound) çalışırsınız.

NVIDIA A100 (SXM4, 40 GB) ile somutlaştıralım:

• Tepe FP16 Tensor çekirdeği hesabı: π_tepe ≈ 312 TFLOP/s
• HBM2e bant genişliği: β ≈ 1.555 TB/s (40 GB model; 80 GB SXM modelde ~2.0 TB/s)
• Sırt noktası: 312 × 10¹² / 1.555 × 10¹² ≈ 200 FLOP/byte

Şimdi standart attention’ın forward IO bütçesini aritmetik yoğunluğa çevirelim. Matris çarpımları 4N²d_k FLOP üretir; Θ(N²) mertebesindeki ana matris trafiği (S yazımı + S/P okuma-yazma geçişleri ≈ 4N² eleman) FP16’da ~8N² bayt taşır. Kaba ama dürüst bir oran:

Tipik d_k = 128 için I ≈ 64 FLOP/byte — A100’ün sırt noktası olan ~200’ün belirgin şekilde altında. Softmax’ın eleman-bazlı geçişleri ise daha da kötüdür: satır maksimumu, üstel ve toplam gibi işlemler bayt başına yalnızca birkaç FLOP yapar (I ≈ 1 mertebesinde) ve saf bellek-bağlıdır. Sonuç: attention, GPU’nun hesaplama gücünü asla doyuramaz; çekirdekler veri bekler.

Zaman tahmini de aynı şeyi söyler. N = 16.384, d_k = 128, tek head, FP16 için:

• Matris çarpımlarının tepe-hesap süresi: 4N²d_k / 312×10¹² ≈ 0,44 ms
• Yaklaşık Θ(N²) HBM trafiğinin süresi: 8N² bayt / 1,555×10¹² ≈ 1,38 ms

Yani saf IO, saf hesabın üç katı — ve bu, softmax’ın eleman-bazlı maliyetlerini ve geri yayılımı hiç saymadan. Standart attention bir hesaplama problemi değil, bir veri taşıma problemidir. FlashAttention’ın hareket noktası tam olarak bu cümledir.

Softmax’ın Kalesi: Neden Blok Blok Parçalanamaz Görünüyor?

Bir matris çarpımını bloklara ayırmak rutindir; tiling, GPU mimarisinin doğasında vardır. Softmax ise satır-özgün (row-coupled) bir normalizasyon operatörüdür ve bu, tiling’i ilk bakışta imkânsız kılar. Tanım gereği softmax:

Paydadaki toplam, satırın tüm elemanlarını gerektirir. Skor matrisinin bir bloğunu SRAM’de hesaplayıp HBM’e yazmadan geçmek isterseniz, o bloğun nihai ağırlığını bilemezsiniz — çünkü normalleştirme sabiti henüz tamamlanmamıştır.

İkinci sorun sayısal kararsızlıktır. FP32’de en büyük temsil edilebilir sayı ~3,4×10³⁸ olduğundan e^x yaklaşık x > 88,7’de taştığında sonsuza (inf) döner; FP16’da bu eşik e¹¹ civarındadır ve çok daha erkendir. Standart stabilite hilesi, her satır için maksimum değerini çekmektir:

Bu dönüşüm matematiksel olarak değişmezdir (pay ve paydaya aynı e^−m çarpanı gelir) ama satırın tamamını görmeden m’yi hesaplayamazsınız. İşte kilit düğüm: tam satır görmek istiyorsanız Θ(N²) bellek malzemeleştirmeniz gerekir; malzemeleştirmemek istiyorsanız normalizasyonu ertelemeniz gerekir.

Online Softmax: Normalizatörü Akışlı Hesaba Çevirmek

Düğümü çözen fikir, Milakov & Gimelshein’in (2018) “online normalizer calculation for softmax” çalışmasında formelleşen ve Rabe & Staats (2021) üzerinden attention’a taşınan gözlemdir: softmax’ın istatistikleri (maksimum ve toplam) bir özyineleme (recurrence) ile adım adım güncellenebilir.

Satırı B bloğa ayırın, her bloğun kendi maksimumu m̃_b ve blok-toplamı ℓ̃_b olsun. Yürüyen (running) maksimum m ve yürüyen toplam ℓ tutarak blokları sırayla eritin:

İlk terim, önceki blokların eski maksimum referansına göre birikmiş toplamını yeni maksimuma yeniden ölçekler; ikinci terim yeni bloğun katkısıdır. Tümevarımla kanıt tek satırdır: her adımda ℓ_b = Σ_{i ≤ blok b} e^{x_i − m_b} korunur; son blokta m_B = m (global maksimum) ve ℓ_B = Σ_i e^{x_i − m} olur. Yani sonuç klasik, sayısal olarak stabil softmax ile matematiksel olarak eşdeğerdir; yaklaşıklık yoktur, yalnızca sıralama farkı vardır. Bir nüansı atlamayın: kayan-nokta toplamada sıralama değiştiği için yuvarlama farkları nedeniyle sonuç bit düzeyinde aynı değildir — pratikte ihmal edilebilir düzeyde yakındır, ama “birebir aynı” iddiası kâğıt üzerinde değil testte doğru olan bir iddiadır.

Aynı özyineleme, paydadaki toplamla birlikte biriktirilen ağırlıklı çıktı toplamına da genelleştirilebilir — ki FlashAttention’ın istediği tam olarak budur. Bir de pratik bir donanım detayı: GPU’lar için en hızlı üstel talimatı doğal tabanlı exp değil, exp2’dir. Bu yüzden FlashAttention, skorları 1/(√d_k · ln 2) ile ön-ölçekler ve tüm üstelleri exp2 cinsinden yazar; ln 2 sabiti yürüyen istatistiklerin ölçeklemesine aynı şekilde taşınır. Sonuç, değişmez bir matematiksel özdeşliktir: e^x = 2^x·log₂e.

FlashAttention: SRAM’ı Algoritmanın Tasarım Dili Yapmak

Dao ve arkadaşlarının 2022’de tanıttığı FlashAttention, yukarıdaki iki parçayı — tiling ve online softmax — GPU bellek hiyerarşisinin fiziksel gerçekleriyle birleştiren IO-farkındalıklı (IO-aware) tam bir attention algoritmasıdır. Anahtar ilke şudur:

N×N skor matrisi HBM’e asla yazılmaz. Her şey, çip üzerindeki hızlı SRAM’de (paylaşılan bellek / L1, A100’de SM başına ~192 KB, tüm çip ölçeğinde ~20 MB) olur biter; HBM ile SRAM arasındaki tek trafiği Q, K, V ve O gibi Θ(Nd) boyutundaki girdi/çıktılar taşır.

Forward geçişin iskeleti

Q ∈ ℝ^N×d satır bloklarına (boyut B_r), K, V ∈ ℝ^B_c×d sütun bloklarına bölünür. Dış döngü K/V blokları, iç döngü Q blokları üzerinde döner (v1’in seçimi; v2’de takas edilir). SRAM’de her adım için şunlar barındırılır:

• Q̃ (B_r×d), K̃ (B_c×d), Ṽ (B_c×d) girdi blokları
• Ara skor bloğu S̃ = Q̃K̃^T (B_r×B_c)
• Yürüyen istatistikler: satır maksimumu m ∈ ℝ^B_r, yürüyen toplam ℓ ∈ ℝ^B_r, yürüyen çıktı biriktirici Õ ∈ ℝ^B_r×d

Bloğun çekirdek güncellemesi (satır bazında, vektörel):

Son satır algoritmanın kalbidir ve iki terim taşır: biri eski blokların yeni maksimum referansına yeniden ölçeklenmiş birikimi, diğeri bugünkü bloğun değer katkısı P̃Ṽ. Sadece P̃ yazmak, normalizasyonun paydasını biriktirirken sayısını (çıktısını) biriktirmemektir; Õ hiçbir zaman doğru O’ya ulaşmaz. Doğru sıralama şudur: P̃ önce Ṽ ile çarpılır (B_r×B_c · B_c×d → B_r×d), sonra yürüyen çıktıya eklenir.

Tüm bloklar bitince tek bir normalizasyon: O = diag(ℓ)⁻¹·Õ. P matrisi hiçbir zaman tam olarak var olmaz; yalnızca SRAM’de doğar, katkısını Õ’ya verir ve yok olur.

Blok boyutları SRAM kapasitesi M (eleman cinsinden) ile kısıtlanır: K̃ ve Ṽ blokları (2·B_c·d), Q̃ ve Õ (2·B_r·d) ile S̃ tamponu (B_r·B_c) toplamı M’i aşmamalıdır. Makalenin pratik seçimi:

Yani K/V bloğu SRAM diliminden türetilir, sorgu bloğu ise aynı dilimin d ile sınırlanmış halidir; analizin geçerli olduğu aralık d ≤ M ≤ Nd koşuluyla ifade edilir.

IO karmaşıklığı ve optimallik kanıtı

FlashAttention’ın teorem düzeyinde iddiası, HBM-SRAM arasındaki toplam veri hareketinin:

olmasıdır; burada d = d_k, M = SRAM kapasitesi. Standart attention’ın Θ(N²) trafiğine oranı d²/M faktörüdür; bu oran 1’den küçük olduğu için FlashAttention daha az veri taşır — kazanımın büyüklüğü M/d² katıdır, d²/M değil. Sayıyla: d = 64 ve M ≈ 10⁵ eleman mertebesinde bir SRAM diliminde M/d² ≈ 24 katlık bir IO azalması söz konusudur. N×N matrisin bir daha asla HBM’e dokunmaması bu azalmanın kaynağıdır.

Dahası makale, bu trafiğin alt sınır olduğunu da kanıtlar: klasik Red-Blue pebble-game tekniğiyle (Hong & Kung, 1981) herhangi bir tam (exact) attention algoritmasının Ω(N²d²/M) HBM erişimi yapmak zorunda olduğu gösterilir. Yani FlashAttention IO anlamında optimaldir; ondan daha az veri taşıyan bir tam attention algoritması (modelleme varsayımları altında) var olamaz.

Hesap tarafında ise iş değişmez: runtime Θ(N²d) FLOP’tur. Burada yaygın bir atıf hatasını düzeltelim: tam attention için karesel zaman alt sınırını FlashAttention makalesi kanıtlamaz. Bu soru ince taneli karmaşıklık (fine-grained complexity) hattının konusudur. Alman & Song (2023) attention’ı belirli yumuşama (decay) koşulları altında kare-altı (subquadratic) sürede hesaplayan tam algoritmalar üretirken, genel durumda koşullu karesel alt sınırlar verdi. Daha yeni bir hat (Gupta ve ark., 2025) sonucu sıcaklığa genelledi: sabit d için Õ(n^2−1/d·polylog(B)) mertebesinde tam algoritma var; oysa d biraz büyümeye başladığında attention SETH altında n^2−o(1) süre gerektiriyor. Yani IO duvarı kırıldı; hesap duvarı, doğru varsayımlara koşullu olarak hâlâ ayakta. FlashAttention hızı azaltmaz; bellek duvarını aşar.

Kağıt üzerindeki sonuçlar

FlashAttention (v1) makalesinin raporladığı, standart implementasyonlara kıyasla: BERT-large eğitiminde %15 duvar-saati hızlanması, GPT-2 eğitiminde 3× hızlanma; GPT-2’de 2× daha uzun dizilerle eğitim kalitesinde ~0,7 perplexity iyileşme, uzun-belge sınıflandırmasında ~6,4 puan artış; Long Range Arena’da 2,4× daha uzun dizilere çıkabilme; Path-X’te 16K bağlamda %61,4 doğruluk, Path-256’da 64K bağlamda %63,1 doğruluk. İkinci gruptaki sayı hız değil kalitedir ve yüzde değil mutlak puandır: %15 daha düşük perplexity değil, ~0,7 perplexity daha iyi. Bunlar yaklaşım (approximation) içermeyen, matematiksel olarak standart attention ile eşdeğer çıktı üreten bir algoritmanın kazançlarıdır — ilk kez “daha akıllı yaklaşık attention” değil, aynı attention ama IO-farkındalıklı olanı mümkündür.

Maskeleme ve Dropout: Pratiği Belirleyen İki Küçük Detay

Somut bir GPT eğitiminin her adımı nedensel (causal) maskeyle çalışır: i. sorgu yalnızca j ≤ i anahtarlarına bakabilir. Bu, N×N matrisin yalnızca alt üçgenini hesaplamak demek olduğu için hesap ve bellek maliyetini yaklaşık yarıya indirir — ama tiling’in doğasında bir incelik yaratır.

FlashAttention blokları kare olduğu için K/V blokları üç durumdan birindedir: (1) maskeden tamamen etkilenen bölge — hiç hesaplanmaz, (2) tamamen geçerli bölge — normal akış, (3) diyagonal kesen bloklar — hesaplanır ama S̃’nin maskelenen girdilerine −∞ verilir. Sıra önemlidir: −∞ doldurması rowmax’tan önce yapılmalıdır; aksi halde yürüyen maksimum bir maske konumundan gelir ve satırın tamamı yanlış ölçeklenir. Etkilenen blokların atlanması, teorik Θ(N²d²/M) ifadesini pratikte ~2× daha ucuz hale getirir; asimptotik mertebe değişmez.

Dropout ise daha sinsi bir sorun: klasik implementasyon P matrisine dropout uygular, fakat FlashAttention P’yi hiç malzemeleştirmediği için maskeyi de saklayamaz. Çözüm deterministik rastgeleliktir: kernel, rastgele bir sayı üretecini (CUDA’da Philox) bir tohum ve sekans ofseti ile sürer; forward’da üretilen maskeyi saklamak yerine, backward’da aynı tohum-ofset çiftiyle olduğu gibi yeniden üretir. Böylece dropout, Θ(N²) maskelik ek bellek yerine sabit sayıda bayt (tohum + ofset) maliyetiyle çalışır. Bu, yeniden hesaplama (recomputation) ilkesinin sadece gradyan için değil, rastgelelik için de geçerli olduğunun zarif bir örneğidir.

Geri Yayılım: Hesabı Unutup Yeniden Kurmak (Recomputation)

Eğitimde asıl bellek maliyeti forward aktivasyonlarının saklanmasıdır: standart attention, geri yayılım için S ve P’yi (Θ(N²)) saklamak zorundadır. FlashAttention’ın geri geçişi bunun yerine yeniden hesaplamayı (recompute) seçer — çünkü hesabın marjinal maliyeti ucuz, veri taşımanınki pahalıdır.

Saklanan tek şey, satır başına iki sayıdır: çıktı O ve log-sum-exp istatistiği:

Böylece forward aktivasyon belleği Θ(N²)‘den Θ(N)‘e düşer. Geri geçişte, saklı L ve dO (başlangıç gradyanı) kullanılarak her P bloğu SRAM’de anında yeniden üretilir: P = exp(S − L) (satır bazında). Gradyan formülleri:

Bu, softmax’ın klasik gradyan özdeşliğinin (d softmax = p ⊙ (dp − ⟨dp⟩)) bloklara ve yeniden ölçeklemeye uyarlanmış halidir; D vektörü, normalizasyon teriminin satır bazlı toplamını temsil eder ve bir forward çıktısı olarak O’dan bedavaya hesaplanır. Toplam FLOP bütçesi forward’a ek ~2 kat daha yeniden hesap içerir (pratikte ~2,5 kat) — ama bu FLOP’lar SRAM’de, bant genişliği duvarına çarpmadan harcandığı için duvar saati (wall-clock) süresi yine de belirgin şekilde kısalır.

FlashAttention-2, -3 ve -4: Utilizasyon Savaşı

Algoritma bir kez IO-optimal olduktan sonra geriye kalan kayıp, GPU’nun içindeki mikro mimari sürtünmelerdir: matris çarpanı olmayan (non-matmul) CUDA çekirdek işleri, yetersiz paralellik ve warp senkronizasyonu.

FlashAttention-2 (Dao, 2023) üç mühendislik darbesiyle FA1’i ~2× hızlandırır ve A100’de tepe FP16 kapasitenin ~%73’üne kadar çıkar:

1. Döngü takası: Dış döngü Q blokları üzerinde döner; böylece farklı bloklar farklı thread-block’lara (ve GPU’lara/sequence’e) paralel dağıtılabilir — uzun dizilerde ve batch=1 eğitiminde darboğaz kırılır.
2. Non-matmul FLOP azaltımı: Yürüyen normalizasyonun bölme ve yeniden ölçekleme işleri iç döngüden çıkarılıp en sona ertelenir; biriktirici Õ, içeride normalize edilmeden toplanır.
3. Daha iyi warp bölümlemesi: Paylaşılan bellekte K/V tüm warp’lar arasında paylaşılan bir bloktur; warp’lar Q satırlarını kendi aralarında böler. Böylece her warp tek bir Q alt bloğunun sorumluluğunu üstlenir, K/V tekrar tekrar okunmaz ve warp’lar arası senkronizasyon ile yinelenen iş azalır.

FlashAttention-3 (Shah ve ark., 2024) NVIDIA Hopper (H100) mimarisinin asenkron birimlerine özel yazılır: WGMMA (warp-group matris çarpma) birimleri hesap yaparken TMA (Tensor Memory Accelerator) birimleri veriyi paralel taşır; bunu düzenlemek için üretici/tüketici (producer/consumer) warp-specialization iskeleti kurulur ve yazılım boru hattıyla (software pipelining) GEMM ile softmax’ın üstel/skala işleri eşzamanlı (overlap) çalışır. Sonuç: FP16’da H100’de FA2’ye göre ~1,5-2× hız ve tepe kapasitenin ~%75’ine varan SM utilisasyonu (ileri geçişte ~740 TFLOP/s mertebesi); FP8’de ise bloklar-arası ölçekleme ve “incoherent processing” (işaret rastgele çarpanları) ile hassasiyet kaybı bastırılarak ~1,2 PFLOP/s mertebesinde çıkar.

FlashAttention-4 (Zadouri ve ark., 2026) aynı savaşın Blackwell (B200) cephesidir; makalenin başlığı iddiayı doğrudan taşır: algoritma ile kernel boru hattının eş tasarımı (co-design). Blackwell’in tam asenkron tcgen05.mma matris birimi ve ayrı tensor memory (TMEM) adres alanı, daha büyük blok (tile) boyutları ile daha derin boru hatları mümkün kılar; 2-CTA MMA modu ve yazılımla taklit edilen (emüle edilen) üstel hesabı ile B200’de BF16’da ~1.613 TFLOP/s, yani tepe kapasitenin ~%71’i raporlanır — cuDNN’in ilgili kernelinden ~1,3×, Triton implementasyonundan ~2,7× hızlı. Not edilmeye değer iki gerçek: (i) FA2→FA3→FA4 boyunca asimptotik ifade hiç değişmedi, yalnızca %73 → ~%75 → ~%71 tepe-utilizasyonu ve sabit faktörler hareket etti; (ii) FA3 Hopper içindi, FA4’ün gerekliliği tam da yeni bir bellek/hesap topolojisinin (TMEM) eski çekirdeği yetersiz kılmasıdır. Bu bölüm, yazının geri kalanından farklı olarak eskiyecek bir katmandır: sayılar nesille değişir, Θ(N²d²/M) değişmez.

Fiziksel Katman: SRAM Hücrelerinden HBM Yığınlarına

Algoritmanın neden SRAM’e sığındığını anlamak için hücre fiziğine inmek gerekir. GPU’nun on-chip belleği 6T SRAM’den oluşur: çapraz bağlı iki eviriciden (inverter) oluşan dört transistör bir bit’i kilitler, iki erişim transistörü (access transistor) ise kelime satırı (wordline) işaretinde hücreyi bit satırlarına (bitline) bağlar.

Hücre başına altı transistör, SRAM’i hacim başına pahalı yapar; DRAM’in tek-transistör + tek-kapasitör (1T1C) hücresiyle yoğunluk yarışamaz. Ama bedeli ödenen şey hızdır: refresh döngüsü yoktur, erişim nanosaniyenin de altındadır ve A100’de ~20 MB’lık toplam SRAM’ın agregat bant genişliği ~19 TB/s mertebesindedir — HBM’in ~12 katı. HBM tarafında ise kapasite bol ama mesafe uzundur: DRAM çekirdekleri silisyum ara-yüzeyine dikey olarak (TSV, through-silicon via) yığılır, paket başına geniş bir arayüzle çipe bağlanır; 40-80 GB’lık alanın bedeli, saniyede ~1,5-2 TB’la tavan yapan bir boru hattıdır.

Dengesizliğin özeti: Çipin üretebildiği hesap, çipin beslenebildiği veriden çok daha hızlı büyüyor. FlashAttention bu eşitsizliği bir algoritma aksiyomu olarak kabul eder ve Θ(N²d²/M) formülündeki M’i —yani SRAM’i— algoritmanın birinci sınıf vatandaşı yapar.

Duvarın Ötesi: KV Cache, Uzun Bağlam ve Yaklaşık Attention

FlashAttention, eğitim ve tekil attention hesabının IO duvarını yıkır; ancak çıkarım sırasında başka bir Θ(N) bellek nesnesi sahneye çıkar: KV cache. Otoregresif üretimde her yeni token, geçmiş tüm tokenların K ve V tensörlerini yeniden hesaplamamak için önbelleğe alınır. Katman başına cache boyutu:

Örnek: 32 katmanlı, d_model = 4096’lık 7B sınıfı bir modelde FP16 ile token başına 2 × 4096 × 32 × 2 bayt ≈ 0,52 MB; 4096 token’lık bir oturum ≈ 2,1 GB. Bağlam uzadıkça KV cache, batch ölçeğinde HBM’i yeniden kuşatır — bu yüzden PagedAttention (vLLM), gruplandırılmış sorgu attention (GQA/MQA) ve MLA gibi yöntemler attention’ın IO probleminin çıkardıma özgü ikinci cephesini hedefler.

Alternatif hat ise O(N²)‘yi asimptotik olarak kırmaya çalışan yaklaşık/lineer attention aileleridir (Performer’ın rastgele özellikler (random features) yaklaşımı, Linformer’ın düşük-rank projeksiyonu, state-space modellerinin tamamen attention-dışı geçişi). FlashAttention’ın konumu burada nettir: kalite ödünü vermeyen, tam sonuç üreten bir IO optimizasyonudur — ve tam da bu yüzden pratikte endüstri standardı olmuştur.

Decode Aşaması: Flash-Decoding

Eğitim sırasında N sorgu bloğu arasında bol paralellik vardır. Çıkarımın decode adımı ise tam tersidir: sorgu uzunluğu 1’dir. Bir thread-block’u tek bir sorgu satırına ayrıldığında, GPU’nun onlarca SM’i işsiz kalır ve asıl iş —uzun bir KV dizisi boyunca tek bir satır okumak— tek bir çekirdeğe yüklenir. Tembel çözüm bunu iki kat kötüleştirir; bağlam uzadıkça attention adım başına Θ(N) okuma yapar ve duvar saati süresi bellek tavanına çarpar.

Flash-Decoding bunu anahtar-değer ekseni boyunca böler. Uzun KV dizesi C parçaya ayrılır; her parçayı ayrı bir SM işler ve her parça kendi yerel üçlüsünü üretir:

Sonra tek bir küçültme (reduction) adımı, online softmax’un yeniden ölçekleme özdeşliğini kullanarak parçaları birleştirir:

Dikkat: bu, forward pass’teki blok bloğun bir-birikimiyle birebir aynı özyinelemedir — yalnızca eksen değiştirmiştir: artık sorgu blokları değil, anahtar-değer blokları arasında paralelleşiyoruz. Aynı log-sum-exp istatistiği geri yayılımda yeniden hesapmanın, decode’da ise böl-ve-birleştirme adımlarının anahtarı olur; bir satırlık iki kayan-nokta sayısının bu kadar şey taşıması, bu algoritmanın en estetik noktasıdır.

Kazanım ölçekte gizlidir: küçük batch ve uzun bağlamda paralelleştirilecek iş olmadığı için iyileşme belirgindir — attention adımının kullanabilir paralelliği pratikte C katına çıkar, yani bölünen KV parça sayısıyla orantılı olarak ölçeklenir. Büyük batch’te ise sorgu ekseni zaten dolu olduğundan fayda mütevazı kalır. Flash-Decoding ile PagedAttention birlikte, uzun-bağlam çıkarımının standart donanım setini oluşturur.

Sonuç: Asimptotik Yeterli Değil, Fiziksel Okuryazarlık Şart

Bu makalenin taşıdığı tek bir mühendislik dersi varsa o da şudur: bir algoritmanın bellek karmaşıklığını hesaplama karmaşıklığından ayrı tasarlamayan her optimizasyon, bellek duvarına çarpmaya mahkumdur. Self-attention’ın O(N²) maliyeti on yıl boyunca “daha az FLOP” ekseninde çözülmeye çalışıldı; kırılma, FLOP’u hiç azaltmadan yalnızca veri hareketini Θ(N²)‘den Θ(N²d²/M)‘ye indiren bir IO-farkındalıklı tasarımla geldi. Online softmax, satırı bütün olarak görme zorunluluğunu bir özyinelemeye; tiling, o özyinelemeyi SRAM’in fiziksel diline çevirdi. Maskeleme, dropout ve Flash-Decoding ise aynı özyinelemenin —kayıtlı tutulan iki sayı— üç ayrı cephede yeniden kullanıldığını gösterdi.

Principia’nın açılış ilkesi budur: Sayılar yalan söylemez; formüller eskimez. Bir sonraki yazıda, bu kez bellek hiyerarşisinin diğer ucuna — HBM’in kendisine ve çipler-arası ölçeklemenin IO fiziğine — ineceğiz.

---

Kaynaklar

• Vaswani ve ark., Attention Is All You Need, NeurIPS 2017 — arxiv.org/abs/1706.03762
• Dao, Fu, Ermon, Rudra, Ré, FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness, NeurIPS 2022 — arxiv.org/abs/2205.14135
• Tri Dao, FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning, ICLR 2024 — arxiv.org/abs/2307.08691
• Shah, Bikshandi, Zhang, Thakkar, Ramani, Dao, FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision, NeurIPS 2024 — arxiv.org/abs/2407.08608
• Zadouri, Hoehnerbach, Shah, Liu, Thakkar, Dao, FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling, 2026 — arxiv.org/abs/2603.05451
• Milakov & Gimelshein, Online Normalizer Calculation for Softmax, NVIDIA 2018 — arxiv.org/abs/1805.02867
• Rabe & Staats, Self-Attention with Concurrent GPU Streaming and Memory-Efficient Exact Softmax, 2021 — arxiv.org/abs/2104.06640
• Wulf & McKee, Hitting the Memory Wall: Implications of the Obvious, ACM CCR 1995
• Williams, Waterman & Patterson, Roofline: An Insightful Visual Performance Model for Multicore Architectures, CACM 2009
• Hong & Kung, I/O Complexity: The Red-Blue Pebble Game, STOC 1981
• Alman & Song, FlashAttention and Longer-Range Transformers: New Fast and Exact Algorithms and Beyond, 2023 — arxiv.org/abs/2306.12150
• Gupta, Huang, Saha, Xu & Ye, Subquadratic Algorithms and Hardness for Attention with Any Temperature, 2025 — arxiv.org/abs/2505.14840
• NVIDIA A100 Tensor Core GPU Mimarisi Beyaz Kağıdı (2020); NVIDIA H100 (2022) ve Blackwell B200 platformu dokümantasyonu