История о том, как штука для рисования треугольников стала самым важным чипом на планете

Андрей Филатов · 2026 · ~20 минут

Я работаю с GPU каждый день — запускаю модели, считаю инференс, жду пока обучится очередной эксперимент. Но долгое время для меня GPU был чёрным ящиком: закинул модель, нажал кнопку, ждёшь. А мне хочется понимать, что будет происходить дальше с миром AI — и я понял, что без понимания железа это невозможно. Я уже пытался предсказать, когда видеогенерация станет реалтаймовой — получилось неплохо, и мне захотелось погрузиться вглубь: а что вообще стоит за GPU?

А ведь решения в мире AI сейчас принимают те же люди, которые проектировали GPU. Если разобраться, почему они принимали именно такие решения на каждом этапе — можно научиться самому прогнозировать, куда всё движется. Ну и заодно построить ментальную модель GPU, которой мне так не хватало.

(а теперь SEO оптимизация) NVIDIA стоит дороже всех компаний мира. H100 продаётся за $30 000. И всё это — из-за процессора, который 25 лет назад рисовал треугольники в компьютерных играх. Как мы сюда попали — разбираемся с нуля. #NVIDIA #GPU #AI

Глава I

Как компьютер рисует картинку

От текста к полигонам

Для меня не секрет, что GPU были созданы для игр — первое что я узнал про видеокарты, это как раз NVIDIA 9800 GT, которые позволяли мне запускать любимые Assassin's Creed и Call of Duty. Но как понадобилось отдельное железо для отрисовки картинок?

Первые компьютеры общались с человеком только текстом — зелёные буквы на чёрном экране. Zork (1977) — один из первых хитов: ты печатаешь open mailbox и получаешь текстовое описание того, что внутри. Такой DnD-симулятор квестов. Я их не люблю, и видимо не я один — потому что довольно быстро появились первые графические игры.

West of House You are standing in an open field west of a white house, with a boarded front door. There is a small mailbox here. > open mailbox_

1977 · текст · 0 пикселей графики

В начале 80-х появились спрайты — заранее нарисованные картинки, которые просто двигаются по экрану. Super Mario Bros (1985), Sonic (1991) — целые яркие миры, собранные из плоских кирпичиков. Это намного круче текстовых игр, но и плата — повышенные вычислительные аппетиты. Впрочем, со спрайтами терпимо: компьютер просто копирует готовые блоки в нужные позиции, 60 раз в секунду. Никаких вычислений «как это должно выглядеть» — только «куда положить картинку».

1985 · спрайты двигаются по экрану — копирование готовых блоков

Но плоские картинки не передают глубину. Wolfenstein 3D (1992) и Doom (1993) красиво это обошли — мир хранился как плоская карта, а движок рисовал столбцы стен по расстоянию. Гениальный трюк, но в Doom нельзя даже посмотреть вверх — третьего измерения буквально нет. Настоящий прорыв — Quake (1996): полностью трёхмерный мир из полигонов. Свободная камера, текстуры, освещение — и всё собрано из сотен треугольников, каждый из которых нужно трансформировать, раскрасить и отрисовать. 30 раз в секунду.

HEALTH 85 AMMO 24

1996 · 3D полигоны · трансформация + освещение каждого пикселя

Один кадр Quake

Давайте заглянем под капот. Quake на экране: игрок в коридоре, стены с текстурами, факел, в дверном проёме — враг. Камера может повернуться куда угодно. Всё это нужно превратить в картинку — 640 на 480 пикселей, 30 раз в секунду. У компьютера 33 миллисекунды на один кадр.

Что происходит за эти 33 миллисекунды? Берём все треугольники сцены, пересчитываем координаты из 3D в плоские экранные. Определяем, какие пиксели попадают внутрь каждого треугольника. И для каждого из 307 200 пикселей считаем цвет: читаем текстуру, считаем освещение, проверяем глубину. Это дофига математических операций — возьми одно число, умножь на второе, сложи с третьим, и так для каждого пикселя.

Анимация — как рисуется один кадр

640×480 · 30 fps · ~300M операций/сек. CPU рисует каждый пиксель последовательно.

Посчитаем. 640 × 480 = 307 200 пикселей. На каждый — 20–50 операций. На один кадр выходит 10–15 миллионов операций. Умножаем на 30 fps — 300–450 миллионов операций каждую секунду.

