Архитектуры AI-чипов: от GPU NVIDIA до TPU, Cerebras и Groq

На Международном симпозиуме по компьютерной архитектуре в 2018 году Джон Хеннесси и Дэвид Паттерсон прочитали туринговскую лекцию «Новый золотой век компьютерной архитектуры».

В 1980-х, когда Хеннесси и Паттерсон вели исследования, за которые позже получили премию Тьюринга, производительность однопоточного процессора росла на 52% в год. К 2018 году, с окончанием действия закона Мура и масштабирования Деннарда, этот показатель упал до 3%.

Возникла потребность в предметно-ориентированных архитектурах (DSA). Показательным примером служил Google TPU v1, уже находившийся в промышленной эксплуатации: 29-кратная пропускная способность по сравнению с CPU на инференсе нейросетей при 80-кратно лучшей энергоэффективности. Заключительный прогноз лекции гласил: «следующее десятилетие увидит кембрийский взрыв новых компьютерных архитектур».

Прогноз сбылся. Сегодня в серьёзной разработке находятся десятки архитектур: GPU, TPU, LPU, NPU, DPU, ASIC, wafer-scale двигатели, реконфигурируемые dataflow-чипы, нейроморфные, фотонные, аналоговые решения. В первую очередь эти архитектуры нацелены на вычисления для ИИ.

Архитектуры, добившиеся реального промышленного развёртывания: GPU (NVIDIA, AMD), систолические матричные ускорители (TPU, Trainium), Cerebras Wafer-Scale Engine и Groq LPU.

NVIDIA — явный лидер; AMD следует за ней, заручившись обязательствами на 6 ГВт со стороны OpenAI и Meta. TPU обучают Gemini и будут обслуживать Anthropic парком до миллиона чипов; сама Anthropic также запускает Claude на более чем миллионе чипов Trainium. Cerebras теперь обслуживает инференс OpenAI; Groq LPU был поглощён NVIDIA за $20 млрд.

Далее рассматриваются эти различные подходы: их философия, архитектура, методы масштабирования (scale-up и scale-out) и программный стек (как программируется чип).

Проблема

Вычисления для ИИ доминирует умножение матриц. Трансформер — это последовательность матричных умножений: проекция Q/K/V, внимание, выходная проекция, FFN — перемежаемая поэлементными операциями: нормализация, активация, остаточные сложения. Обучение передовой модели выполняет порядка 10²⁵ операций умножения-накопления (матричные умножения — это последовательность таких операций).

Форма этих матричных умножений зависит от рабочей нагрузки. Обучение прогоняет батч последовательностей через каждый слой вперёд, распространяет ошибку назад и обновляет веса, причём тысячи токенов проходят через одну и ту же весовую матрицу одновременно. Prefill — это фаза приёма промпта при инференсе: вся входная последовательность проецируется через модель за один проход, до того как сгенерирован первый выходной токен. И обучение, и prefill складывают множество токенов против одной весовой матрицы, поэтому математика каждого слоя — это большое умножение матрица-на-матрицу (GEMM) с высокой арифметической интенсивностью (вычисления являются узким местом). Decode — авторегрессивен: модель выдаёт по одному токену за раз, каждый обусловлен всеми предыдущими, и токен N+1 не может начаться, пока не сгенерирован токен N. За шаг проецируется лишь один токен, поэтому каждое матричное умножение становится произведением матрица-на-вектор (GEMV). Генерация одного токена требует полного прохода по всем весам модели плюс полного чтения KV-кэша для внимания. Арифметическая интенсивность падает на порядки по сравнению с prefill.

Системы инференса восстанавливают часть этой интенсивности, объединяя токены в батчи, чтобы превратить GEMV обратно в GEMM: continuous batching складывает шаги decode многих пользователей вместе, спекулятивное декодирование складывает K черновых токенов на запрос и верифицирует их за один проход, а multi-token prediction встраивает тот же трюк внутрь самой модели. Это повышает загрузку матричных блоков и увеличивает соотношение операций к байтам. При continuous batching каждый запрос пользователя всё равно читает собственный KV-кэш, поэтому decode с длинным контекстом смещается от ограничения по пропускной способности весов к ограничению по пропускной способности KV-кэша.

Архитектурная задача здесь — перемещение чисел туда, где происходит матричное умножение, достаточно быстро. Это известно как стена памяти: вычислительная мощность росла экспоненциально, пропускная способность памяти — нет.

Каждая архитектура предлагает свою стратегию для выигрыша в игре по перемещению данных. Понимание чипа сводится к четырём вопросам: где живут данные, как они перемещаются к вычислительным блокам, как устроены сами вычислительные блоки, и как чипы общаются друг с другом в масштабе.

NVIDIA GPU

NVIDIA GPU — это массивно-параллельный процессор. Философия такова: программируемый чип с тысячами потоков, оркестрируемый хост-процессором и доступный через CUDA, — правильная машина для распараллеливаемых нагрузок. Каждое поколение добавляет ускоряющие примитивы поверх программируемых Streaming Multiprocessor, не меняя модель программирования. Один и тот же чип обучает трансформеры, обслуживает инференс, рендерит графику и выполняет научные симуляции (ускоренные вычисления).

Родословная

2006 — Tesla (G80): первый CUDA-совместимый GPU; унифицированные шейдеры и модель исполнения SIMT.
2010 — Fermi (GF100): первая настоящая вычислительная архитектура: унифицированные кэши L1/L2, два варп-планировщика, IEEE-754 FP64.
2012 — Kepler (K20, K40): SMX, динамический параллелизм, Hyper-Q; GPU может запускать собственную работу.
2014 — Maxwell (M40): переработанный SM с ~2× производительностью на ватт по сравнению с Kepler.
2016 — Pascal (P100): NVLink 1.0, HBM2, нативная пропускная способность FP16; первый GPU, спроектированный явно для глубокого обучения.
2017 — Volta (V100): первые Tensor Core; независимое планирование потоков.
2018 — Turing (T4): Tensor Core второго поколения с INT8/INT4; первые RT Core.
2020 — Ampere (A100): Tensor Core третьего поколения с TF32 и структурированной разреженностью; разделение Multi-Instance GPU.
2022 — Hopper (H100, H200, GH200): Tensor Core четвёртого поколения, FP8, Transformer Engine; HBM3, TMA, кластеры тред-блоков, асинхронный wgmma.
2024 — Blackwell (B100, B200, GB200): Tensor Core пятого поколения с FP4, Tensor Memory (TMEM), двухчиповый чиплетный GPU, NVLink 5.
2025 — Blackwell Ultra (B300, GB300): обновление посреди цикла: ~1,5× пропускной способности FP4, 288 ГБ HBM3e. Настроен под рассуждения с длинным контекстом.
2026 — Rubin (Rubin, VR200, Rubin CPX): HBM4, Transformer Engine третьего поколения, пара с CPU Vera, дезагрегированный prefill через Rubin CPX.
2027 — Rubin Ultra: четырёхчиповый пакет GPU, 1 ТБ HBM4e на пакет. Развёрнут в стойках NVL576 Kyber на 600 кВт при 100 ПетаFLOPS FP4 на GPU.

Архитектура

