Стэнфордский курс CS149: устройство архитектуры GPU и основы CUDA

Лекция Стэнфордского университета в рамках курса CS149 посвящена архитектуре графических процессоров (GPU) и основам параллельного программирования с использованием платформы CUDA. Преподаватель подробно разбирает эволюционный путь видеокарт от узкоспециализированных игровых чипов до мощных вычислительных систем, обеспечивших NVIDIA статус триллионной компании. На конкретных примерах демонстрируются принципы работы модели SPMD, иерархия памяти и внутреннее устройство стриминговых мультипроцессоров.

🎮 От Quake до триллионной империи: Эволюция графических процессоров 0:17

История современных параллельных вычислений неразрывно связана с индустрией компьютерных игр. Изначально чипы, производимые такими компаниями, как NVIDIA и AMD, создавались ради одной практической задачи — обеспечения высокой частоты кадров в играх уровня Quake. Однако стремление к реалистичной графике заставило инженеров наращивать количество вычислительных ядер, что в итоге привело к тектоническому сдвигу в индустрии: сегодня NVIDIA превратилась в технологического гиганта с капитализацией более $1 трлн, далеко обогнав Intel, а за право обладания её ускорителями борются ведущие мировые компании.

Как подчеркивает лектор, базовые концепции параллелизма за последние десятилетия не изменились. В современных GPU используются все те же знакомые разработчикам инструменты:

• Многоядерность (multicore)
• Векторизация (SIMD)
• Многопоточность (multithreading)

Главное отличие кроется исключительно в масштабе реализации этих идей. Показательно, что популярный компилятор ISPC (Intel SPMD Program Compiler) фактически стал симметричным ответом CPU-индустрии на появление CUDA: разработчики ISPC стремились перенести удобную модель программирования графических чипов на классические процессоры.

📐 От треугольников к пикселям: Как графика заставила чипы измениться 2:17

Чтобы понять архитектуру GPU, необходимо обратиться к основам трехмерной графики. Исторически главная задача графического чипа заключалась в том, чтобы взять математическое описание трехмерной сцены (геометрию поверхностей, расположение источников света, параметры виртуальной камеры) и смоделировать поведение световых лучей, выдав на выходе готовую картинку. Уже к 2015 году флагманские видеокарты легко справлялись с этой задачей на скорости 60 кадров в секунду.

В основе вычислений лежит обработка полигональных сеток, чаще всего состоящих из треугольников. Процесс рендеринга делится на два ключевых этапа:

1. Проецирование трехмерных вершин (vertices) треугольников на двухмерную плоскость экрана.
2. Вычисление цвета для каждого пикселя, попавшего внутрь спроецированного треугольника.

В начале 2000-х годов программисты осознали, что мир состоит из слишком сложных и разнообразных материалов — от глянцевого металла до человеческой кожи, сквозь которую свет проникает и рассеивается. Стало очевидно, что жестко зашитые алгоритмы окрашивания неэффективны, и индустрия перешла к концепции программируемых пиксельных шейдеров. Для каждого пикселя на экране запускалась независимая функция, рассчитывающая финальный RGB-цвет на основе текстурных карт и нормалей поверхности.

Потребность обрабатывать изображения в разрешении 4K с частотой обновления от 60 Гц вынудила инженеров массово добавлять в чипы новые вычислительные блоки (ALU) и ядра для параллельного обсчета миллионов независимых пикселей. В этот же период (около 20 лет назад) классические процессоры Intel уперлись в «тепловой барьер» на частоте около 4 ГГц и исчерпали возможности суперскалярной архитектуры по поиску параллелизма на уровне инструкций (ILP) [7:57–8:24]. Рост производительности CPU затормозился, в то время как GPU продолжали наращивать мощность пропорционально закону Мура за счет параллельной архитектуры.

🛠️ Эпоха «хаков» и появление концепции Compute Mode 8:52

Остановка роста производительности CPU побудила исследователей из Стэнфорда и других университетов применить мощность видеокарт для неграфических расчетов. Это породило эпоху масштабных программных «хаков».

Чтобы запустить на GPU симуляцию физики частиц или фолдинг белка, ученые обманывали видеокарту. Они принудительно отрисовывали на весь экран два огромных треугольника, чтобы заставить чип создать миллионы потоков для обсчета пикселей. Внутри пиксельного шейдера вместо расчета цвета прописывались формулы симуляции, а финальный результат в формате RGBA интерпретировался разработчиками как координаты XYZ в пространстве [9:57–10:11].

В 2004 году в Стэнфорде был запущен исследовательский проект Brook, призванный превратить этот «хак» в легитимную модель потокового программирования. Язык Brook позволял применять функции к коллекциям данных в дата-параллельном стиле. На уровне пользователя код выглядел чисто, но под капотом компилятор по-прежнему переводил его в графические инструкции отрисовки треугольников.