Pentium 200 МГц — топовый процессор 1996 года — выдавал около 200 миллионов инструкций в секунду. Уже на грани. Но в реальности всё ещё хуже: каждый пиксель лезет в текстуру — изображение, натянутое на треугольник. Адрес непредсказуем, в кэше данных почти наверняка нет. Кэш-промах — процессор стоит и ждёт 50–100 наносекунд, пока память ответит. Десятки потерянных тактов. На каждый пиксель. Реальная производительность падает в 2–3 раза. Прям мега грустно. CPU явно не справляется — нужно отдельное железо для графики.

Зашито в кремний

Рынок это понял быстро. В середине 90-х появились первые графические ускорители: 3dfx Voodoo, ATI Rage, S3 ViRGE. Идея простая — вынести графические операции с CPU на отдельный чип. Но транзисторов мало, жили бедно, поэтому операции напрямую впаивали в кремний. Voodoo умел растеризировать и накладывать текстуры — и только это. Трансформация вершин? Всё ещё на CPU. ATI Rage умел и то, и другое, но только с одним источником света. S3 ViRGE вообще оказался медленнее программного рендеринга. Каждый чип — фиксированный набор функций. Хочешь эффект, которого в чипе нет — мечтай.

Для разработчика — кошмар. Каждая карта говорит на своём языке: свои регистры, свои команды. Написал игру под 3dfx — на ATI она не работает. Вышла новая карта — переписывай всё заново.

OpenGL (1992) и DirectX (1995, эхх, сколько его приходилось ставить) частично решили проблему совместимости: разработчик вызывает стандартную функцию вроде DrawTriangle, а драйвер переводит её в команды конкретного чипа. Один код → любое железо.

Один вызов → разные команды для разного железа DrawTriangle(A, B, C) DirectX / OpenGL — единый API 3dfx Voodoo REG_TRI_V0, SUBMIT → свой протокол ATI Rage CMD_SETUP, VERT_DATA → другой протокол S3 Virge S3_TRI_BEGIN, XY → третий протокол Драйвер переводит один вызов API в команды конкретного чипа.

Но проблема fixed-function никуда не делась. Каждый ускоритель умел только то, что в него зашили. Хочешь новый эффект — жди следующее поколение чипов. И ни один из них не покрывал весь конвейер целиком: часть работы всё равно падала на CPU.

Нужен был другой подход.

Глава II

Проектируем GPU с нуля

Ок, хотим спроектировать ускоритель — надо сначала понять, что именно ускоряем. Рендеринг кадра — это конвейер. Сначала каждую вершину треугольника пересчитываем из 3D в экранные координаты — матричное умножение 4×4. Потом определяем, какие пиксели попадают внутрь треугольника — растеризация. И для каждого пикселя считаем цвет: текстура, освещение, глубина.

И вот ключевое наблюдение, из которого родился GPU: на каждом этапе — тысячи операций, и все они независимы. Один пиксель не ждёт результата другого. Одна вершина не зависит от соседней. Формула одинаковая, данные разные, всё независимо. Когда это слово встречается — стоит подумать о параллелизме. Осталось спроектировать процессор, который умеет именно это.

Почему не CPU

Внутри CPU — нажми на блок

Предсказание ветвлений Перестановка инструкций Кэш L1 / L2 / L3 ALU ~80% площади — управление · ~20% — вычисления
← нажми на блок чтобы узнать что он делает и почему для графики это лишнее

А как с параллелизмом у CPU? Спойлер — всё плохо. Большая часть транзисторов CPU — это не вычисления, а управление. Предсказание ветвлений, переупорядочивание инструкций, кэши. Это делает CPU невероятно умным для сложных последовательных задач — как гений-одиночка. Но это как заставить профессора математики считать на калькуляторе. Он справится, но это не его сильная сторона. Нам не нужно предсказывать ветвления — формула для каждого пикселя одинаковая. Не нужно переставлять инструкции — все потоки делают одно и то же. Не нужны огромные кэши — данные всё равно непредсказуемы. Чисто вычислительная часть CPU занимает ~20% площади — ALU, арифметико-логическое устройство. Единственный блок, который реально считает.

Ну вы как умный человек подумаете — мало вычислений? Давайте сделаем больше вычислений. Забьём весь чип ALU и всё будет по кайфу. И будете правы.