NVIDIA GPU — это группа ориентированных на пропускную способность ядер, глубокая иерархия памяти для их питания и ровно столько планирующей логики, сколько нужно для одновременного выполнения тысяч потоков. Ядра — это Streaming Multiprocessor, повторённые более 100 раз на пакет: 80 на V100, 108 на A100, 132 на H100, 148 на B200, 160 на B300, 224 на Rubin. Внутри каждого SM — одинаковый набор: четыре SM Sub-Partition, каждый со своим варп-планировщиком, диспетчером, регистровым файлом на 16k×32-бит, скалярными ядрами CUDA, специальным функциональным блоком для трансцендентных функций и приватным портом в Tensor Core своего SM. Четыре партиции делят блок L1/shared memory и TMA. Потоки группируются в варпы по 32, исполняющиеся в жёсткой синхронности SIMT; десятки резидентных варпов на партицию позволяют планировщику скрывать простои памяти/арифметики переключением между ними.

Компьют

CUDA Core — исходная вычислительная мощность, и для ИИ они по-прежнему обслуживают всё, что не является матричным умножением: активации, остаточные сложения, нормализацию, адресную арифметику. Но блок трансформера — это ~99% FLOPs на матричные умножения, поэтому подавляющая вычислительная мощность идёт от Tensor Core.

Эти ядра выполняют слитное умножение-накопление матриц на небольших плитках: D = A·B + C. Полное матричное умножение разбивается на выходные плитки: чтобы получить одну выходную плитку, ядро проходит общую внутреннюю размерность K, беря A из полосы строк левой входной матрицы и B из полосы столбцов правой, и сворачивает каждое частичное произведение в накопитель. C хранит текущую частичную сумму, D — обновлённое значение, переносимое на следующий шаг. После завершения внутреннего цикла D — это одна готовая плитка полной выходной матрицы; всё матричное умножение строится из множества таких плиточных MMA.

Формы плиток записываются как M × N × K, где M × N — размер выходной плитки, а K — сколько внутренней размерности инструкция сворачивает за один запуск; остальная часть оси K матричного умножения обходится внутренним циклом ядра. Накопитель «липкий» на протяжении этого цикла: выход D каждого MMA становится входом C следующего MMA, поэтому уравнение на самом деле выглядит как C ← A·B + C на месте: последовательные инструкции сворачивают свои частичные произведения в одно и то же хранилище, пока ось K не будет полностью пройдена.

Первое поколение блоков V100 (8 на SM) выполняло варп-уровневый MMA 16×16×16 FP16. Третье поколение A100 добавило TF32, BF16, FP64-умножение и структурированную разреженность 2:4. Четвёртое поколение H100 добавило нативный FP8 и подняло абстракцию с варпа до группы варпов: 128 кооперирующихся потоков, запускающих асинхронный wgmma формы 64×256×16, работающий в фоне, пока запускающие варпы загружают следующую плитку. Пятое поколение B200 пошло ещё дальше: MMA на двух SM формы 256×256×16 с операндами, разделёнными между парой SM, нативный FP4 и выделенный скретчпад Tensor Memory (TMEM) на 256 КБ на SM, хранящий плитки-накопители вместо того, чтобы перетекать в регистровый файл. Шестое поколение Rubin расширяет пропускную способность FP4, добавляет нативный FP6 и работает в паре с Transformer Engine третьего поколения, который выполняет адаптивное микроблочное масштабирование NVFP4 аппаратно, оставляя метаданные квантования плитки на пути Tensor Core, а не через CUDA Core.

Что остаётся неизменным во всех шести поколениях — матричное умножение живёт внутри иерархии тредов/варпов, но число потоков, требуемых для его запуска, сократилось, а сам запуск отделился от исполнения. Инструкция Volta mma.sync коллективна на уровне варпа и синхронна: все 32 потока варпа выполняют её вместе, каждая линия держит регистровые фрагменты A, B и накопителя D, и варп блокируется до завершения. Hopper-овский wgmma.mma_async расширяет исполнителя до группы варпов из 128 потоков, перемещает B в дескриптор общей памяти (A становится опциональным: либо регистры, либо дескриптор, на выбор ядра) и возвращается немедленно: матричное умножение выполняется в фоне, пока группа варпов ставит в очередь следующую плитку, а завершение отслеживается через wgmma.commit_group / wgmma.wait_group.

Blackwell-овский tcgen05.mma завершает эту миграцию: A присоединяется к B в дескрипторах общей памяти (либо A берётся напрямую из TMEM), а накопитель D попадает в TMEM, а не в регистровый файл. Поскольку все операнды сняты с линий, координировать запуск на уровне отдельных потоков больше не нужно, поэтому единственный поток запускает инструкцию и сразу возвращается, а о завершении сигнализирует mbarrier, которого ждёт варп-потребитель. Остальная часть варпа и сам запускающий поток тем временем свободны для другой работы. Вариант с парой CTA масштабирует ту же модель на два SM: по одному потоку на каждом SM в спаренном кластере запускают согласованные MMA, разделяющие операнды между собой, составляя плитку 256×256×16 на двух SM под тем же асинхронным завершением через mbarrier, только поднятым до барьера уровня кластера, чтобы пара оставалась синхронизированной.

Матричное умножение стало одновременно и крупнее, и легче для запускающих потоков: инструкция, начинавшаяся как 32 линии, действующие в лок-степ, теперь ближе к единой команде, управляемой дескриптором, диспетчеризуемой изнутри модели варпов, но больше не исполняемой ею.

Именно это разделение делает ядра внимания трансформера эффективными на GPU. Варп может выполнять softmax, применять маску или предзагружать следующую плитку, пока матричное умножение находится в полёте; наложение матричного умножения и окружающей поэлементной работы — структура каждого современного ядра внимания (FlashAttention-3, FA4), и оно зависит от того, что матричная инструкция не блокирует варп.

Память

Иерархия на чипе — это аппаратно управляемые кэши на каждом уровне, с программными подсказками поверх. За пределами чипа — HBM: 32 ГБ HBM2 на V100, 80 ГБ HBM3 на H100, 192 ГБ HBM3e на B200, 288 ГБ на B300, 288 ГБ HBM4 на Rubin. Между HBM и SM находится L2-кэш уровня чипа: 6 МБ на V100, 40 МБ на A100, 50 МБ на H100, 60 МБ на B200 (разделён на два банка по 30 МБ через двухчиповый пакет, с контролем локальности, позволяющим закреплять «горячие» плитки за ближним чипом). Внутри каждого SM 256 КБ унифицированной L1/SMEM разделяются при запуске ядра между аппаратно управляемым L1 и программируемым скретчпадом. Регистровый файл — ещё около 256 КБ на SM, разрезанный на четыре части по партициям.

Blackwell добавляет пятый уровень: TMEM, 256 КБ на SM, выделенные под накопители MMA и адресуемые только Tensor Core, снимая давление резидентности операндов с общего регистрового файла.

Перемещение между уровнями постепенно отделялось от варпа. До Ampere загрузка плитки была синхронной: каждый поток выполнял собственную глобальную загрузку, варп блокировался до тех пор, пока каждый фрагмент не попадал в регистры, и второй проход копировал их в общую память; каждая плитка сжигала линии варпа на адресную арифметику и ожидание. Ampere ввёл cp.async: асинхронные копии по потокам из HBM в SMEM, полностью минуя регистры, причём варп фиксирует группы копий в полёте и ждёт только тогда, когда потребителю действительно нужны данные. Hopper заменил это на TMA, выделенный движок DMA: один поток отправляет многомерный дескриптор плитки (базовый адрес, ведущая размерность, свизл), движок сам занимается адресной арифметикой и пишет в общую память, а о завершении сигнализирует mbarrier. Весь варп освобождается от выпуска загрузок и адресной математики; ядро просто ставит дескрипторы в очередь. TMA также поддерживает мультикаст на уровне кластера: одно чтение HBM разветвляется на каждый SM в кластере тред-блоков, превращая то, что раньше было N отдельными загрузками, в одну. Blackwell снова расширяет TMA: прямые загрузки в TMEM, так что плитки-накопители поступают без промежуточного этапа в SMEM. Тенденция — на одну операцию меньше для варпа на плитку, поколение за поколением.

