Hacker News Digest

• Новое волокно с полым сердечником (DNANF) показало рекордное затухание 0,091 дБ/км на 1550 нм — в 1,5 раза лучше кремниевых волокон (0,14 дБ/км).
• Скорость передачи выше на 45 %; низкие потери (<0,2 дБ/км) сохраняются в диапазоне 66 ТГц (700–2400 нм).
• Удалось подавить три механизма потерь: утечку, рассеяние на границе и микроизгиб; испытано на 15 км.
• Технология обещает увеличить дальность без усилителей, снизить стоимость бита и открыть новые длины волн для связи, лидаров и высокомощной лазерной доставки.

• Hollow-core fiber (HCF) передаёт свет почти со скоростью света в вакууме (~c), в отличие от обычного стекловолокна (~2/3 c), снижая задержку на 30–45 %.
• Ключевое преимущество — не пропускная способность, а латентность: важно для HPC, AI-кластеров, онлайн-игр и трансконтинентальных линий.
• Потери 0,2 дБ/км и полоса 66 ТГц уже достигнуты, но сплайсинг и механическая надёжность пока требуют спецоборудования и персонала.
• Существующее стекловолокно не устарело: для большинства «последней мили» и домашних 1–10 Гбит/с оно всё ещё в десятки раз недогружено.
• Массовое развертывание HCF начнётся с сверхдлинных магистралей и HPC-центров, а до FTTH доберётся не раньше чем через десятилетия, если вообще доберётся.

Flash Attention на 5090 в CUDA C++

Цель — научиться писать attention-ядро на CUDA C++, чтобы использовать MXFP8/NVFP4 MMA для sm120, чего нет в Triton.
Код: learn-cuda/07_attention.

Бенчмарк (bs=1, heads=8, q=4096, kv=8192, BF16, 5090@400 W, CUDA 12.9, SOL 209.5 TFLOPS):

ядро	TFLOPS	%SOL
F.sdpa (Flash)	186.73	89.13
F.sdpa (CuDNN)	203.61	97.19
flash-attn	190.58	90.97
v1 (basic)	142.87	68.20
v2 (swizzle)	181.11	86.45
v3 (2-stage)	189.84	90.62
v4 (ldmatrix.x4)	194.33	92.76
v5 (pipe)	197.74	94.39

---

Алгоритм Flash Attention 2

Псевдокод:

scale = DIM**-0.5
for b, tile_Q:
    tile_O = 0
    tile_Q = load(Q[b, tile_Q])
    for tile_KV:
        tile_K = load(K[b, tile_KV])
        tile_S = tile_Q @ tile_K.T * scale
        online_softmax(tile_S)  # in-place
        tile_V = load(V[b, tile_KV])
        tile_O += tile_S @ tile_V
    store(O[b, tile_Q])

head_dim=128 помещается в регистры.

---

v1 — базовая версия

1. G2S: cp.async.ca.shared.global 128-битными транзакциями.
2. S2R: ldmatrix для Q, K, V → 8×8 фрагменты.
3. Softmax online:
   • m = max(m_prev, m_curr)
   • d = d_prev * exp(m_prev - m) + Σ exp(S - m)
   • O = O_prev * (d_prev/d) * exp(m_prev - m) + (exp(S - m)/d) @ V

---

v2 — swizzled shared memory

• 128-битные банки → конфликты при 8×8 tile.
• Swizzle K и V по 32-битным строкам; Q оставляем линейно.
• +40 % пропускной способности.

---

v3 — 2-stage pipeline

• Двойной буфер: пока вычисляем S/P@V, асинхронно грузим следующий KV.
• cp.async.commit_group() + cp.async.wait_group(1).
• +5 % к SOL.

---

v4 — ldmatrix.x4

• Одна инструкция ldmatrix.x4 загружает 4×8×8 фрагмента K/V за раз.
• Снижает инструкций на 25 %.
• +2 % к SOL.

---

v5 — улучшенный pipeline

• 3-4 стадии:
  1. prefetch KV
  2. compute S
  3. compute P@V
  4. write-back O
• __pipeline_wait_prior(N) + __pipeline_commit().
• +2 % к SOL.

---

Что дальше

• Использовать TMA (cp.async.bulk) и NVFP4/MXFP8 MMA.
• Поддержка head_dim > 128 (FlashMLA).

• Пользователи удивлены, что RTX 5090 даёт всего 209 TFLOPS BF16 — менее 10 % от серверного Blackwell B200 (2250 TFLOPS), но при цене ~$30-40 k за B200 соотношение цена/производительность почти сравнялось.
• Обсуждают, что NVIDIA с 4090 и далее искусственно ограничивает тензорные ядра игровых карт для ML-операций FP8/FP16.
• У 5090 выше TDP, чем у 4090, и можно ограничить мощность лишь до 70 % (4090 — до 50 %), что мешает апгрейду для ML-станций.
• Появились вопросы о поддержке Flash Attention на 5090/5080 и о нативной компиляции под Blackwell в PyTorch 2.7.
• Участники спорят, стоит ли вкладываться в Triton, если нужны фирменные типы NVFP4/MXFP8, которых там пока нет.