Шаг 1: Множим ALU

Берём CPU, выкидываем всё «умное», оставляем только ALU. На освободившуюся площадь можно уместить не одно ALU с обвязкой, а десятки голых. Каждый проще, каждый медленнее — но их много, и они работают одновременно. Не один умный процессор — целая армия простых.

ветвления перестановка кэши ALU ~20% вычисления ~95% вычисления

Шаг 2: Одна команда на всех

Окей, у нас десятки ALU на одном кристалле. Но сразу столкнёмся с проблемой — блин, а как ими управлять? Если каждому дать свой декодер команд, свой счётчик инструкций, свою управляющую логику — мы опять потратим кучу транзисторов на управление. Тут погорячились — вернулись к тому, от чего убежали.

Но вспомним нашу задачу: все пиксели считаются одной и той же формулой. Цвет пикселя 0 — как цвет пикселя 31, просто входные данные разные. Значит, можно проще: один декодер команд на группу ALU. Читает инструкцию один раз — все ALU в группе выполняют её одновременно, каждый со своими данными. Это SIMD — Single Instruction, Multiple Data. Все делают одно и то же.

В GPU NVIDIA такая группа — warp, 32 потока. Почему 32? Компромисс: больше группа — дешевле управление, но больше проблем когда потоки хотят делать разные вещи. Итого: один декодер, 32 ALU, одна инструкция — все считают параллельно. Проблему управления решили.

Но есть нюанс. Что если в коде if/else? При рендеринге Quake проверяем, попадает ли пиксель в тень. Часть из 32 пикселей в тени, часть на свету — а декодер-то один, он не может выполнять разные инструкции одновременно.

Решение элегантное: GPU выполняет обе ветки на всех 32 ALU, но с маскированием. Прогоняет ветку «в тени» — все 32 считают, но результаты пикселей на свету отбрасываются. Потом «на свету» — то же самое наоборот. Два такта вместо одного, часть работы впустую. Это warp divergence. Для графики терпимо — соседние пиксели обычно в одной ветке.

Warp divergence: if (в тени?) Такт 1: ветка «в тени» 12 впустую Такт 2: ветка «на свету» 20 впустую 20 в тени + 12 на свету → 2 такта вместо 1 Потеря: 37.5% производительности на этом if

Шаг 3: Прячем задержки

У нас массив ALU, объединённых в warps. Они считают быстро. Но есть узкое место: память. Когда шейдер обращается к текстуре — а это происходит для каждого пикселя — данные нужно прочитать из видеопамяти. А видеопамять медленная по сравнению с процессором: запрос и ответ занимают 200–400 тактов. Всё это время 32 ALU в warp стоят и ждут. Ничего не считают. Чистый простой. (Кстати, проблемы с HBM у H100 взялись не из ниоткуда — перегонять данные с памяти на чип было дорого ещё тогда.)

CPU решает это кэшами — держит данные поближе. Но мы же выкинули большие кэши, когда проектировали GPU. Как быть?

Решение красивое. Вместо того чтобы прятать задержку — GPU заполняет её полезной работой. На одном процессорном блоке (SM — Streaming Multiprocessor) GPU держит в регистрах состояние не одного warp, а десятков. Переключение — за один такт, без накладных расходов. Warp A запросил текстуру и ждёт? GPU переключается на warp B. Потом C. Потом D. К моменту, когда данные для A приходят из памяти — очередь как раз до него дошла. Простои, чтобы от них избавиться — GPU просто делает другую работу, пока ждёт.

Latency hiding: переключение warps Без: warp A ждёт: 200 тактов warp A ждёт: 200 тактов ~25% С: A B C D A B C D A ~100% 4 warps по 32 потока. Пока A ждёт — B, C, D считают. В реальности: 32–64 warps на SM. GPU никогда не простаивает. Формула: нужно warps ≥ (latency / compute_time). 200 тактов / 60 тактов ≈ 4 warps минимум

Чем больше warps в полёте — тем лучше скрывается задержка. Задержка 200 тактов, один warp считается 60 — значит, нужно минимум 4 в очереди. В реальности GPU держит 32–64 warps на одном SM и почти никогда не простаивает. Важнейший принцип — запоминаем.

CPU прячет задержку кэшами — хранит данные рядом. GPU — переключением потоков: пока один ждёт, другой считает.