Специализация варпов

Программная идиома эпохи Hopper — специализация варпов: внутри одного блока некоторые варпы выступают производителями, выполняя загрузки TMA одну за другой; другие — потребителями, запускающими wgmma на только что прибывших плитках. Синхронизация между ними больше не старый барьер __syncthreads() уровня всего SM; это mbarrier (барьеры памяти в общей памяти) и асинхронные транзакционные барьеры, привязанные к завершениям TMA, позволяющие тонкое взаимодействие производитель/потребитель на уровне варпа, а не блока. Паттерн, ставший эталоном для каждого современного ядра внимания (FlashAttention-3, ping-pong GEMM в CUTLASS, ядро FA4 для Blackwell), един: конвейер производителей на базе TMA питает конвейер потребителей wgmma через общую память и TMEM, с рукопожатиями mbarrier и кластерами тред-блоков (Hopper и новее), объединяющими несколько SM в один кооперативный вычислительный блок, поверх которого естественно ложится MMA на двух SM в Blackwell.

Численные форматы

FP32 был историческим значением по умолчанию; Volta принесла FP16 с накоплением в FP32 и трюками масштабирования потерь, сделавшими его пригодным для обучения; Ampere добавила TF32 (диапазон FP32, мантисса FP16, прямая замена для матричного умножения FP32), BF16 и структурированную разреженность 2:4, удваивающую эффективную пропускную способность на прореженных весах. Hopper ввёл нативный FP8 в вариантах E4M3 и E5M2 в паре с Transformer Engine, автоматически масштабирующим активации слой за слоем, удерживая их в динамическом диапазоне FP8. Blackwell снова вдвое снизил точность с FP4 и добавил микромасштабные форматы MX (общие экспоненты на уровне блока, восстанавливающие большую часть точности, потерянной на FP4), вместе с Transformer Engine второго поколения, перенаправляющим конвейер автомасштабирования на FP4. Transformer Engine третьего поколения в Rubin добавляет NVFP4 (ужесточённый вариант FP4 от NVIDIA) и нативный FP6 с более агрессивной разреженностью. Сама компоновка чипа теперь часть истории численных форматов: B100/B200/B300 — это два чипа на пределе размера кристалла, сшитые связью NV-HBI на ~10 ТБ/с и представленные программному обеспечению как один логический GPU с 8 стеками HBM на пакет; Rubin расширяет чиплетную схему до двух кристаллов на ~336 млрд транзисторов с 8 стеками HBM4. Каждое поколение выигрывает примерно 2× пропускной способности на ватт, вдвое урезая биты и восстанавливая точность более тонкой схемой масштабирования, и всё чаще — присоединяя больше кремния в пакет.

Ставки

• Ставка 1: Программируемость. Рабочая нагрузка — движущаяся цель (варианты внимания, новые архитектуры моделей), поэтому каждый блок остаётся программируемым, а разработчик пишет CUDA. Даже специализированные блоки доступны через эту модель, а не как фиксированные функциональные блоки.
• Ставка 2: Скрытие задержек массовой многопоточностью. Задержка непредсказуема и зависит от данных, поэтому она скрывается не статическим расписанием, а массовым избытком потоков — до 64 резидентных варпов на SM — с аппаратным варп-планировщиком, выбирающим готовый варп каждый такт.
• Ставка 3: Матричное умножение, обёрнутое в варп. Матричный блок даёт подавляющую вычислительную мощность, но он должен жить за той же абстракцией варпа/потока, что и всё остальное, поэтому он обёрнут в mma.sync → wgmma → tcgen05.mma, а не выставлен как фиксированный функциональный конвейер. Это позволяет одному ядру сливать матричное умножение, softmax и поэлементные операции за один проход.
• Ставка 4: Асинхронная иерархия памяти. Иерархия памяти сделана явной и программно управляемой, а не неявной и планируемой компилятором. L2-кэш сохраняется, но SMEM и TMEM выставлены как именованные скретчпады, а поверх наложена асинхронная машинерия: TMA для массовых копий, TMEM для накопителя матричного умножения, mbarrier для рукопожатия производитель/потребитель. Иерархия программно конвейеризована внутри программируемого ядра, а не статически спланирована компилятором против скретчпада с известной задержкой.
• Ставка 5: Амортизированный налог SIMT. Каждый транзистор, потраченный на варп-планировщик, регистровый файл или когерентный кэш — это транзистор, не потраченный на MAC; налог принимается, но окупается двумя путями: Tensor Core теперь достаточно велик, чтобы амортизировать логику SIMT на гораздо большем числе MAC, а блоки вроде TMEM обменивают часть гибкости общего назначения на плотность MAC.

Масштабирование

Есть два режима масштабирования: scale-up и scale-out.

Scale-up — объединение нескольких GPU в один когерентный домен памяти. Любой GPU может загружать или записывать HBM любого другого GPU напрямую через NVLink с наносекундными задержками: одно адресное пространство, без явных передач.

Scale-out — объединение этих доменов в сеть на уровне стойки и кластера. Данные пересекают границы через явный RDMA с микросекундными задержками: раздельные адресные пространства, но десятки тысяч чипов на кластер.

ИИ-инфраструктура использует оба режима: требовательные к пропускной способности коллективные операции (тензорный параллелизм, маршрутизация экспертов MoE) остаются внутри домена scale-up; параллелизм данных и конвейерный параллелизм пересекают фабрику scale-out.

Scale-up

Стек scale-up — это NVLink плюс NVSwitch. NVLink реализует кэш-когерентную фабрику между GPU, так что загрузка или запись на одном GPU может нацеливаться на HBM другого GPU, а аппаратура сама занимается трансляцией адресов и когерентностью. Но сам по себе NVLink точка-точка: одна связь соединяет ровно два чипа. NVSwitch — выделенный кроссбар-чип, к которому подключается каждый GPU, маршрутизирующий трафик так, чтобы каждый GPU мог одновременно общаться с любым другим на полной пропускной способности NVLink, неблокирующе и по схеме «все-со-всеми».

Вместе они определили базовую плату HGX на 8 GPU, объединяющую восемь модулей H100 SXM с хостами x86 (AMD EPYC или Intel Xeon) через PCIe Gen5. Hopper также поставлялся в форме, спаренной с Grace: GH200 Grace Hopper Superchip соединял один ARM-процессор Grace с одним H100 через NVLink-C2C на скорости 900 ГБ/с, устраняя переход хост-устройство через PCIe. Модули масштабировались в пары GH200 NVL2 и стоечные конфигурации GH200 NVL32. Blackwell делает такое спаривание стандартом. Модуль GB200 объединяет один Grace с двумя B200 через NVLink-C2C, а NVL72 сшивает 36 таких модулей в единый жидкостно охлаждаемый домен scale-up: 72 GPU, 36 процессоров Grace, 13,5 ТБ HBM и 17 ТБ LPDDR5X как одно плоское когерентное адресное пространство. Rubin делает следующий шаг в два этапа. NVL144 выходит в 2026 году как обновление поколения Rubin внутри той же стойки класса Oberon: 72 пакета Rubin, обозначенные как 144 GPU по новой системе подсчёта NVIDIA, с HBM4 и NVLink 6, удваивающими пропускную способность на пакет. Настоящий скачок на уровне стойки — Rubin Ultra в 2027 году: NVL576 упаковывает 144 четырёхчиповых пакета Rubin Ultra в новое шасси Kyber, получая 576 чипов GPU в одном когерентном домене.

