Графические процессоры, или GPU, стали неотъемлемой частью современных вычислений, значительно расширяя возможности в сфере обработки графики, научных расчётов и глубинного обучения. Их архитектура и принципы работы кардинально отличаются от традиционных центральных процессоров (CPU), что позволяет выполнять огромное количество операций параллельно. Чтобы понять, как именно достигается такая высокая производительность GPU и почему оптимизация их работы требует особых стратегий, важно изучить основные технические аспекты, лежащие в основе их устройства. В основе GPU лежит необычное соотношение вычислительной мощности и пропускной способности памяти. Например, NVIDIA A100 — одна из самых продвинутых моделей GPU — способна выполнять порядка 19.
5 триллиона операций с плавающей точкой в секунду (TFLOPS), в то время как её пропускная способность памяти составляет «всего» около 1.5 терабайт в секунду (TB/s). Это создает явный дисбаланс: устройство вычисляет гораздо быстрее, чем успевает получить данные из глобальной памяти. За это отвечает иерархия памяти, внутри которой данные проходят несколько стадий обработки, оптимизированных для снижения задержек. Глобальная память, или VRAM, это большой и многочисленный, но медленный по сравнению с вычислительными модулями ресурс.
В примере с A100 VRAM объемом около 40 ГБ и пропускной способностью 1.5 TB/s выступает в роли основного хранилища данных. Для ускорения доступа GPU организует более быстрые области памяти ближе к вычислительному ядру — shared memory (SRAM) и регистры. Каждый Streaming Multiprocessor (SM) — базовый вычислительный блок внутри GPU — содержит около 192 КБ общей памяти и отдельные регистры для каждого потока. Эти ресурсы предоставляют чрезвычайно быструю локальную память, что позволяет повторно использовать загруженные данные множество раз, снижая обращение к медленной глобальной памяти.
GPU управляет миллионами потоков, сгруппированных в варпы по 32 потока. Вычисления внутри варпа происходят синхронно — все 32 потока выполняют одну и ту же инструкцию одновременно. Это принцип lockstep execution влияет на то, как стоит организовывать код: разветвления и ветвления в исполнении приводят к снижению эффективности, поскольку разные ветви должны исполняться по очереди. Производительность алгоритма на GPU определяется двумя основными режимами — память-зависимый (memory-bound) и вычислительно-зависимый (compute-bound). В первом случае узким местом является пропускная способность памяти: устройство простаивает, ожидая загрузки данных из глобальной памяти при низкой арифметической интенсивности.
Примером служит элементарное поэлементное сложение массива, где операций вычислений мало по отношению к объему загружаемых данных. Во втором случае ограничивающим фактором становится максимальная скорость вычислений SM, когда данные максимально эффективно загружены и используется высокая арифметическая интенсивность — большое число операций на единицу данных. Типичный пример — умножение матриц, где многократное переиспользование данных из быстрой локальной памяти значительно улучшает производительность. Арифметическая интенсивность (AI) — это ключевой показатель, отражающий отношение числа вычислительных операций к количеству загруженных байтов из глобальной памяти. Для GPU с учетом архитектуры NVIDIA A100 граница между режимами находится примерно на отметке 13 FLOPs/байт.
Если AI ниже, производительность лимитируется памятью; если выше — вычислительной мощностью. Для повышения AI используются две популярные стратегии — фьюжн (fusion) и тайлинг (tiling). Фьюжн заключается в объединении последовательно выполняемых простых операций в один более сложный вычислительный процесс, что позволяет сократить промежуточные чтения и записи данных в глобальную память. Например, операция, состоящая из сложения и затем применения функции активации ReLU, выгоднее выполнять в одном ядре, чтобы промежуточный результат не записывался и не читался из VRAM. Тайлинг применяется преимущественно к сложным операциям с потенциалом высокой арифметической интенсивности, например, к умножению матриц.
Вместо того чтобы заставлять каждый поток читать полный набор данных для одного элемента, задача разбивается на небольшие блоки (тайлы), которые загружаются в быструю shared memory. Потоки внутри блока сотрудничают, используя данные в shared memory многократно, значительно снижая необходимость повторного обращения к медленной глобальной памяти. При загрузке данных из глобальной памяти важно обеспечение коалесцированного доступа — когда 32 потока варпа читают последовательный 128-байтный блок памяти одним транзакцией. Это повышает пропускную способность, снижая накладные расходы на множественные мелкие операции чтения. Однако кэш shared memory построен на 32 банках, которые могут одновременно обслуживать только 32 различные адреса.
Если несколько потоков варпа обращаются к одному банку, возникает конфликт, приводящий к сериализации запросов и снижению пропускной способности. Это особенно актуально при работе с матрицами, где доступы к столбцам являются по природе разреженными в памяти. Для решения этой проблемы используется трюк «загрузки с транспонированием»: данные читаются из VRAM с коалесцированным доступом построчно, а при записи в shared memory матрица транспонируется. Это устраняет конфликт банков при последующих вычислениях, так как все потоки обращаются к разным банкам, что позволяет выполнять операции максимально параллельно. Вычислительная часть в рамках блока состоит из повторяющегося цикла загрузки, синхронизации и вычисления, где потоки сначала загружают части матриц, затем синхронизируются через барьер __syncthreads(), чтобы гарантировать завершение загрузки, и после приступают к частичным вычислениям.
Такой подход обеспечивает максимальное переиспользование данных и увеличивает арифметическую интенсивность. Для дальнейшего повышения производительности применяются методы встраивания нескольких выходных значений в локальные регистры одного потока. Вместо того, чтобы каждый поток вычислял единственный элемент, нередко он вычисляет блок элементов, что позволяет увеличить размер тайла и AI без увеличения числа потоков, поддерживая при этом приемлимую нагрузку на shared memory и регистры. Еще одной важной темой является управление занятостью (occupancy) — количеством одновременно активных варпов на Streaming Multiprocessоре. Высокая занятость позволяет скрывать задержки длительных операций, переключаясь на готовые к исполнению варпы.
Однако увеличение использования регистров и shared memory на поток снижает число одновременно запущенных блоков, что может негативно сказаться на общей производительности. Оптимизация достигается балансировкой этих факторов. Экономия энергопотребления и повышение эффективности достигается за счет избегания дивергенций внутри варпа, когда ветвление кода требуется выполнять поочередно, теряя параллелизм. Поэтому код на GPU стремятся писать максимально ветвлено-независимым, не используя сложных ветвлений или заменяя их на арифметические операции с условными выборами. Техники квантования (quantization) играют все большую роль, позволяя снизить разрядность чисел с FP32 до FP16, BF16 и даже INT8.
Это сразу уменьшает объем передаваемых данных из глобальной памяти и повышает AI, а также задействует специализированные аппаратные блоки с более высокой пропускной способностью и скоростью. Так, например, квантование может повысить теоретическую производительность на порядок, открывая новые горизонты для оптимизации ML и научных задач. Однако реальные показатели производительности, которые достигаются при выполнении вычислений, часто заметно ниже пиковых. Это связано с особенностями кода, характеристиками используемых инструкций, энергопотреблением и управлением ресурсами. Вынужденный padding при разборе данных, несовершенная векторизация и прочие аппаратные накладные расходы влияют на итоговый результат.