Осознав потенциал этого направления, компания NVIDIA в 2007 году представила официальную альтернативу графическому конвейеру — режим вычислений (compute mode) и язык CUDA [13:49–15:06]. Если операционная система на CPU запускает потоки последовательно, поочередно настраивая регистры для каждой программы, то CUDA предложила концепцию bulk-запуска: программист пишет одну функцию (ядро/kernel) и просит GPU запустить сразу $N$ ее независимых копий [15:19–15:32]. Это классическое воплощение модели программирования SPMD (Single Program, Multiple Data).

🏗️ Архитектурные абстракции CUDA: Потоки, блоки и сетки 16:54

В терминологии CUDA базовая единица параллелизма называется «потоком CUDA» (CUDA thread), что концептуально ближе к понятию экземпляра программы (program instance) в ISPC, работающего на векторной линии, нежели к тяжеловесному аппаратному потоку ОС.

Запуск CUDA-ядра осуществляется с помощью специального синтаксиса с тройными угловыми скобками, где указывается конфигурация создаваемых потоков. Для удобства работы с матрицами, трехмерной графикой и тензорами в ИИ идентификаторы потоков могут быть многомерными (2D или 3D). Это позволяет избежать лишних операций деления при расчете адресов в памяти.

Потоки CUDA организуются в иерархическую структуру:

• Потоки (Threads): Индивидуальные исполнители вычислений.
• Блоки потоков (Thread Blocks): Группы потоков, которые гарантированно выполняются на одном физическом ядре и могут взаимодействовать друг с другом.
• Сетка (Grid): Совокупность всех блоков, созданных за один запуск ядра.

Например, для поэлементного сложения матриц размером 12 на 6 разработчик может создать блоки размером 4 на 3 потока [19:46–20:00]. Внутри кода каждый поток определяет свой глобальный индекс, используя встроенные переменные blockIdx (индекс блока), blockDim (размер блока) и threadIdx (индекс потока внутри блока).

Преподаватель обращает особое внимание на поведение SPMD-модели при некратных размерах данных. Если размер матрицы составляет 11 на 5, а размер блоков — 4 на 3, то округление при выделении блоков приведет к созданию избыточного количества потоков. В таком случае в коде критически важна условная конструкция (if), проверяющая выход за границы массива. Без этой проверки поток произведет несанкционированную запись в чужую память, что приведет к аварийному завершению программы (segfault).

💾 Иерархия памяти и оптимизация через Shared Memory 25:53

В базовом режиме работы CPU (хост) и GPU (девайс) имеют раздельные адресные пространства. Попытка разыменовать указатель CPU внутри CUDA-кода вызовет ошибку. Процесс вычислений на GPU классически состоит из четырех шагов:

1. Выделение памяти на GPU через cudaMalloc.
2. Копирование данных из DRAM процессора в DRAM видеокарты через cudaMemcpy по шине PCIe.
3. Запуск CUDA-ядра для обработки данных на GPU.
4. Копирование результатов обратно в память CPU.

Поскольку шина PCIe относительно медленная, этот процесс напоминает модель передачи сообщений (message passing) между изолированными узлами [28:14–28:39].

Помимо глобальной памяти девайса, в CUDA реализована строгая иерархия внутренних хранилищ. Каждый поток имеет приватную локальную память (стек). Но ключевой элемент оптимизации — это общая память блока (Shared Memory), помечаемая модификатором __shared__ [33:26–37:18].

На примере одномерной свертки (1D convolution) лектор демонстрирует силу этого инструмента. В наивной реализации каждый поток считывает три соседних элемента из медленной глобальной памяти, что порождает избыточные дублирующие запросы от соседних потоков [35:31–36:13]. Оптимизированный алгоритм устроен иначе:

• Блок потоков (например, из 128 потоков) выделяет общий массив __shared__ под 130 элементов (с учетом граничных элементов для свертки).
• Каждый поток кооперативно загружает из глобальной памяти ровно один элемент и записывает его в Shared Memory, а первые два потока дозагружают крайние значения [38:27–38:40].
• Вызывается функция барьерной синхронизации __syncthreads(), останавливающая потоки блока до тех пор, пока весь массив не будет заполнен [38:54–39:22].
• Потоки параллельно рассчитывают свертку, считывая данные из быстрой Shared Memory, которая физически является управляемым L1-кэшем на кристалле.

Отвечая на вопрос из аудитории, лектор подчеркивает: удаление __syncthreads() гарантированно сломает программу [40:44–40:58]. Без барьера быстрые потоки начнут вычисления до того, как медленные потоки успеют записать свои элементы в общую память, что приведет к состоянию гонки (race condition).

⚡ Аппаратная архитектура: Внутри стримингового мультипроцессора 46:26