Такая плотность держится на пассивной меди. Фабрика NVLink в NVL72 работает через 5184 кабеля с слепым сочленением через объединяющую панель (~2 мили кабелей на стойку, без ретаймеров внутри кабеля, SerDes живёт прямо на GPU и коммутационных ASIC), обеспечивая ~130 ТБ/с пропускной способности «все-со-всеми» на 72 GPU. По оценке NVIDIA, выбор меди экономит примерно 20 кВт на стойку по сравнению с оптическим эквивалентом, требующим подключаемых трансиверов на каждой связи. Медь делает идею «стойка как один GPU» экономически осуществимой: на дистанциях меньше 2 метров она по-прежнему выигрывает по энергопотреблению, стоимости и целостности сигнала на доллар; за этим пределом биты приходится передавать по стеклу.

NVL144 остаётся внутри Oberon, и медь продолжает работать, поскольку число пакетов (72) не изменилось относительно NVL72; кабели не нужно удлинять, только передавать быстрее по SerDes поколения 6. NVL576 в Rubin Ultra сохраняет тот же медный предел, изменив форму стойки: новый форм-фактор Kyber примерно вдвое выше Oberon и упаковывает все 576 чипов GPU в один корпус, спроектированный так, чтобы каждый путь NVLink оставался в пределах досягаемости пассивной меди даже при 144 четырёхчиповых пакетах и десятках тысяч кабелей.

Scale-out

Стек scale-out происходит из приобретения Mellanox. В отличие от NVLink, фабрики scale-out не когерентны: узлы сохраняют раздельные адресные пространства, и данные пересекают границы только через явный RDMA, инициируемый программно, обычно обёрнутый в коллективные операции NCCL вроде all-reduce или all-to-all. Эталонный кластер — DGX SuperPOD: восемь стоек NVL72, сшитых через InfiniBand Quantum-X800, дают 576 GPU Blackwell под одним планировщиком, а обучающие кластеры масштабируются дальше плиткованием SuperPOD. Rubin SuperPOD в 2026 году сохраняет тот же паттерн из 8 стоек с NVL144 (получая 1152 GPU на SuperPOD вместо 576). Rubin Ultra в 2027 году масштабирует схему на порядок: стойки Kyber по 576 чипов GPU каждая, сшитые через Quantum-X Photonics CPO, объединяя тысячи GPU под одним планировщиком.

У каждого GPU есть собственная сетевая карта ConnectX в эту фабрику. Узлы Blackwell работают на ConnectX-8 со скоростью 800 Гбит/с на GPU — на порядок меньше пропускной способности NVLink на GPU, и задержки растут с наносекунд до микросекунд. Rubin переходит на ConnectX-9 со скоростью 1,6 Тбит/с на GPU, удваивая пропускную способность scale-out на GPU по мере роста домена scale-up на стойку с 72 до 576 GPU. Рядом с каждой сетевой картой стоит DPU BlueField, добавляющий ARM-ядра и ускорители для разгрузки хранилища, сети и безопасности с хост-процессора. Для клиентов, предпочитающих Ethernet InfiniBand, есть Spectrum-X — альтернатива на базе lossless-Ethernet, настроенная под ИИ-трафик.

Переход от меди к стеклу происходит на границе стойки. Внутри NVL72 позвоночник — медный; как только связь должна пересечь стойки на скорости 800 Гбит/с, она становится оптической. Пассивный медный DAC ограничен примерно 1,5–2 метрами на скорости 200 Гбит/с на канал — существенно меньше межстоечной дистанции, поэтому сегодняшний позвоночник SuperPOD работает через подключаемые трансиверы OSFP-RHS, каждый модуль несёт собственный лазер, модулятор, фотодетектор и DSP. Позвоночник SuperPOD, расходящийся на тысячи GPU, в оптических терминах — это десятки тысяч подключаемых модулей, потребляющих десятки киловатт только на лазеры трансиверов.

С Rubin этот оптический слой сворачивается внутрь коммутационного ASIC. Quantum-X Photonics (InfiniBand) и Spectrum-X Photonics (Ethernet) заменяют подключаемые модули совместно упакованной оптикой: лазеры, модуляторы и фотодетекторы соединены с пакетом коммутатора через технологию TSMC COUPE. NVIDIA заявляет о примерно 4-кратном сокращении числа лазеров и 3,5-кратном снижении энергопотребления связи по сравнению с эквивалентом на подключаемых OSFP-модулях. Чиплетная логика, превратившая GPU в двухчиповый пакет и уложившая HBM рядом, теперь проявляется и на уровне сети: вертикальная интеграция вычислений, памяти и фотоники на одной подложке.

Недавно открытый NVLink Fusion сделал саму фабрику scale-up открытой: сторонние процессоры и XPU теперь могут присоединяться к доменам NVLink, позволяя гиперскейлерам строить полузаказные стойки вокруг интерконнекта NVIDIA, не разрабатывая собственную когерентную фабрику с нуля.

Программное обеспечение

CUDA — естественная модель программирования для массивно-параллельного процессора. Разработчик пишет ядро (код, исполняемый один раз на поток) и запускает его на тысячах потоков, организованных в блоки и варпы; программист решает, чем они делятся, когда синхронизируются и какую часть задачи каждый обрабатывает. Именно поэтому абстракция почти не изменилась за восемнадцать лет, и любое ядро CUDA, написанное с 2007 года, всё ещё скомпилируется и запустится на Blackwell.

Эта преемственность — и ров, и ограничение одновременно. Каждое новое поколение добавляет новое оборудование (Tensor Core, TMA, TMEM) поверх той же модели ядер и варпов, выставленной как интринсики в PTX и SASS: mma.sync, wgmma.mma_async и так далее. NVIDIA не может радикально переосмыслить SM, потому что слишком много кода от него зависит; взамен каждая инвестиция в софт CUDA накапливается из поколения в поколение.

Поверх PTX выстроен стек, создававшийся два десятилетия. cuBLAS и cuDNN для математики и примитивов глубокого обучения; CUTLASS, кодирующий десятилетия экспертизы по GEMM в шаблонном C++; TensorRT-LLM для paged attention, потокового батчинга и спекулятивного декодирования; привязки фреймворков через PyTorch, Triton и JAX.

FlashAttention, одна из важнейших алгоритмических переработок в современном ИИ, тайлит внимание, чтобы избежать материализации матрицы O(N²). Четыре поколения этого алгоритма (от FA1 до FA4) каждый раз вручную оптимизировались под последний кремний NVIDIA (FA3 под асинхронные конвейеры Hopper, FA4 под Blackwell), а порты на другое оборудование отставали на месяцы или годы.

Большая часть этого стека написана людьми, которым NVIDIA не платит. Ров — не сама CUDA; это два десятилетия сторонних ядер, библиотек и инструментов, а также миллионы разработчиков, изучивших этот API за это время.

NVIDIA также поставляет человеческую экспертизу вместе с кремнием. Компания встраивает десятки собственных инженеров внутрь передовых лабораторий и команд гиперскейлеров, пишущих ядра под каждую новую архитектуру модели и настраивающих их под каждое новое поколение кремния. Что бы лаборатория ни захотела обучить в следующем месяце, это, как правило, работает на NVIDIA гораздо быстрее, чем на других платформах. Поэтому отказ от NVIDIA — это не просто переписывание ядер и библиотек. Это переобучение ментальных моделей целого инженерного коллектива и потеря инженеров NVIDIA, которые сегодня сидят внутри здания.