Подробнее об архитектуре GPU, warps и latency hiding — в бесплатном курсе Cornell Virtual Workshop: Understanding GPU Architecture.

Мы только что создали GPU

Оглянемся. За три шага мы спроектировали новый тип процессора. Выкинули из CPU всё лишнее для графики — предсказание ветвлений, перестановку инструкций, огромные кэши. Забили площадь десятками простых ALU. Объединили их в группы с общим декодером — SIMD. Проблему медленной памяти решили переключением между потоками. Это и есть GPU.

Именно так рассуждали инженеры NVIDIA в конце 90-х. Складываем всё вместе — и в 1999 году появляется GeForce 256. Они сами назвали его «первым в мире GPU». 17 миллионов транзисторов, 120 МГц, 10 миллионов полигонов в секунду. Главное — аппаратный Transform & Lighting: трансформации вершин и расчёт освещения впервые целиком переехали с CPU на видеокарту. Один чип — весь конвейер рендеринга.

NVIDIA GeForce 256 — первый в мире GPU, 1999

NVIDIA GeForce 256 (1999) · 17M транзисторов · 120 МГц · 10M полигонов/с · фото: Wikimedia Commons, public domain

Глава III

От шейдеров к CUDA

Шейдеры: от зашитого к программируемому

Мы нанесли удар по скорости вычислений — больше транзисторов, больше ALU, красота. Но была другая проблема: всё зашито в кремний. Хочешь другой эффект? Жди новый чип.

Почему вообще зашивали? Бюджет транзисторов — 17 миллионов, каждый на счету. Вот конкретный пример. Берём один пиксель на стене в Quake рядом с факелом. Чтобы вычислить его цвет, чип проводит пиксель через цепочку зашитых операций:

текстура серый кирпич N · L угол к факелу → где светлее (R·V)ⁿ блик — яркое пятно от света × Kd × цвет кирпича → коричневый + ambient + фоновый свет → тени не чёрные готовый пиксель pixel = texture × (N·L) + Ks×(R·V)ⁿ + ambient 🔒 зашито в кремний — формулы и порядок не поменять ~200 транзисторов, 1 такт · программируемый: ~50 000, 3-4 такта

Проблема конкретная. Та формула, которую мы только что видели — она считает освещение стены. Свет падает, отражается, получается блик. Для кирпичной стены в Quake — отлично. Но вода отражает и преломляет свет совсем иначе. Кожа рассеивает свет под поверхностью. Металл — резкий блик, а мех — мягкий. Для каждого материала нужна своя формула. А в чипе зашита одна. И зашить все варианты в кремний невозможно — их слишком много.

стена блик на поверхности N·L + (R·V)ⁿ зашита в чип вода отражение + преломление reflect + refract(η) нет в чипе кожа свет проходит внутрь scatter(L, depth) нет в чипе в чипе зашита одна формула — а нужно три разных

Но рост числа транзисторов помог: 17М (1999) → 57М (2001) → 107М (2002). Появился бюджет для эксперимента.

И таким образом появились шейдеры — небольшие программы, которые пишет разработчик, а GPU исполняет для каждой вершины или пикселя. Разница принципиальная: в fixed-function транзисторы соединены проводами в определённом порядке — чип всегда выполняет ровно ту цепочку, которую впаяли на заводе. В шейдерной модели на чипе стоит универсальный ALU + память для инструкций. Перед рендерингом загружаешь туда свою программу — любую. Тот же ALU может посчитать и стену, и воду, и кожу — просто загрузи другую формулу. Заскейлились:

SM 1.0 2001 · GeForce 3 if/else до 128 инструкций 2-3× медленнее fixed-function SM 2.0 2002 · Radeon 9700 if / else 256 инструкций Реальная логика SM 3.0 2004 · GeForce 6 if / else / for / while до 65 000 инструкций Полноценный процессор 2001 2002 2004

Рост программируемости: от 128 инструкций без ветвлений до полноценного процессора за 3 года.

Shader Model 1.0 (2001, GeForce 3) — до 128 инструкций, нет ветвлений. По сути тот же fixed-function, только можно переставить шаги. Программируемый блок в 2-3 раза медленнее зашитого — гибкость стоит дорого.

Shader Model 2.0 (2002, Radeon 9700) — 256 инструкций, впервые if/else. Уже можно писать реальную логику, overhead меньше.