Массовый запуск миллионов потоков в коде не означает, что чип мгновенно создает под них физические контексты. Распределением блоков по реальным ядрам занимается аппаратный планировщик GPU, реализованный на уровне кремния.

На примере архитектуры Volta V100 лектор объясняет устройство базового вычислительного блока GPU — Streaming Multiprocessor (SM). Архитектура субъядра SM включает в себя:

• Блок выборки и декодирования инструкций (Fetch and Decode).
• 16-wide SIMD ALУ для операций с плавающей запятой (FP32), а также блоки для целочисленных вычислений и тригонометрии [51:14–52:06].
• Массив регистров, хранящий контексты выполнения потоков.

Принципиальное отличие GPU от CPU заключается в реализации концепции implicit SIMD (неявного SIMD). На CPU векторные инструкции генерируются компилятором на этапе сборки. На GPU каждый поток формально имеет собственный счетчик команд (Program Counter, PC). Аппаратная логика на лету сравнивает PC тридцати двух последовательных потоков. Если их PC совпадают, чип объединяет их и выполняет одну инструкцию над вектором из 32 элементов за один такт на SIMD-юните [53:43–54:09].

Группа из 32 потоков, выполняющихся синхронно, называется варпом (warp). Варп — это неделимая единица планирования на уровне железа, скрытая от глаз программиста. Так как ширина SIMD-векторов в Volta составляет 16 элементов, а размер варпа — 32 потока, физически выполнение одной инструкции варпа занимает два последовательных такта процессора. Планировщик использует эти такты, чтобы переключаться между обработкой инструкций разного типа (FP32, INT, Load/Store), минимизируя простои декодера [56:12–56:38].

Если в коде возникает ветвление (if/else) и потоки внутри одного варпа расходятся по разным веткам (divergence), GPU применяет маскирование: сначала выполняются инструкции для одной части потоков, а затем — для оставшейся, что снижает эффективность вычислений [1:03:11–1:04:15]. NVIDIA сохраняет абстракцию независимых PC для каждого потока именно ради того, чтобы иметь возможность менять внутреннюю ширину SIMD в будущих поколениях чипов без необходимости перекомпиляции пользовательского софта.

📊 Масштаб вычислений и риски взаимной блокировки 1:00:35

Архитектура одного стримингового мультипроцессора (SM) в чипе Volta V100 способна одновременно хранить контексты для 64 варпов, что эквивалентно 2048 живым CUDA-потокам на одном SM. Физически чип V100 содержит 80 таких мультипроцессоров. Путем несложных математических расчетов (80 SM × 64 варпа × 32 потока) мы получаем ошеломляющую цифру: видеокарта способна удерживать на кристалле 163 000 параллельных потоков одновременно. Для сравнения, топовые процессоры Intel того времени поддерживали всего по два гиперпотока на ядро. Суммарная пиковая производительность V100 достигает 12.7 терафлопс.

Однако ручное управление конфигурацией блоков накладывает жесткие ограничения. Если программист создаст ядро, требующее 2048 потоков на один блок, а целевой GPU старого поколения аппаратно поддерживает максимум 256 контекстов на SM, программа завершится ошибкой на этапе запуска, так как блок физически невозможно разместить на ядре [1:05:27–1:05:52].

В финальной части лекции преподаватель разбирает фундаментальное правило CUDA: блоки потоков должны быть полностью независимы друг от друга. В стандартной CUDA отсутствует вытесняющая многопоточность (preemption) — поток нельзя принудительно «снять» с ядра до завершения его работы.

Это приводит к двум важным сценариям:

1. Гистограммы и атомарные операции (Безопасно): Если потоки из разных блоков параллельно пишут данные в одну область глобальной памяти через atomicAdd, программа отработает корректно [1:16:19–1:17:05]. Нам не важен порядок выполнения блоков; они могут конкурировать за доступ к памяти, но вычисления завершатся успешно [1:17:21–1:17:34].
2. Зависимые блоки (Опасно / Крах системы): Если разработчик создаст два блока, где Блок 1 пишет флаг готовности в память, а Блок 0 в цикле ожидает изменения этого флага, программа может намертво заблокировать GPU [1:17:48–1:18:02]. Если у аппаратного планировщика хватит ресурсов запустить только Блок 0, он займет все вычислительные мощности ядра и будет бесконечно ждать Блок 1, который даже не может начать выполнение из-за нехватки ресурсов. Возникнет неразрешимый дедлок (deadlock).

Разработчик CUDA имеет право использовать барьеры и синхронизацию только внутри одного блока, поскольку потоки одного блока гарантированно находятся на одном SM в один и тот же момент времени. Любые попытки синхронизировать разные блоки между собой до окончания работы ядра нарушают философию архитектуры и ведут к нестабильной работе ПО.