Google TPU

TPU — это машина для умножения матриц. Философия в том, чтобы вместо программируемого чипа, способного запустить любую массивно-параллельную нагрузку, сосредоточиться на одном примитиве (плотное умножение матриц на большом систолическом массиве) и позволить компилятору XLA заранее спланировать каждый такт и каждый байт памяти. Никакого аппаратного планировщика, никакого кэша, никаких потоков/варпов. Каждое поколение наращивает под — тысячи чипов, соединённых через интерконнект ICI в одну когерентную машину. У TPU нет амбиций рендерить графику или выполнять научные симуляции; он существует, чтобы обучать и обслуживать нагрузки Google (поиск, перевод, рекомендации, Gemini) эффективнее на ватт, чем любая универсальная альтернатива.

Родословная

2015 — TPU v1: первый промышленный ASIC для глубокого обучения; только инференс INT8 через PCIe.
2017 — TPU v2: первый TPU, способный к обучению; переключил MXU с INT8 на BF16, установил связку из двух TensorCore + HBM.
2018 — TPU v3: первый TPU с жидкостным охлаждением; удвоил MXU и HBM относительно v2; поды на 1024 чипа.
2020 — TPU v4 (v4, v4i): первые реконфигурируемые оптические коммутаторы (Palomar); SparseCore; и BF16, и INT8; поды на 4096 чипа.
2023 — TPU v5 (v5e, v5p): v5e для эффективности, v5p для производительности; v5p имеет в 3,3× больше INT8 FLOPs и в 2,2× больше пропускной способности HBM, чем v4; поды на 8960 чипов.
2024 — Trillium (v6e): первый MXU 256×256; 4,7× пиковых FLOPS v5e при схожем энергопотреблении; обучил Gemini 2.0.
2025 — Ironwood (v7): создан для инференса моделей рассуждений; добавляет нативный FP8; суперпод на 9216 чипов при 42,5 ЭкзаFLOPS FP8.
2026 — TPU v8 (8t, 8i): 8t для обучения, 8i для инференса; добавляет нативный FP4; суперпод на 9600 чипов при 121 ЭкзаFLOPS FP4 (8t).

Архитектура

Чип TPU — это движок матричного умножения, обёрнутый ровно в столько кремния, сколько нужно для его питания. Единица вычислений — TensorCore: флагманские чипы начиная с v2 несут по два на пакет; чипы, настроенные на эффективность (v4i, v5e, v6e), несут по одному. Внутри каждого TensorCore — одинаковый пятикомпонентный набор: один или несколько MXU для матричной математики, VPU для поэлементной математики, скалярный блок, управляющий процессом, XLU для кросс-линейных редукций и присоединённый блок транспонирования/перестановки, плюс очереди накопителей, питающие и опустошающие MXU. Начиная с v4 каждый чип также несёт выделенные dataflow-движки SparseCore за пределами TensorCore (4 на чип на v4, v5p и Ironwood; 2 на чип на Trillium), явно выделенные для поглощения нагрузки поиска по эмбеддингам, для которой систолический массив был неподходящей формой. Каждый блок сидит на единой плоскости выпуска VLIW, управляемой Core Sequencer, заполняющим все восемь функциональных слотов 322-битного пакета каждый такт. Нет промаха кэша инструкций, нет варп-планировщика, нет движка внеочередного исполнения, нет предсказателя переходов: компилятор — это планировщик, а сэкономленная площадь кремния тратится на больше MAC.

TensorCore

MXU — это систолический массив. v1 поставлялся с одним 256×256 инференс-массивом INT8; v2 стал первым TPU, способным к обучению, и ввёл ячейки 128×128 с умножением BF16 и накоплением FP32 (INT8 вернулся в MXU начиная с v4 с эквивалентной пропускной способностью). Число ячеек на TensorCore росло оттуда: 1 MXU на v2 → 2 на v3 → 4 на v4/v5e/v5p. Trillium вернулся к 256×256 (65 536 ячеек умножения-накопления на массив за такт), и Ironwood, 8t и 8i все сохранили форму 256×256.

Чтобы вычислить C = A × B, значения матрицы B предзагружаются по одному весу на ячейку: dataflow с фиксированными весами (weight-stationary) — выбор, отличающий TPU от массивов с фиксированным выходом (output-stationary) в других местах. Активации входят с левого края, распространяются по одной колонне за такт, умножаются на резидентный вес в каждой ячейке, а частичные суммы текут вниз в очереди накопителей внизу. Как только данные вошли в массив, обращений к памяти больше не происходит: каждый вес переиспользуется для каждой прошедшей активации, каждая активация переиспользуется 128 (или 256) раз вдоль строки. Повторное использование данных зашито в кремний, а не арбитрируется кэшем. Доминирующая стоимость вычислений — не само умножение (несколько пикоджоулей), а чтение и запись памяти при энергозатратах в 100–1000 раз выше на обращение; систолический массив устраняет эту стоимость конструктивно. Компромисс — недозаполнение: матричное умножение 128×128 на массиве 256×256 тратит впустую 75% кремния, поэтому XLA тайлит, дополняет и планирует размерности до кратных 128 (или 256 на v6e+), а код модели пишется с учётом этих квантов.

VPU — второй по значимости вычислительный движок, но во многом более интересный микроархитектурный объект: каждый TPU — это 2D-векторная машина, а не 1D-SIMD машина. Регистровый файл VPU хранит 2D VREG. На v4/v5p форма — (8, 128): 128 линий в ширину, 8 сублиний в глубину, 32 (v4) или 64 (v5p) регистра на ядро, с 4 независимыми ALU с плавающей точкой на пару (линия, сублиния). Ось линий совпадает с шириной входа систолического массива, поэтому число линий, предположительно, расширилось до 256 вместе с MXU на Trillium и Ironwood; Google не публиковала размеры VPU после v5p. Ось сублиний позволяет VPU прогонять плитки через MXU со скоростью одно матричное умножение за X тактов (где X — размерность сублиний). Большая часть ускорения в современных программах TPU происходит от наложения VPU и MXU: квантование, layernorm, softmax, активация и сложение смещений — всё это выполняется на VPU в те же такты, когда MXU выполняет матричное умножение позади них. Кросс-линейные редукции (неудобный случай для любого 2D-векторного ISA) обрабатываются XLU: медленно, дорого и известное узкое место компилятора. Преобразования layout, не совпадающие с 2D-формой, поглощаются выделенным блоком транспонирования/перестановки, избавляя от обхода через память.

Скалярный блок — самый маленький блок и, пожалуй, самый значимый: однопоточный ALU с двойной выдачей для целых чисел, с 32 32-битными регистрами и 4 КиБ SMEM для управляющего состояния, в паре с Imem, хранящей программу. Это единственный блок, выполняющий выборку инструкций; каждый такт он вытягивает 322-битный пакет VLIW, исполняет свои два скалярных слота локально (адресная арифметика, счётчики циклов, переходы, проверки регистров синхронизации) и диспетчеризует оставшиеся шесть слотов остальной части чипа: 2 векторных ALU (VPU), 2 векторных загрузки/сохранения (DMA HBM↔VMEM), 2 матричных (push/pop очереди MXU). Синхронизация между блоками явная: флаги синхронизации отслеживают занятость конвейеров MXU и VPU, и компилятор вставляет проверки барьеров, а не аппаратура отслеживает зависимости. Скалярный блок — то, что делает остальную часть TensorCore похожей на dataflow с фиксированной функцией: каждый такт одно место решает, что происходит с восемью вещами, и нет динамического буфера переупорядочивания, способного отменить плохое решение.