Shader Model 3.0 (2004, GeForce 6) — до 65 000 инструкций, циклы, динамическое ветвление. Полноценный процессор. Overhead стал настолько маленьким, что драйверы начали эмулировать fixed-function через шейдеры — зашитые блоки буквально потеряли смысл.

Нафиг тогда в железо операции зашивать, если программно почти так же быстро и ещё миллион вещей умеет? Гибкость победила.

GPU из «калькулятора с кнопками» превратился в «программируемый калькулятор».

Следующий логичный шаг — все блоки на чипе стали одинаковыми. 128 универсальных блоков, планировщик сам раскидывает задачи. Утилизация близка к 100%.

И тут инженеры NVIDIA осознали кое-что поинтереснее. Если все 128 блоков умеют любую математику — это уже не видеокарта. Это массивно-параллельный вычислитель общего назначения. Игровая оптимизация случайно создала суперкомпьютер на видеокарте.

CUDA: как к этому вообще добраться

Ок, у нас суперкомпьютер на видеокарте. Но попасть на него можно было только через графический API — OpenGL или DirectX. Хочешь посчитать уравнение? Данные закодируй как текстуру, вычисление оформи как шейдер, результат прочитай обратно как картинку. Учёный буквально притворяется что рисует картинку, чтобы GPU посчитал ему физику. Напоминает Infinite Storage Glitch — проект, где ребята кодируют файлы в пиксели видео и заливают на YouTube как бесплатное облачное хранилище. Платформа для видео, но технически стримит любые байты.

Решение появилось довольно быстро. В 2004 году Ian Buck в Стэнфорде сделал Brook — первый язык для GPU без графического API. Корявый, ограниченный — но доказал что идея работает. NVIDIA его наняла, и уже в 2007 вышла CUDA.

Что изменилось принципиально? Вместо "закодируй данные как текстуру, притворись что рисуешь" — три простые идеи:

три идеи CUDA

1. Просто напиши функцию — помечаешь её __global__ и говоришь "запусти на 10 000 потоков". Всё. GPU сам раскидает по ядрам.

2. Иерархия потоков — threads собираются в blocks, blocks в grid. Ты думаешь о структуре задачи, а не о железе. А помните warps из Главы II? Вот как это связано: ты говоришь "блок из 256 потоков", а GPU внутри нарезает его на warps по 32 — те самые группы с общим декодером и SIMD. Программист управляет блоками, GPU управляет warps.

3. Потоки общаются — внутри блока есть быстрая shared memory. Потоки могут обмениваться данными напрямую, не гоняя их через медленную глобальную память.

что пишет программист → что делает GPU Grid block 0 block 1 block 2 ... ты задаёшь Block (256 threads) warp 0 — threads 0..31 warp 1 — threads 32..63 ...warp 7 — threads 224..255 shared memory GPU нарезает на warps SM (железо) 32 ALU = 1 warp за такт (SIMD) переключение warps = latency hiding всё из Главы II программист: grid → blocks · GPU: blocks → warps → SIMD на ALU ты думаешь о задаче, GPU думает о железе
CUDA: просто пиши код __global__ void solve(float *data) { int i = threadIdx.x; data[i] = f(data[i]); } → nvcc → GPU 128 ядер ответ без текстур, без шейдеров, без OpenGL — просто математика

По факту — пишешь нормальный C-код, помечаешь функцию __global__, компилируешь через nvcc и запускаешь. Те же 128 ядер считают твою задачу. Без текстур, без треугольников, без притворства. Просто математика.

Железо не изменилось — изменилось то, как ты с ним разговариваешь. И учёные это оценили быстро, потому что паттерн тот же что в графике: одна формула на миллионы точек.

одна формула × миллионы точек — тот же паттерн что в графике молекулы AMBER: 100× быстрее 1 GPU = 32 CPU ядра сейсмика недели → часы нефтяные компании Monte Carlo финансовые симуляции миллионы сценариев везде один паттерн: возьми формулу, примени к миллионам точек параллельно через год после выхода CUDA — десятки научных групп уже считали на GPU

GPU стал универсальным вычислителем. Но каждая новая задача обнажала слабые места — и требовала новых решений в железе.

Глава IV

От вычислителя к AI-чипу