Память

Иерархия памяти на чипе устроена по той же идее, что и вычислительная сторона: нет кэшей, каждый уровень управляется программно. За пределами чипа — HBM (16 ГБ на v2/v5e, 32 ГБ на v3/v4/v6e, 95 ГБ на v5p, 192 ГБ на Ironwood, 216–288 ГБ на поколении v8), а на чипе — вручную выстроенный уровень явно адресуемых скретчпадов. Ближе всего к вычислениям — VMEM, векторный скретчпад, питающий VPU и входные очереди MXU, размером 32 МиБ на v4, 128 МиБ на v5e, растянутый до 384 МиБ на настроенном под инференс v8i именно для того, чтобы вместить целый KV-кэш на чипе. Над ним — CMEM, введённый с v4 на 128 МиБ: более медленная и большая область SRAM между HBM и VMEM, поглощающая промежуточные результаты слитых операций. У скалярного блока — собственный SMEM (~10 МиБ для управляющего состояния на v4) и крошечный скалярный регистровый файл. Каждый тензор в программе закреплён за одним уровнем на этапе компиляции; проход назначения буферов XLA планирует DMA между уровнями так, чтобы данные прибывали как раз перед тактом, который их потребляет. Аппаратура не делает ни предвыборки, ни вытеснения, ни когерентности; когда компилятор всё делает правильно, массив никогда не простаивает; когда ошибается — запасного пути нет.

SparseCore

Блок за пределами TensorCore, нарушающий систолическую форму, — это SparseCore, введённый с v4. Рекомендательные и ранжирующие модели живут на поиске по эмбеддингам (миллиарды индексов в огромных таблицах), и паттерн доступа — противоположность плотному матричному умножению: нерегулярный, косвенный, «все-со-всеми». Систолический массив 256×256 — именно неправильная форма для этого. SparseCore — это dataflow-процессор с 16 вычислительными плитками и выделенными скретчпадами SPMEM, расположенный рядом с TensorCore и поглощающий примитивы scatter, gather и segmented-reduce плюс зависимый от данных трафик «все-со-всеми», порождаемый шардированными таблицами эмбеддингов. Это даёт ускорение в 5–7× на моделях с интенсивным использованием эмбеддингов при ~5% площади и энергопотребления кристалла. v4 поставлялся с 4 SparseCore на чип, v5p сохранил это число, Trillium снизил до 2, а Ironwood вернулся к 4 (по 2 на чиплет в своей двухчиповой компоновке). Инференс-чип v8i (Zebrafish) полностью убирает SparseCore и заменяет его CAE (Collectives Acceleration Engine) на чиплете ввода-вывода: другая проблема (коллективные редукции во время авторегрессивного декодирования), та же идея (выделить небольшой ускоритель из основного ядра для поглощения нагрузки, для которой систолический массив — неправильная форма).

Численные форматы

TPU v1 был только инференсом INT8; v2 переключился на BF16 как канонический формат обучения: тот же динамический диапазон, что у FP32, вдвое меньше памяти, без трюков масштабирования потерь. v4 вернул нативную поддержку INT8. Ironwood затем добавил нативную поддержку FP8 (и E4M3, и E5M2) для примерно 2-кратной пропускной способности BF16 в той же площади. v8 добавляет нативный FP4 плюс блочно-масштабное умножение прямо внутри MXU, что устраняет накладные расходы деквантования на VPU, которые ещё платил Ironwood. Стохастическое округление аппаратно поддерживается на каждом современном TensorCore: решения об округлении принимаются младшими битами мантиссы, действующими как вероятность, что сохраняет ожидаемое значение накоплений низкой точности на протяжении длинных прогонов обучения и является одной из тех мелких деталей, что позволяют BF16/FP8 закрыть разрыв в точности с FP32.

На границе чипа находятся сами порты ICI (4 порта на чипах с 2D-тором v2/v3/v5e/v6e, 6 на флагманах с 3D-тором v4/v5p/v7/8t), и сетевая карта DCN для scale-out. С точки зрения чипа порты ICI выглядят просто как ещё один набор движков DMA, на который может нацелиться Core Sequencer внутри пакета VLIW: отправка удалённого тензора — тот же класс инструкции, что и передача VMEM-в-HBM, и компилятор рассматривает коллективные операции как часть того же общего расписания, что строится для вычислений и локальной памяти.

Ставки

• Ставка 1: Систолический массив. В нагрузке доминирует матричное умножение, поэтому кремний тратится на систолический массив.
• Ставка 2: Программные скретчпады. Вычисления дёшевы, а память дорога, поэтому данные переиспользуются прямо в проводах массива, а кэши заменяются программно управляемыми скретчпадами.
• Ставка 3: Планирование компилятором. Нагрузка статически предсказуема, поэтому планирование переносится в компилятор: выдача VLIW, никакой спекуляции, никакого внеочередного исполнения, никакого динамического планировщика.
• Ставка 4: Кремний только под MAC. Энергопотребление важнее пика, поэтому удаляется каждый транзистор, не выполняющий умножение-сложение: каждый тег кэша, каждый предсказатель переходов, каждый буфер переупорядочивания.
• Ставка 5: Выделенные движки вне массива. Плотный массив матричного умножения — неправильная форма для некоторых реальных нагрузок (эмбеддинги, коллективные операции), поэтому выделяются небольшие специализированные движки (SparseCore, CAE), а не искажается основное ядро под них.

Масштабирование

История scale-up у TPU — обратная по отношению к NVIDIA. Там, где NVLink + NVSwitch делают HBM любого другого GPU похожим на локальную память (аппаратно управляемое когерентное адресное пространство), ICI Google — это передача сообщений. Нет семантики удалённой загрузки, нет когерентности кэша, нет кроссбара. Каждая многочиповая операция — явная коллективная операция, компилируемая XLA. Домен scale-up связан не коммутационной фабрикой, а тором (чипы, подключённые напрямую к соседям, с оборачиванием краёв) и сшит на границе стойки оптическими коммутаторами.

Scale-up: чипы соединяются напрямую друг с другом в 2D- или 3D-торе через ICI. XLA генерирует SPMD-коллективы, тесно хореографирующие тысячи TPU как одну программу. Нет когерентности, но огромная пропускная способность разреза при низкой задержке.

Scale-out: поды соединяются в сеть через фабрику дата-центра: гораздо больше чипов, чем помещается в один домен ICI, но при меньшей пропускной способности на чип. Сегодня: Virgo обрабатывает трафик TPU восток-запад (v8t+), Jupiter — север-юг. Multislice + Pathways оркестрируют SPMD между подами.

Scale-up

Связи ICI выходят прямо из кристалла TPU: высокоскоростные последовательные линии, прямое медное подключение внутри куба на 64 чипа (расположение 4×4×4 в одной стойке с жидкостным охлаждением), оптика между кубами. Совокупная пропускная способность ICI на чип выросла с ~250 ГБ/с на v2 до 1,2 ТБ/с в обоих направлениях на Ironwood, и вдвое больше на v8t. Топология чередуется по поколениям: 2D-тор на чипах, настроенных на эффективность (v2, v3, v5e, v6e), 3D-тор на флагманах (v4, v5p, v7, v8t).