Ок, у нас универсальный параллельный вычислитель с нормальным интерфейсом. Победа? Не совсем. Каждый раз когда учёные и инженеры начинали реально использовать GPU — они находили новую проблему. И каждое поколение архитектуры — это ответ на конкретную боль предыдущего.

Fermi → Maxwell: GPU учится быть надёжным (2010–2014)

Первая проблема оказалась неожиданной. Учёные запустили физическую симуляцию — а результат на GPU отличается от CPU. Не потому что алгоритм неправильный, а потому что в памяти случайно перевернулся бит. Для игр это незаметно — ну мигнул пиксель. Для науки — катастрофа, расчёт неверный.

три поколения за 4 года — GPU взрослеет Fermi (2010) ECC-память бит перевернулся? чип заметит и поправит наука доверяет GPU Kepler (2012) Dynamic Parallelism GPU сам создаёт потоки не ждёт CPU CPU больше не бутылочное горлышко Maxwell (2014) энергоэффективность больше производительности на каждый ватт можно ставить в дата-центры GPU стал надёжным, самостоятельным и экономным

Pascal: одной карточки мало (2016)

Модели росли быстрее памяти. Одна карточка — 12–16 ГБ, а нужно больше. Что делать? Поставить несколько рядом. Но они общались через PCIe — универсальный разъём на материнской плате, через который подключается вообще всё. Данные шли через CPU: GPU → CPU → GPU, всего ~16 ГБ/с. Pascal сделал прямое соединение GPU-GPU без крюка через CPU — NVLink, 160 ГБ/с, в 10 раз шире. Несколько карточек стали работать почти как одна.

Pascal (2016): соединяем карточки PCIe (было) ~16 ГБ/с данные идут через CPU: GPU → CPU → GPU универсальный, но медленный NVLink (стало) 160 ГБ/с · 10× шире прямое соединение: GPU → GPU напрямую специализированный и быстрый HBM2 память на подложке память стоит рядом с чипом а не на отдельной плате шире канал, меньше задержка несколько карточек работают почти как одна

Volta: возврат специализации (2017)

А вот тут легендарный разворот. Помните, в Главе III мы убрали фиксированные операции с GPU ради гибкости? А теперь NVIDIA возвращает их обратно — но уже для другой задачи. ML — это по сути умножение матриц, миллионы раз подряд. Обычное CUDA-ядро делает одну операцию за такт. Целая матрица 4×4 — 64 такта. Расточительно.

Volta (2017): возврат специализации CUDA-ядро (было) a × b + c за такт матрица 4×4 = 64 такта по одной операции универсально, но медленно для ML Tensor Core (стало) D = A×B+C за 1 такт целая матрица 4×4 сразу 64 операции за такт впаяно в кремний — fixed-function для матриц спираль 1999: fixed-function для освещения 2017: fixed-function для матриц та же идея, новый уровень Tesla V100 стал стандартом data-центров

Ampere: два мира на одном чипе (2020)

Tensor Cores работали отлично, но мир изменился — появились модели на 175 миллиардов параметров (GPT-3). И выяснилось что тренировка и inference — совсем разные задачи. Тренировка хочет максимум вычислений, inference хочет минимум задержки. Один чип должен уметь и то и другое.

A100: три новых трюка TF32 скорость FP16 точность ~FP32 без изменения кода Tensor Cores 3-го поколения Sparsity 2:4 3 0 7 0 2 из 4 = нули аппаратный пропуск 2× ускорение бесплатно MIG 1 GPU → 7 инстансов изоляция для inference A100: первый GPU, который проектировали одновременно для training и inference

Ampere A100 решил это тремя трюками. TF32 — считает быстро как FP16, но точно почти как FP32, и главное без изменения кода. Sparsity 2:4 — если в матрице половина нулей, GPU пропускает их аппаратно, 2× ускорение бесплатно. MIG — один GPU можно порезать на 7 изолированных кусков, на каждом своя задача. Первый чип, который проектировали одновременно для training и inference.

Hopper: эра больших моделей (2022)

LLM — GPT, LLaMA — оказались прожорливыми по-новому. Разные слои модели нуждаются в разной точности, но FP16 везде — расточительно: тратим память и bandwidth на точность которая не нужна.

Hopper H100: чип для LLM Transformer Engine FP8 ↔ FP16 динамически чип сам решает где точнее, а где сэкономить 2× throughput без потери качества HBM3 80 ГБ · 3.35 ТБ/с больше памяти, шире канал модели побольше влезают NVLink 4.0 900 ГБ/с ещё шире связь между карточками кластеры из тысяч GPU до 9× быстрее A100 на LLM — тот самый чип, за которым все охотятся

Blackwell: чип не может расти вечно (2024)

Один монолитный чип упёрся в физику — чем больше кристалл, тем больше дефектов при производстве. Плюс 700W на чип — дата-центр из тысяч таких потребляет как небольшой город.

Blackwell B200: за пределами одного кристалла два кристалла 10 ТБ/с между ними один чип не влезает — склеили два 208 млрд транзисторов FP4 в Tensor Cores ещё менее точный формат для inference хватает, экономия памяти огромная FP32 → FP16 → FP8 → FP4 результат 3× training 15× inference vs Hopper 192 ГБ HBM3e · 8 ТБ/с 700W на чип — дата-центр потребляет как небольшой город

Что дальше: Feynman (2028)

NVIDIA уже показала роадмап. Следующий большой шаг — Feynman (2028). Горизонтально расти некуда — растём вертикально.

Feynman (2028): что дальше 3D-стэкинг чипы друг на друга кристаллы укладываются вертикально, не рядом TSMC 1.6nm (A16) кастомная память NVIDIA делает свою HBM возможно >1 ТБ на GPU память под свои нужды не зависят от Samsung/SK Hynix оптический NVLink свет вместо электричества связь между карточками на скорости света меньше энергии, выше bandwidth горизонтально расти некуда — растём вертикально

Вся эволюция одним взглядом

год проблема решение 2010 GPU считает неправильно — случайные ошибки в памяти защита от ошибок, как контрольная сумма 2016 одной карточки мало, несколько не могут быстро общаться соединили карточки быстрым каналом (NVLink, 10× шире) 2017 ML = умножение матриц миллионы раз, ядра считают по одной операции впаяли блоки которые умножают целую матрицу за раз 2020 training и inference — разные задачи, один чип должен уметь оба чип для обоих задач + можно порезать на части 2022 LLM тратят память впустую — высокая точность где не нужно чип сам решает где считать точнее, а где сэкономить 2024 один чип не может расти бесконечно — дефекты, энергия склеили два кристалла, добавили сверхнизкую точность FP4 2028 горизонтально расти некуда чипы стакаются вертикально, своя память, оптическая связь каждое поколение — ответ на конкретную проблему предыдущего

Глава V

От картиночек до почти AGI

Мы прошли путь от штуки, которая рисует треугольники в Quake, до чипов, на которых тренируют модели с сотнями миллиардов параметров. GPU начинался как специализированный ускоритель для одной задачи, стал универсальным вычислителем — а теперь снова специализируется, только уже под матричные умножения. И именно эта новая специализация породила вопрос: а может, GPU — не единственный ответ?

Проблема inference

Когда LLM генерирует текст — один токен за раз. На каждый токен модель читает все свои веса из памяти: миллиарды параметров. Вычислений на прочитанный байт — крошечное количество. Для полной утилизации H100 нужно ~300 операций на каждый прочитанный байт, а реально получается ~1. Тысячи ядер и Tensor Cores простаивают, ожидая данные из HBM. Утилизация 5–15%. Купил карточку за $30K, а используешь на 10% — обидно.

Главная проблема: как не ждать память. TPU, Groq, Cerebras — все решают именно это, но по-разному.

В GPU тысячи ядер, и каждое на каждом шаге бегает в общую память за данными. Посчитал — записал обратно. Посчитал — записал. Универсально, но ядра постоянно ждут пока память ответит.

Google TPU: данные текут, а не стоят

TPU устроен принципиально иначе. Давай пошагово на примере. В GPU каждое ядро на каждом шаге лезет в память: забрал число, умножил, записал обратно, забрал следующее. В TPU так:

Шаг 1. Загружаем веса из памяти в решётку — один раз. Каждый элемент решётки запоминает свой вес. Элемент [0,0] запомнил W₁, элемент [0,1] запомнил W₂, и так далее. Веса "застыли" в решётке.

Шаг 2. Пускаем входные данные X₁ слева в первый элемент. Он умножает X₁ на свой W₁ и передаёт результат соседу справа. Не в память — напрямую соседу.