Элемент без аналога у NVIDIA — Palomar OCS: оптический коммутатор на 3D-MEMS, стоящий между кубами. Крошечные зеркала физически поворачиваются, чтобы отобразить любое входное волокно на любой выходной порт. Суперпод v4 использует 48 коммутаторов Palomar, чтобы соединить 64 куба (4096 чипов) в один 3D-тор; v5p и Ironwood масштабируют ту же схему дальше. Реконфигурация занимает миллисекунды, не наносекунды, но это нормально, потому что OCS — коммутатор каналов: топология выбирается при старте задания, работает неделю, затем перенастраивается под следующую нагрузку. Три проблемы сворачиваются в один компонент: реконфигурация топологии под нагрузку (скрученные торы дают до 70% лучшую пропускную способность разреза), нарезка подподов по требованию, и отказоустойчивость (когда чип выходит из строя, OCS оптически подменяет запасной куб, и прогон продолжается без потери домена ICI).

Это делает суперпод единицей масштабирования scale-up: эквивалентной по роли NVL72 от NVIDIA, но на два порядка больше. v4 насчитывал 4096 чипов; v5p — 8960; Ironwood (TPU v7) — 9216 чипов, организованных как 144 куба по 64, представляя 1,77 ПБ HBM (~68 ПБ/с) и 42,5 ЭкзаFLOPS FP8 как один когерентный домен ICI.

TPU 8t (Sunfish) растягивает это до 9600 чипов, 2 ПБ HBM (~62 ПБ/с) и 121 ЭкзаFLOPS FP4. TPU 8i (Zebrafish) имеет 1024 чипа, ~295 ТБ HBM (8,8 ПБ/с) и ~10 ЭкзаFLOPS FP4. 8i заменяет тор новой иерархической высокорадиксной топологией под названием Boardfly (кольцо из 4 чипов → группа из 8 плат → до 36 групп, связанных через OCS), сокращая вдвое задержку «все-со-всеми». Это спроектировано для инференса MoE. 3D-тор превосходен, когда коллективные операции — соседские (ring all-reduce использует каждую связь каждый такт), но маршрутизация экспертов MoE — противоположный паттерн, «все-со-всеми»: каждый чип отправляет уникальные фрагменты каждому другому, а задержка кругового пути ограничена самой длинной парой прыжков. 3D-тор на 1024 чипа имеет диаметр в 16 прыжков; иерархия Boardfly кольцо → группа → OCS сжимает это до 7.

Scale-out

Вплоть до TPU v7 scale-out работал через единую фабрику: Jupiter, полностью оптическую на позвоночнике с 2022 года через Apollo OCS, того же семейства 3D-MEMS, что и Palomar, масштабированную на всё здание. Google использует один и тот же примитив (оптическая коммутация каналов) на каждом слое от стойки до позвоночника дата-центра; это архитектурная подпись, которой ни у кого больше нет. Jupiter сегодня несёт 13 Пбит/с пропускной способности разреза на здание.

С TPU 8t scale-out разделился на две фабрики. Трафик TPU-TPU восток-запад переехал на Virgo, выделенную фабрику ускорителей; Jupiter сохранил роль север-юг: доступ к хранилищу, общие вычисления, масштабирование между площадками. Virgo — это плоская, двухуровневая, неблокирующая топология, построенная на высокорадиксных коммутаторах: любой TPU не более чем в двух коммутаторах от любого другого. Один кластер Virgo связывает более 134 000 TPU 8t при 47 Пбит/с пропускной способности разреза (в 4× выше пропускной способности на чип и на 40% ниже ненагруженной задержки, чем у предыдущего поколения DCN), с многоплоскостной изоляцией отказов и субмиллисекундной телеметрией, позволяющей планировщику убивать отстающие задачи до того, как они сорвут шаг обучения. Архитектурный выигрыш в том, что каждый слой теперь может эволюционировать независимо: scale-up, scale-out восток-запад и фронтенд могут итерироваться в разном темпе, не переподключая друг друга.

Пропускная способность scale-out на чип составляет порядка 100 Гбит/с на Ironwood и в 4× больше на v8t, но всё ещё на два порядка меньше, чем ICI на чип. Этот разрыв в пропускной способности диктует партиционирование: тензорный параллелизм и маршрутизация экспертов MoE остаются внутри ICI; параллелизм данных и конвейерный параллелизм пересекают фабрику scale-out.

Фреймворк Multislice от Google, встроенный в XLA, позволяет одной SPMD-программе охватывать несколько срезов в разных подах; компилятор генерирует иерархические коллективные операции (ring all-reduce внутри каждого среза, редукция более высокого уровня между ними). Структура — именно тот трюк для скрытия разрыва пропускной способности ICI/DCN: как можно больше работы остаётся внутри среза через быстрый ICI, оставляя только межсрезовый остаток платить за медленную фабрику.

Поверх этого стоит Pathways. Там, где планировщики в стиле NCCL + Slurm + Megatron управляют SPMD от множества контроллеров, Pathways управляет всем заданием от одного клиента и виртуализирует несколько «островов» (подов со своими доменами ICI), соединённых через DCN. Он выполняет групповое планирование, эластичное обучение (когда срез выходит из строя, OCS перестраивает топологию, и Pathways возобновляет работу с последнего чекпоинта на новой форме) и межрегиональную оркестрацию. Gemini Ultra стала первой передовой моделью, обученной на нескольких дата-центрах; Pathways сшивает их в одно синхронное SPMD-задание.

Философия: компилятор — это планировщик, тор — это топология, а оптический коммутатор — универсальная реконфигурируемая подложка на каждом слое от стойки до дата-центра.

Программное обеспечение

Стек TPU управляется компилятором там, где CUDA управляется ядрами. На GPU разработчик пишет ядро, а фреймворк связывает ядра вместе; работа компилятора в основном локальна. На TPU разработчик пишет числовую программу на JAX, и XLA отвечает за всё, что ниже: какие операции сливаются, где живёт каждый тензор, как он раскладывается по 2D-векторным регистрам, когда выдаются DMA из HBM в VMEM, как планируются 322-битные пакеты VLIW, как программа шардируется по тысячам чипов. Нет аппаратного запасного пути: ни варп-планировщика, ни кэша, ни движка внеочередного исполнения, чтобы сгладить плохое расписание. Компилятор — это система. Компромисс — центральный для этой архитектуры: XLA приближается к теоретическому пределу без ручной настройки, но закрыть оставшийся разрыв сложнее.

Путь компиляции: JAX → JAXpr → StableHLO → HLO → LLO → пакеты VLIW. JAX трассирует функцию Python в типизированное функциональное IR (JAXpr) под jit, снижает её до StableHLO (стандартизированный OpenXLA, версионированный набор операций из ~100 статически типизированных примитивов, который теперь генерируют все фронтенды), который XLA принимает как HLO и прогоняет через свой конвейер проходов: слияние операций (сворачивание поэлементных + редукционных + матричных операций в одно ядро, чтобы промежуточные результаты никогда не попадали в HBM), назначение раскладки (решение о 2D-тайлинге каждого тензора, чтобы он поступал в MXU без транспонирования: существенно сложнее, чем на 1D-SIMD машинах, поскольку и регистры, и входы систолического массива двумерны), назначение буферов (каждый тензор закрепляется за VMEM, CMEM или HBM с предвычисленными окнами перекрытия), SPMD-партиционирование и, наконец, VLIW-планировщик, заполняющий все восемь слотов каждого пакета. HLO снижается до LLO (Low-Level Optimizer), TPU-специфичного IR, а LLO генерирует финальный поток VLIW. Хорошо скомпилированная программа накладывает друг на друга систолическое исполнение MXU, поэлементную математику VPU и DMA HBM↔VMEM в одном и том же пакете каждый такт.