Шаг 3. Сосед получает результат, прибавляет свой W₂×X₂, передаёт дальше. Следующий — то же самое. На выходе справа — готовый результат строки матрицы.

Шаг 4. Пока первая строка "проходит" через решётку, сзади уже запускается вторая. Элементы не простаивают ни одного такта.

Итого: GPU обращается к памяти на каждую операцию. TPU — один раз на всю матрицу. Разница в количестве обращений огромная, и именно поэтому для матричных умножений TPU быстрее.

GPU HBM (внешняя память) каждое ядро бегает в память и обратно посчитал → в память → забрал TPU (systolic array) вход ×+ ×+ ×+ ×+ выход вход ×+ ×+ ×+ ×+ выход данные текут от соседа к соседу в память только на входе и выходе конвейер: посчитал → передал соседу GPU: ядра бегают в память · TPU: данные текут через ядра для матричных умножений — конвейер эффективнее

Главная ставка Google — масштаб: TPU Pod из тысяч чипов, соединённых быстрой сетью. Потенциал: если ты в экосистеме Google Cloud — может быть дешевле и быстрее GPU. Но ты привязан к Google, своё не поставишь.

Groq LPU: убрали память вообще

Groq решил проблему ещё радикальнее. Помните из Главы II — HBM это внешняя память, стоит рядом с чипом но физически отдельно. Каждое обращение — сотни тактов задержки. SRAM — память прямо внутри чипа, на том же кристалле, доступ за 1-2 такта. Groq подумал: если при inference мы просто читаем веса по порядку и модель влезает в SRAM — зачем вообще ходить во внешнюю память? Убрали HBM, всё на чипе. Задержка исчезла, скорость inference детерминистическая. Но ограничение очевидное: SRAM конечный, большие модели не влезают.

GPU vs Groq: где живут данные GPU чип SRAM мало HBM 80 ГБ 200+ тактов задержки ядра ждут память Groq LPU SRAM — всё на чипе модель + вычисления рядом 1-2 такта · нет HBM ничего не ждёт убрали бутылочное горлышко — убрали внешнюю память

Cerebras WSE: чип размером с пластину

Cerebras зашёл с третьей стороны: если проблема в том что данные далеко от вычислений — сделаем чип гигантским, чтобы всё влезло. Один кристалл размером с целую кремниевую пластину — 850 000 ядер, 44 ГБ SRAM. Зачем нарезать пластину на маленькие чипы и потом соединять обратно, если можно оставить как есть? Потенциал: невероятная плотность. Но производство сложное и экосистема маленькая.

обычный чип vs Cerebras: размер решает обычный GPU пластину нарезают на чипы используют один, остальные — другим Cerebras WSE 850K ядер · 44 ГБ SRAM · один кристалл зачем нарезать и соединять обратно, если можно оставить целиком?
три подхода к одной проблеме Google TPU systolic array данные текут через решётку без обращения к памяти ставка: масштаб (TPU Pod) Groq LPU только SRAM, без HBM вся модель на чипе 0 обращений к внешней памяти ставка: скорость inference Cerebras WSE чип = целая пластина 850K ядер, 44 ГБ SRAM данные и вычисления рядом ставка: плотность вычислений GPU — не единственный ответ. Но пока самый универсальный.

Что из этого следует

В начале этого поста GPU был для меня чёрным ящиком. Закинул модель, нажал кнопку, ждёшь. Теперь — нет. Мы прошли от рисования треугольников в Quake до чипов, на которых тренируют модели, приближающие нас к AGI. И за каждым решением — не магия, а конкретная инженерная проблема и конкретный компромисс.

Для меня главный вывод такой: когда разбираешь что-то до базовых принципов — перестаёшь бояться сложного. GPU казался невероятно запутанной штукой, а оказался набором остроумных решений под ограничения физики. Параллелизм — потому что пиксели независимы. SIMD — потому что формула одинаковая. Latency hiding — потому что память медленная. Tensor Cores — потому что ML это матрицы. Любая технология так устроена — разбери до первопринципов и станет понятно. И это навык, который переносится на всё остальное.

Но знать архитектуру — половина дела. Вторая — уметь писать код, который её учитывает. Как самому написать CUDA kernel. Какие библиотеки и решения существуют. Как понять, что твой код использует GPU на 10%, и что с этим делать.

Об этом — в следующем посте.

Источники