Многочиповое исполнение — это SPMD: одна программа, шардированные данные, иерархические коллективные операции, генерируемые GSPMD (сейчас заменяемым Shardy, MLIR-нативным преемником, становящимся стандартом по умолчанию в начале 2026 года). Пользователь выражает шардирование декларативно с помощью аннотаций Mesh + PartitionSpec на нескольких ключевых тензорах; компилятор распространяет шардирование по остальному графу и вставляет all-reduce, all-gather и reduce-scatter там, где меняется раскладка. Когда компилятор выбирает неправильную коллективную операцию, shard_map опускает пользователя до ручного SPMD (код по устройствам с явными локальными формами и явными коллективными операциями), компонуемого внутри jit, так что одно ядро может быть вручную партиционировано без отказа от автопартиционирования везде ещё. Это обратное от идиомы PyTorch: FSDP и DeepSpeed оборачивают модель в рантайм, выдающий коллективные операции на границах модулей; GSPMD/Shardy партиционирует весь граф как задачу компилятора.

Pallas — это лазейка: язык написания ядер для JAX, в целом эквивалент Triton для TPU. Ядра Pallas пишутся на Python во вкусе JAX, снижаются через Mosaic (бэкенд TPU на основе MLIR) до LLO и встраиваются обратно в HLO как пользовательская операция. Он существует, потому что XLA не всегда может синтезировать оптимум для новых вариантов внимания, слитой диспетчеризации MoE или чего-либо, требующего ручного тайлинга VMEM и планирования DMA: оптимизация класса FlashAttention, где выигрыш в расписании, а не в алгебре. Pallas:Mosaic-GPU нацелен на H100/Blackwell с тем же фронтендом, так что автор ядра может написать один раз и снизить на любую из подложек. Уровень библиотек над этим полностью нативен для JAX: Flax NNX для модулей, Optax для оптимизаторов, Orbax для асинхронного распределённого чекпоинтинга, Grain для входных конвейеров, Tunix для post-training/RL, Qwix для квантования. Эталонные стеки обучения Google (MaxText для LLM, включая MoE класса DeepSeek-V3, и MaxDiffusion для Flux, Wan 2.1) находятся наверху, на чистом JAX; Pathways — под ними, доступен пользователю как pathwaysutils, так что единый клиент Python может управлять заданием на тысячах чипов и нескольких подах-островах, не отказываясь от модели программирования JAX.

Путь PyTorch реален, но второстепенен. torch_xla использует механизм LazyTensor: каждая операция PyTorch записывается в граф HLO, компилируемый на следующем барьере, с кэшированием скомпилированного артефакта по хешу формы графа. PyTorch/XLA 2.x добавил аннотации шардирования в стиле GSPMD, интеграцию torch.compile через бэкенд XLA, мост к JAX и (PyTorch/XLA 2.7) сборки на ABI C++11 с существенно более быстрой трассировкой. Разрыв с JAX реален (примитивы JAX отображаются на StableHLO чище, а сложные стратегии параллелизма покрыты лучше), поэтому vLLM TPU (на базе плагина tpu-inference, анонсированного на Cloud Next 2025) снижает каждую модель, определённую хоть в JAX, хоть в PyTorch, через унифицированный путь JAX→XLA. TorchTPU, анонсированный в апреле 2026 года, — ответ Google: нативный опыт PyTorch с eager-режимом, torch.distributed и torch.compile через XLA, призванный заменить torch_xla.

По сравнению с CUDA, экосистема TPU централизована, а не разрослась. Почти всё ниже фреймворка (XLA, JAX, Flax, Optax, Pallas, MaxText, Pathways, Shardy, Mosaic) открыто самим Google и эволюционирует в ногу с кремнием. Сторонних ядер значительно меньше, чем накопила CUDA за десятилетия; ров тоньше там, где нагрузка выглядит необычно, и глубже там, где она похожа на Gemini. Недавняя формулировка «кодизайна ИИ-стека» для Ironwood (v7) — явное обрамление этой идеи: чип, фабрика ICI, OCS, XLA, Pathways, Pallas, MaxText, vLLM и Pathways выпущены вместе как один продукт, а v8t/v8i продолжают ту же модель под единым путём снижения tpu-inference. Triton и torch.compile сужают разрыв на стороне NVIDIA (управление ядрами и управление компилятором сходятся), но философские полюса всё ещё реальны: на TPU компилятор — единственный интерфейс, который имеет значение; на GPU компилятор — лишь один из нескольких.

AMD GPU

GPU AMD Instinct построены на другой ставке, чем NVIDIA: там, где NVIDIA каждое поколение расширяет то, что может делать каждый SM, AMD с момента GCN (2012) держит Compute Unit консервативным и реинвестирует в пакет: соответствует или превосходит современный флагман NVIDIA по ёмкости HBM каждое поколение начиная с 2021 года; первый 3D-стековый датацентровый GPU (CDNA 3); первый когерентный CPU+GPU APU (MI300A); и открытая экосистема (ROCm, HIP, OCP MX, UALink).

Родословная

2018 — Vega 20 (MI50, MI60): первый 7-нм GPU; векторная пропускная способность FP64 1:2. Последний Instinct семейства GCN перед CDNA/RDNA.
2020 — CDNA (MI100): первые матричные ядра MFMA; графическая фиксированная логика полностью убрана. Нативный BF16.
2021 — CDNA 2 (MI210, MI250, MI250X): первый MCM Instinct через двухчиповый пакет GCD; полноскоростное матричное FP64.
2023 — CDNA 3 (MI300A, MI300X): первый 3D-стековый чиплетный GPU: XCD гибридно соединены с IOD через TSV; FP8; Infinity Cache; когерентный CPU+GPU APU на MI300A; питал суперкомпьютер El Capitan.
2024 — обновление CDNA 3 (MI325X): те же вычисления, обновлённая HBM3E: 256 ГБ при 6,0 ТБ/с.
2025 — CDNA 4 (MI350X, MI355X): нативные FP4/FP6 с микромасштабированием OCP MX; FP64 на CU урезан примерно вдвое; первое поколение, склонившееся в сторону плотности ИИ над HPC.
2026 — CDNA Next (MI430X, MI440X, MI455X): HBM4; стойка Helios (флагманский MI455X на 72 GPU через UALoE при запуске, нативный UALink с 2027 года): первый ответ AMD на NVL72.

Архитектура

Терминология AMD и NVIDIA: Compute Unit (CU) ↔ Streaming Multiprocessor (SM); SIMD ↔ SM Sub-Partition; SIMD Lane ↔ CUDA Core (FP32 ALU); Wavefront (wave64) ↔ Warp (warp32); Matrix Core ↔ Tensor Core; MFMA ↔ mma.sync / wgmma / tcgen05.mma; VGPR/SGPR ↔ регистровый файл; LDS (Local Data Share) ↔ SMEM (Shared Memory); Infinity Fabric ↔ NVLink.

Там, где архитектурные амбиции NVIDIA живут внутри каждого SM (новые матричные примитивы, новая асинхронная машинерия, новые хранилища операндов каждое поколение), амбиции AMD живут между CU, в том, сколько их можно связать в единый когерентный пакет. Сам CU консервативен: четыре 16-линейных SIMD, один общий скалярный блок, 64 КБ Local Data Share...