Чипы NVIDIA Fermi появятся в небольших объемах в марте

Последние выходные я потратил на освоение программирования CUDA и SIMT. Это плодотворно проведённое время закончилось почти 700-кратным ускорением моего «рейтрейсера на визитке» [1] — с 101 секунд до 150 мс. Такой приятный опыт стал хорошим предлогом для дальнейшего изучения темы и эволюции архитектуры Nvidia. Благодаря огромному объёму документации, опубликованному за долгие годы «зелёной» командой, мне удалось вернуться назад во времени и вкратце пройтись по удивительной эволюции её потоковых мультипроцессоров. В этой статье мы рассмотрим: Год Поколение Серия Кристалл Техпроцесс Самая мощная карта
===========================================================================
2006 Tesla GeForce 8 G80 90 nm 8800 GTX
2010 Fermi GeForce 400 GF100 40 nm GTX 480
2012 Kepler GeForce 600 GK104 28 nm GTX 680
2014 Maxwell GeForce 900 GM204 28 nm GTX 980 Ti
2016 Pascal GeForce 10 GP102 16 nm GTX 1080 Ti
2018 Turing GeForce 20 TU102 12 nm RTX 2080 Ti

Тупик

Вплоть до 2006 года архитектура GPU компании NVidia коррелировала с логическими этапами API рендеринга[2].

GeForce 7900 GTX, управлявшаяся кристаллом G71, состояла из трёх частей, занимавшихся обработкой вершин (8 блоков), генерацией фрагментов (24 блоков), и объединением фрагментов (16 блоков). Кристалл G71. Обратите внимание на оптимизацию Z-Cull, отбрасывающую фрагмент, не прошедший бы Z-тест. Эта корреляция заставила проектировщиков угадывать расположение «узких места» конвейера для правильной балансировки каждого из слоёв. С появлением в DirectX 10 ещё одного этапа — геометрического шейдера, инженеры Nvidia столкнулись со сложной задачей балансировки кристалла без знания того, насколько активно будет использоваться этот этап. Настало время для перемен.

Tesla

Nvidia решила проблему роста сложности при помощи «объединённой» архитектуры Tesla, выпущенной в 2006 году. В кристалле G80 больше не было различий между слоями. Благодаря возможности выполнения вершинного, фрагментного и геометрического «ядра», потоковый мультипроцессор (Stream Multiprocessor, SM) заменил все существовавшие ранее блоки. Уравновешивание нагрузки выполнялось автоматически, благодаря замене выполняемого каждым SM «ядра» в зависимости от требований конвейера. «Фактически, мы выбросили всю шейдерную архитектуру NV30/NV40 и с нуля создали новую, с новой общей архитектурой универсальных процессоров (SIMT), в которой также были введены новые методологии проектирования процессоров».

Джона Албен (интервью extremetech.com)

Больше не имеющие возможности выполнять инструкции SIMD «блоки шейдеров» превратились в «ядра», способные выполнять по одной целочисленной инструкции или по одной инструкции с float32 за такт. SM получает потоки в группах по 32 потока, называемых warp. В идеале все потоки одного warp выполняют одновременно одну и ту же инструкцию, только для разных данных (отсюда и название SIMT). Многопотоковый блок инструкций (Multi-threaded Instruction Unit, MT) занимается включением/отключением потоков в warp-е в случае, если их указатель инструкций (Instruction Pointer, IP) сходится/отклоняется. Два блока SFU помогают выполнять сложные математические вычисления, например, обратный квадратный корень, sin, cos, exp и rcp. Эти блоки также способны выполнять по одной инструкции за такт, но поскольку их только два, скорость выполнения warp-а делится на четыре. Аппаратная поддержка float64 отсутствует, вычисления выполняются программно, что сильно влияет на скорость выполнения. SM реализует свой максимальный потенциал, когда способен скрывать задержки памяти благодаря постоянному наличию диспетчеризируемых warp-ов, но также когда поток в warp-е не отклоняется (управляющая логика удерживает его на одном пути выполнения инструкций). Состояния потоков хранятся в 4-килобайтных файлах регистров (Register File, RF). Потоки, занимающие слишком большое пространство в стеке, снижают количество возможных потоков, которые могут выполняться одновременно, понижая при этом производительность.

Кристаллом-флагманом поколения Tesla был 90-нанометровый G80, представленный в GeForce 8800 GTX. Два SM объединены в кластер обработки текстур (Texture Processor Cluster, TPC) вместе с текстурным блоком (Texture Unit) и кешем Tex L1.

Обещалось, что G80 с 8 TPC и 128 ядрами генерирует 345,6 гигафлопс[3]. Карта 8800 GTX была в своё время чрезвычайно популярна, она получила замечательные отзывы и полюбилась тем, кто мог себе её позволить.

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

G80, установленный в 8800 GTX. Render Output Units (ROP) занимаются выполнением сглаживания. Вместе с Tesla компания Nvidia представила язык программирования C для Compute Unified Device Architecture (CUDA) — надмножество языка C99. Это понравилось энтузиастам GPGPU, приветствовавшим альтернативу обмана GPU при помощи текстур и шейдеров GLSL. Хотя в этом разделе я в основном рассказываю о SM, это была только одна половина системы. В SM необходимо передавать инструкции и данные, хранящиеся в памяти GPU. Чтобы избежать простоев, GPU не пытаются минимизировать переходы в память при помощи больших кешей и прогнозирования, как это делают CPU. GPU пользуются задержкой, насыщая шину памяти для удовлетворения потребностей ввода-вывода тысяч потоков. Для этого кристалл (например, G80) реализует высокую пропускную способность памяти при помощи шести двусторонних шин памяти DRAM. GPU пользуются задержками памяти, в то время как CPU скрывают их при помощи огромного кеша и логике прогнозирования.

Fermi

Tesla была рискованным ходом, оказавшимся очень успешным. Она была настолько успешной, что стала фундаментом для GPU компании NVidia на следующие два десятка лет. «Хотя с тех пор мы, конечно же, внесли серьёзные архитектурные изменения (Fermi была серьёзным изменением архитектуры системы, а Maxwell стал ещё одним крупным изменением в проектировании процессоров), фундаментальная архитектура, представленная нами в G80, и сегодня осталась такой же [Pascal]».

Джона Албен (интервью extremetech.com)

В 2010 году Nvidia выпустила GF100, основанный на совершенно новой архитектуре Fermi. Внутренности её последнего чипа подробно описаны в технической документации Fermi[4]. Модель выполнения по-прежнему основана на warp-ах из 32 потоков, диспетчеризируемых в SM. NVidia удалось удвоить/учетверить все показатели только благодаря 40-нанометровому техпроцессу. Благодаря двум массивам из 16 ядер CUDA, SM теперь мог одновременно диспетчеризировать два полу-warp-а (по 16 потоков). При том, что каждое ядро выполняло по одной инструкции за такт, SM по сути был способен исключать по одной инструкции warp за такт (в четыре раза больше, чем у SM архитектуры Tesla). Количество SFU также увеличилось, однако не так сильно — мощность всего лишь удвоилась. Можно прийти к выводу, что инструкции такого типа использовались не очень активно. Присутствует полуаппаратная поддержка float64, при которой комбинируются операции, выполняемые двумя ядрами CUDA. Благодаря 32-битном АЛУ (в Tesla оно было 24-битным) GF100 может выполнять целочисленное умножение за один такт, а из-за перехода от IEEE 754-1985 к IEEE 754-2008 имеет повышенную точность при работе с конвейером float32 при помощи Fused Multiply-Add (FMA) (более точного, чем используемое в Tesla MAD). С точки зрения программирования, объединённая система памяти Fermi позволила дополнить CUDA C такими возможностями C++, как объект, виртуальные методы и исключения.

Благодаря тому, что текстурные блоки стали теперь SM, от концепции TPC отказались. Она была заменена кластерами Graphics Processor Clusters (GPC), имеющими по четыре SM.

И последнее — SM теперь одарён Polymorph Engine, занимающимся получением вершин, преобразованием окна обзора и тесселяцией.

Карта-флагман GeForce GTX 480 на основе GF100 рекламировалась, как содержащая 512 ядер и способная обеспечить 1 345 гигафлопс[5].

GF100, установленный в GeForce GTX 480. Обратите внимание на шесть контроллеров памяти, обслуживающих GPC.

Kepler

В 2012 году Nvidia выпустила архитектуру Kepler, названную в честь астролога, наиболее известного открытием законов движения планет. Как обычно, взглянуть внутрь нам позволила техническая документация GK104[6]. В Kepler компания Nvidia значительно улучшила энергоэффективность кристалла, снизив тактовую частоту и объединив частоту ядер с частотой карты (ранее их частота различалась вдвое). Такие изменения должны были привести к снижению производительности. Однако благодаря вдвое уменьшившемуся техпроцессу (28 нанометров) и замене аппаратного диспетчера на программный, Nvidia смогла не только разместить на чипе больше SM, но и улучшить их конструкцию. Next Generation Streaming Multiprocessor (SMX) — это монстр, почти все показатели которого были удвоены или утроены. Благодаря четырём диспетчерам warp-ов, способным на обработку целого warp-а за один такт (Fermi мог обрабатывать только половину warp-а), SMX теперь содержал 196 ядер. Каждый диспетчер имел двойную диспетчеризацию, позволявшую выполнять вторую инструкцию в warp-е, если она была независима от текущей исполняемой инструкции. Двойная диспетчеризация была не всегда возможна, потому что один столбец из 32 ядер был общим для двух операций диспетчеризации. Такая схема усложнила логику диспетчеризации (к этому мы ещё вернёмся), но благодаря выполнению до шести инструкций warp-ов за такт SMX обеспечивал удвоенную производительность по сравнению с SM архитектуры Fermi.

Заявлялось, что флагманская NVIDIA GeForce GTX 680 с кристаллом GK104 и восемью SMX имеет 1536 ядер, достигающими 3 250 гигафлопс[7]. Элементы кристалла стали настолько запутанными, что мне пришлось убрать со схемы все подписи.

GK104, установленный в GeForce GTX 680. Обратите внимание на полностью переделанные подсистемы памяти, работающие с захватывающей дух частотой 6 ГГц. Они позволили снизить количество контроллеров памяти с шести до четырёх.

Maxwell

В 2014 году Nvidia выпустила GPU десятого поколения под названием Maxwell. Как говорится в технической документации GM107[8], девизом первого поколения архитектуры стали «Максимальная энергоэффективность и чрезвычайная производительность на каждый потреблённый ватт».

Карты позиционировались для «ограниченных в мощности сред, таких как ноутбуки и PC с малым форм-фактором (small form factor, SFF)». Важнейшим решением стал отказ от структуры Kepler с количеством ядер CUDA в SM, не являющимся степенью двойки: некоторые ядра стали общими и вернулись в работе в режиме половины warp-ов.

Впервые за всю историю архитектуры SMM имел меньше ядер, чем его предшественник: «всего» 128 ядер. Согласование количества ядер и размера warp-ов улучшило сегментацию кристалла, что привело к экономии площади и энергии. Один SMM 2014 года имел столько же ядер (128), сколько вся карта GTX 8800 в 2006 году.

Второе поколение Maxwell (описанное в технической документации GM200[9]) значительно повысило производительность, сохранив при этом энергоэффективность первого поколения.

Техпроцесс оставался на уровне 28 нанометров, поэтому инженеры Nvidia не могли для повышения производительности прибегнуть к простой миниатюризации. Однако уменьшение количества ядер SMM снизило их размер, благодаря чему на кристалле удалось разместить больше SMM. По сравнению с Kepler, второе поколение Maxwell удвоило количество SMM, при этом всего на 25% увеличив площадь кристалла. В списке усовершенствований также можно найти упрощённую логику диспетчеризации, позволившую снизить количество избыточных повторных вычислений диспетчеризации и задержку вычислений, что обеспечило повышение оптимальности использования warp-ов. Также на 15% была увеличена частота памяти.

Изучение структурной схемы Maxwell GM200 уже начинает напрягать глаза. Но мы всё равно внимательно его исследуем. Флагманская карта NVIDIA GeForce GTX 980 Ti с кристаллом GM200 и 24 SMM обещала 3072 ядер и 6 060 гигафлопс[10].

GM200, установленный в GeForce GTX 980 Ti.

Pascal

В 2016 году Nvidia представила Pascal. Техническая документация GP104[11] оставляет ощущение дежавю, потому что Pascal SM выглядит точно так же, как Maxwell SMM. Отсутствие изменений SM не привело к стагнации производительности, потому что 16-нанометровый техпроцесс позволил разместить больше SM и снова удвоить количество гигафлопс.

Среди других серьёзных улучшений была система памяти, основанная на совершенно новой GDDR5X. 256-битный интерфейс памяти благодаря восьми контроллерам памяти обеспечивал скорости передачи в 10 гигафлопс, увеличив на 43% пропускную способность памяти и снизив время простоя warp-ов.

Флагман NVIDIA GeForce GTX 1080 Ti с кристаллом GP102 и 28 TSM обещал 3584 ядер и 11 340 гигафлопс[12].

GP104, установленный в GeForce GTX 1080.

Turing

Выпуском в 2018 году Turing компания Nvidia произвела свой «крупнейший за десять лет архитектурный шаг вперёд»[13]. В «Turing SM» появились не только специализированные ядра Tensor с искусственным интеллектом, но и ядра для трассировки лучей (rautracing, RT).

Такая фрагментированная структура напоминает мне многослойную архитектуру, существовавшую до Tesla, и это ещё раз доказывает, что история любит повторения. Кроме новых ядер, в Turing появилось три важные особенности.

Во-первых, ядро CUDA теперь стало суперскалярным, что позволяет параллельно выполнять инструкции с целыми числами и с числами с плавающей запятой. Если вы застали 1996 год, то это может напомнить вам об «инновационной» архитектуре Pentium компании Intel.

Во-вторых, новая подсистема памяти на GDDR6X, поддерживаемая 16 контроллерами, способна теперь обеспечивать 14 гигафлопс. В-третьих, потоки теперь не имеют общих указателей инструкций (IP) в warp-е. Благодаря появившейся в Volta диспетчеризации Independent Thread Scheduling, каждый поток имеет собственный IP.

В результате этого SM способны гибче настраивать диспетчеризацию потоков в warp-е без необходимости как можно более быстрого их схождения.

Флагманская карта NVIDIA GeForce GTX 2080 Ti с кристаллом TU102 и 68 TSM имеет 4352 и достигает 13 45 гигафлопс[14]. Я не стал рисовать структурную схему, потому что она выглядела бы как размытое зелёное пятно.

Что ждёт нас дальше

По слухам, следующая архитектура под кодовым названием Ampere будет объявлена в 2020 году.

Так как Intel доказала на примере Ice Lake, что по-прежнему существует потенциал миниатюризации при помощи 7-нанометрового техпроцесса, почти нет сомнения в том, что Nvidia использует его для дальнейшего уменьшения SM и удвоения производительности.

Терафлопс/с для каждого кристалла/карты Nvidia (источник данных: techpowerup.com). Интересно будет посмотреть, как Nvidia продолжит эволюцию идеи кристаллов, имеющих три типа ядер, выполняющих разные задачи. Увидим ли мы кристаллы, целиком состояние из Tensor-ядер или RT-ядер? Любопытно.

Справочные материалы

[ 1] Источник: Revisiting the Business Card Raytracer [ 2] Источник: Fermi: The First Complete GPU Computing Architecture [ 3] Источник: NVIDIA GeForce 8800 GTX (techpowerup.com) [ 4] Источник: Fermi (GF100) whitepaper [ 5] Источник: NVIDIA GeForce GTX 480 [ 6] Источник: Kepler (GK104) whitepaper [ 7] Источник: NVIDIA GeForce GTX 680 [ 8] Источник: Maxwell Gen1 (GM107) whitepaper [ 9] Источник: Maxwell Gen2 (GM200) whitepaper [10] Источник: NVIDIA GeForce GTX 980 Ti [11] Источник: Pascal (GP102) whitepaper [12] Источник: NVIDIA GeForce GTX 1080 Ti [13] Источник: Turing (TU102) whitepaper [14] Источник: NVIDIA GeForce GTX 2080 Ti

Ферми (микроархитектура) — Fermi (microarchitecture)

Ферми это кодовое имя для графический процессор (GPU) микроархитектура разработан Nvidia, впервые выпущенный в розницу в апреле 2010 года, как преемник Тесла микроархитектура.

Это была основная микроархитектура, использованная в GeForce 400 серии и GeForce 500 серии. За этим последовало Кеплер, и использовался вместе с Кеплером в GeForce 600 серии, GeForce 700 серии, и GeForce 800 серии, в последних двух только в мобильный GPU.

На рынке рабочих станций Fermi нашел применение в Quadro x000, модели Quadro NVS, а также в Nvidia Tesla вычислительные модули. Все графические процессоры Fermi для настольных ПК производились по 40 нм, мобильные графические процессоры Fermi — по 40 и 28 нм.

Fermi — старейшая микроархитектура от NVIDIA, получившая поддержку API рендеринга Microsoft Direct3D 12 feature_level 11.

Архитектура названа в честь Энрико Ферми, итальянский физик.

Обзор

Рис. 1. Архитектура NVIDIA FermiУсловные обозначения цифрами: оранжевый — планирование и отправка; зеленый — исполнение; светло-голубой — регистры и тайники. Снимок графического процессора GF100 внутри видеокарт GeForce GTX 470

Графические процессоры Fermi (GPU ) имеют 3,0 миллиарда транзисторов, а их схема схематически представлена на рис.1.

Потоковый мультипроцессор (SM): состоит из 32 CUDA ядра (см. разделы «Многопроцессорная потоковая передача» и «Ядро CUDA»).
Глобальный планировщик GigaThread: распределяет блоки потоков между планировщиками потоков SM и управляет переключением контекста между потоками во время выполнения (см. Раздел «Планирование деформации»).
Хост-интерфейс: подключает графический процессор к процессору через шину PCI-Express v2 (пиковая скорость передачи 8 ГБ / с).
DRAM: поддерживается до 6 ГБ памяти GDDR5 DRAM благодаря возможности 64-разрядной адресации (см. Раздел «Архитектура памяти»).
Тактовая частота: 1,5 ГГц (не выпущена NVIDIA, но оценена Insight 64).
Пиковая производительность: 1,5 Тфлопс.
Частота глобальной памяти: 2 ГГц.
DRAM пропускная способность: 192 ГБ / с.

Потоковый мультипроцессор

Каждый SM имеет 32 ядра CUDA одинарной точности, 16 блоков загрузки / хранения, четыре блока специальных функций (SFU), блок высокоскоростной встроенной памяти объемом 64 КБ (см. Подраздел L1 + Shared Memory) и интерфейс для кеш-памяти L2 ( см. подраздел L2 Cache).

Загрузка / хранение единиц

Позволяет рассчитывать адреса источника и назначения для 16 потоков за такт. Загружать и сохранять данные из / в тайник или же DRAM.

Единицы специальных функций (SFU)

Выполняйте трансцендентные инструкции, такие как грех, косинус, обратная величина и квадратный корень. Каждый SFU выполняет одну инструкцию на поток за такт; деформация выполняется более восьми часов. Трубопровод SFU отделен от блока диспетчеризации, что позволяет блоку диспетчеризации передавать данные другим исполнительным блокам, пока SFU занят.

Ядро CUDA

Целочисленное арифметико-логическое устройство (ALU): Поддерживает полную 32-битную точность для всех инструкций, в соответствии с требованиями стандартного языка программирования. Он также оптимизирован для эффективной поддержки 64-битных операций и операций повышенной точности.

Модуль с плавающей запятой (FPU)

Реализует новый стандарт IEEE 754-2008 с плавающей запятой, обеспечивая слитное умножение-сложение (FMA) инструкция для арифметических операций с одинарной и двойной точностью. На SM за такт может выполняться до 16 операций умножения-сложения с двойной точностью.[1]

Полиморф-двигатель

Этот раздел пуст. Вы можете помочь добавляя к этому. (Сентябрь 2019)

Слитное умножение-сложение

Слитное умножение-сложение (FMA) выполняет умножение и сложение (например, A * B + C) с одним заключительным шагом округления без потери точности при сложении. FMA более точен, чем выполнение операций по отдельности.

Планирование деформации

В архитектуре Fermi используется двухуровневая распределенная нить планировщик.

Каждый SM может выдавать инструкции, использующие любые два из четырех зеленых столбцов выполнения, показанных на схематическом рис. 1. Например, SM может смешивать 16 операций из 16 ядер первого столбца с 16 операциями из 16 ядер второго столбца, или 16 операций. из единиц загрузки / сохранения с четырьмя из SFU или любые другие комбинации, указанные программой.

Обратите внимание, что 64-битная плавающая точка операции занимают оба первых двух столбца выполнения. Это означает, что SM может одновременно выполнять до 32 операций с плавающей запятой одинарной точности (32-бит) или 16 операций с плавающей запятой двойной точности (64-бит).

GigaThread Engine

Механизм GigaThread планирует блоки потоков для различных SM

Планировщик двойной деформации

На уровне SM каждый планировщик деформации распределяет основы из 32 потоков по своим исполнительным модулям. Потоки планируются в группах по 32 потока, называемых деформациями.

Каждый SM имеет два планировщика деформации и два блока диспетчеризации команд, что позволяет одновременно выдавать и выполнять две деформации.

Планировщик двойной деформации выбирает две деформации и выдает по одной инструкции от каждой деформации группе из 16 ядер, 16 единиц загрузки / сохранения или 4 SFU.

Большинство инструкций может быть выдано двойным; две целочисленные инструкции, две инструкции с плавающей запятой или сочетание инструкций с целыми числами, с плавающей запятой, загрузки, сохранения и SFU могут выполняться одновременно.Двойная точность инструкции не поддерживают двойную отправку с любыми другими операциями.[нужна цитата ]

Спектакль

Теоретическая одинарная точность вычислительная мощность графического процессора Fermi в GFLOPS вычисляется как 2 (операций на инструкцию FMA на ядро CUDA за цикл) × количество ядер CUDA × тактовая частота шейдера (в ГГц).

Обратите внимание, что предыдущее поколение Тесла мог параллельно передавать MAD + MUL на ядра CUDA и SFU, но Fermi потерял эту способность, так как он может выдавать только 32 инструкции за цикл на SM, что позволяет полностью использовать только 32 ядра CUDA.

[2] Следовательно, невозможно использовать SFU для выполнения более 2 операций на ядро CUDA за цикл.

Теоретическая мощность обработки с двойной точностью графического процессора Fermi составляет 1/2 от производительности с одинарной точностью на GF100 / 110. Однако на практике эта мощность двойной точности доступна только на профессиональных Quadro и Тесла карты, а потребитель GeForce карты ограничены до 1/8.[3]

объем памяти

Кэш L1 на SM и унифицированный кеш L2, который обслуживает все операции (загрузка, сохранение и текстура).

Регистры

Каждый SM имеет 32 КБ 32-битных регистров. Каждый поток имеет доступ к своим регистрам, а не к регистрам других потоков. Максимальное количество регистров, которое может использоваться ядром CUDA, составляет 63.

Число доступных регистров постепенно уменьшается с 63 до 21 по мере увеличения рабочей нагрузки (и, следовательно, требований к ресурсам) с увеличением количества потоков.

Регистры имеют очень высокую пропускную способность: около 8000 ГБ / с.

Встроенная память, которая может использоваться либо для кэширования данных для отдельных потоков (переполнение регистров / кэш L1) и / или для обмена данными между несколькими потоками (общая память).

Эта память объемом 64 КБ может быть сконфигурирована либо как 48 КБ общей памяти с 16 КБ кеш-памяти L1, либо как 16 КБ общей памяти с 48 КБ кеш-памяти L1.

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

Он обеспечивает доступ с малой задержкой (10-20 циклов) и очень высокой пропускная способность (1600 ГБ / с) для умеренных объемов данных (таких как промежуточные результаты в серии вычислений, одна строка или столбец данных для матричных операций, строка видео и т. Д.). Дэвид Паттерсон говорит, что эта общая память использует идею локального блокнот[4]

Локальная память

Под локальной памятью подразумевается область памяти, используемая для хранения «пролитых» регистров. Распространение регистров происходит, когда блоку потока требуется больше регистрового хранилища, чем доступно на SM.

Локальная память используется только для некоторых автоматических переменных (которые объявлены в коде устройства без каких-либо квалификаторов __device__, __shared__ или __constant__).

Как правило, автоматическая переменная находится в регистре, за исключением следующего: (1) массивы, которые компилятор не может определить, индексируются с постоянными величинами; (2) большие структуры или массивы, которые занимали бы слишком много места в регистрах; Любая переменная, которую компилятор решает передать в локальную память, когда ядро использует больше регистров, чем доступно на SM.

Кэш L2

Унифицированный кэш L2 размером 768 КБ, совместно используемый 16 модулями SM, который обслуживает всю загрузку и хранение из / в глобальную память, включая копии в / из хоста ЦП, а также запросы текстур. Подсистема кэша L2 также реализует атомарные операции, используемые для управления доступом к данным, которые должны совместно использоваться блоками потоков или даже ядрами.

Глобальная память

Доступен для всех потоков, а также для хоста (ЦП). Высокая задержка (400-800 циклов).

Распаковка / сжатие видео

Видеть Nvidia NVDEC (ранее назывался NVCUVID), а также Nvidia PureVideo и Nvidia NVENC.

Чипы Ферми

GF100
GF104
GF106
GF108
GF110
GF114
GF116
GF118
GF119
GF117

Общий

Н. Бруквуд, «NVIDIA решает загадку вычислений на GPU».
П.Н. Гласковский, «Ферми от NVIDIA: первая полная вычислительная архитектура на GPU».
Н. Уайтхед, А. Фит-Флореа, «Точность и производительность: соответствие с плавающей точкой и IEEE 754 для графических процессоров NVIDIA»., 2011.
С.Ф. Оберман, М. Сиу, «Высокопроизводительный многофункциональный интерполятор с эффективным использованием площади», Proc. из 17-й симпозиум IEEE по компьютерной арифметике, Кэп-Код, Массачусетс, США, 27–29 июля 2005 г., стр. 272–279.
Р. Фарбер, «Проектирование и разработка приложений CUDA», Морган Кауфманн, 2011 г.
Примечание по применению NVIDIA «Настройка приложений CUDA для Fermi».

внешняя ссылка

Архитектура NVIDIA Fermi в решениях Orange Owl

NVIDIA Fermi – процессор из трех миллиардов транзисторов

Осень 2009 года принесла некоторое оживление на рынок графических адаптеров. В сентябре компания AMD презентовала видеокарты ATI Radeon HD 5870 и ATI Radeon HD 5850 на основе процессоров RV870.

Сразу же стало понятно, что до появления нового флагмана компании NVIDIA именно эти видеокарты являются самими производительными из всех однопроцессорных графических адаптеров.

По сравнению со своими предшественниками — ATI Radeon HD 4870 (которые сами являются весьма удачными продуктами компании AMD) – видеокарты ATI Radeon HD 5870 обладают вдвое более высокой производительностью.

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

Модернизация видеочипов именно в этом направлении привела к вполне ожидаемому резкому увеличению и количества транзисторов, из которых состоят интегральные схемы – если конструкция процессора ATI Radeon HD 4870 предусматривала использование 0,956 млрд транзисторов, то в случае ATI Radeon HD 5870 их количество составило уже 2,15 млрд. Но даже эта астрономическая цифра меркнет перед новым продуктом NVIDIA, анонсированного представителями компании в ходе конференции GPU Technologies Conference – графическим процессором NVIDIA Fermi.

Разработчики из Калифорнии основательно подошли к созданию графического процессора нового поколения – микрочипы Fermi (кстати, их «старое» кодовое обозначение – GT300) состоят из более чем трех миллиардов транзисторов.

Эта цифра сразу на 40% выше количества транзисторов у процессора RV870, при том, что они являются продуктами одного поколения и изготавливаются по одному технологическому процессу – 40-нм на мощностях тайваньской компании TSMC.

Если сравнивать процессоры Fermi с решениями предыдущего поколения GT200, то преимущество в количестве транзисторов и вовсе двукратное – 1,4 млрд против 3 млрд.

Увеличение количества транзисторов вполне предсказуемо сказалось на характеристиках процессоров: по сравнению с GT200 увеличено до 512 количество вычислительных блоков, увеличена до 384 бит разрядность интерфейса графической памяти (шесть 64-разрядных блоков), реализована поддержка памяти стандарта GDDR5 максимальным объемом аж до 6 Гб. Надеемся, что все читатели помнят о поддержке процессорами GT200 стандарта GDDR3, а значит, переход на более скоростную графическую память позволит заметно увеличить возможности соответствующей подсистемы видеокарт. Ожидалась и аппаратная поддержка DirectX 11, которая была реализована разработчиками. Но на этом сюрпризы не заканчиваются, ведь дополнительные транзисторы «потрачены» не только на простое увеличение количества исполнительных блоков, как сделали инженеры AMD/ATI. В отличие от их продукта, графический процессор имеет заметно переработанную архитектуру, в которой реализовано значительное количество интересных и передовых (для графических процессоров как класса) нововведений.

На «верхнем» уровне архитектуры графических процессоров существенных качественных отличий не наблюдается. С этой позиции Fermi можно рассматривать всего лишь как масштабированную версию графических процессоров GT200.

Но как только мы обращаем свое внимание на «нижние» уровни архитектуры, на ее фундамент, то сразу же появляются существенные нововведения инженеров NVIDIA. Первое на что необходимо обращать внимание – на графические ядра, которые ранее обозначались как потоковые процессоры (SP — Streaming Processor).

Сейчас разработчики вместо привычного термина перешли на использования термина CUDA-ядер (CUDA Core). В случае графических интегральных микросхем G80 и GT200 разработчики объединяли по восемь потоковых процессоров в единые группы – потоковые мультипроцессорные SM-блоки (SM — Streaming Multiprocessor).

Похожая организация сохранена и для процессоров Fermi, за тем лишь исключением, что теперь в единый блок объединяются не восемь, а тридцать два ядра.

В зависимости от конкретной реализации, в основу графических процессоров NVIDIA Fermi будут входить до шестнадцати SM-блоков, состоящих из 32 CUDA-ядер, оптимизированных для работы с вычислениями общего назначения. В результате и получаются 512 вычислительных ядра, которые и являются базой для высочайшей производительности процессоров.

В дополнение к упомянутым CUDA-ядрам, в состав мультипотоковых блоков входят и дополнительные вычислительные элементы. Речь идет о блоках Special Function Unit (SFU), основная область применения которых — трансцендентальная математика и интерполяция.

Впрочем, блоки SFU сложно назвать главными исполнительными элементами графического процессора. На это указывает и тот факт, что разработчики решили лишь удвоить количество этих компонентов для каждого из SM-блоков графического процессора – до четырех штук.

К сожалению, более подробную информацию об основных исполнительных компонентах архитектуры NVIDIA Fermi разработчики пока предпочли не разглашать.

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

Поэтому пока практически невозможно сказать, насколько интереснее архитектура Fermi применительно к компьютерным играм по сравнению со своими предшественниками в лице GT200.

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

В случае видеочипов предыдущих поколений разработчики также оснащали SM-блоки кэш-памятью первого и второго уровней фиксированным объемом 24 Кб и 256 Кб соответственно. Но использовалась она лишь для хранения «текстурных» данных, к тому же, блок кэш-памяти уровня L1 распределялся между тремя потоковыми мультипроцессорными блоками.

Графические процессоры следующего поколения оснащены уже универсальной кэш-памятью, причем каждому SM-блоку соответствует собственный блок кэш-памяти первого уровня. Но что самое важное, эта память является конфигурируемой.

Что это означает? Каждый SM-блок имеет доступ к 64 Кб памяти, которая может быть разделена непосредственно на кэш-память первого уровня и разделяемую память, причем разделена двумя способами: 48 Кб/16 Кб, либо 16 Кб/48 Кб. Такой подход позволяет решить сразу несколько задач.

Во-первых, возможность различной конфигурации кэш и разделяемой памяти означает полную совместимость с приложениями, оптимизированными для работы с графическими процессорами GT200 с их памятью объемом 16 Кб.

Во-вторых, при выполнении вычислений общего назначения, специализированная традиционная «текстурная» кэш-память оказывается неэффективной – эта проблема решена в случае Fermi с ее универсальной кэш-памятью первого и второго уровней.

В-третьих, возможность различной конфигурации кэш-памяти позволяет с максимальной эффективностью организовать работу программного обеспечения, оптимизированного для работы с большим объемом кэш-памяти – до 48 Кб. И последнее нововведение – увеличение объема универсальной кэш-памяти второго уровня до 768 Кб, которое позволяет существенно повысить производительность при работе с так называемыми атомарными операциями, часто используемыми при вычислениях общего назначения. Согласно заверениям разработчиков, эффективность работы повышается по сравнению с GT200 в 4 – 20 раз (!).

Как архитектура NVIDIA эволюционировала от Теслы до Тьюринга

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

ГодАрхитектурасерияУмеретьЛитографический процессНаиболее представительный график

2006	Тесла	GeForce 8	G80	90nm	8800 GTX
2010	ферми	GeForce 400	GF100	40nm	GTX 480
2012	Кеплер	GeForce 600	GK104	28 нм	GTX 680
2014	Максвелл	GeForce 900	GM204	28 нм	980 GTX Ti
2016	Паскаль	GeForce 10	GP102	16nm	1080 GTX Ti
2018	Тьюринг	GeForce 20	TU102	12nm	RTX 2080 Ti

Пат: эпоха до NVIDIA Tesla

До прихода Tesla в 2006 году NVIDIA GPU / ГРАФИЧЕСКИЙ ПРОЦЕССОР дизайн был соотнесен с логическими состояниями API рендеринга. В GeForce GTX 7900 Работает на матрице G71 и состоит из трех секций (одна предназначена для обработки вершин (8 единиц), другая для генерации фрагментов (24 единицы) и другая для объединения этих (16 единиц)).

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

Пришло время менять как вы разрабатываете свою архитектуру .

Архитектура NVIDIA Tesla

NVIDIA решила проблему увеличения сложности с помощью своей первой «унифицированной» архитектуры Tesla в 2006 году. умереть G80 был больше нет различий между слоями.

Потоковые мультипроцессоры (SM) заменили все предыдущие диски благодаря их способности выполнять обработку вершин, генерацию сегментов и соединение сегментов без различия в одном ядре.

Таким образом, кроме того, нагрузка автоматически уравновешивается путем обмена «ядер», выполняемых каждым SM, в зависимости от потребностей каждого момента.

https://www.youtube.com/watch?v=YE1vuELvX7g\u0026t=11s

Таким образом, шейдерные блоки теперь являются «ядрами» (больше не совместимыми с SIMD), которые способны самостоятельно обрабатывать целочисленные инструкции или инструкции с плавающей запятой (SM получают потоки группами по 32 потока, называемыми деформациями).

В идеале все потоки в деформации будут выполнять один и тот же оператор одновременно только с разными данными (отсюда и название SIMT).

Блок многопоточных (MT) инструкций отвечает за включение и отключение потоков на каждой деформации в случае, если указатели инструкций (IP) сходятся или различаются.

Две СФУ единицы (вы можете увидеть их на диаграмме выше) отвечают за сложные математические вычисления, такие как обратные квадратные корни, синусы, косинусы, exp и rcp.

Эти блоки также способны выполнять одну инструкцию для каждого тактового цикла, но поскольку их всего две, скорость выполнения делится на четыре для каждого из них (то есть для каждого четырех ядер используется один SFU).

Аппаратная поддержка вычислений float64 отсутствует, и они выполняются программно, что значительно снижает производительность.

SM работает с максимальным потенциалом, когда задержку памяти можно устранить, всегда имея программируемые деформации в очереди выполнения, но также и тогда, когда поток деформации не имеет расхождений (это то, для чего нужен поток управления, который поддерживает их всегда в одном и том же путь инструкций). Файл журнала ( 4KB РФ ), где хранятся состояния потоков, а потоки, которые потребляют слишком много очереди выполнения, уменьшают количество их в этом журнале, что также снижает производительность.

«Флагманским» кристаллом этой архитектуры NVIDIA Tesla был G90, основанный на 90-нанометровом литографическом процессе, представленный в знаменитой GeForce 8800 GTX.

Два SM сгруппированы в кластер текстурного процессора (TPC) вместе с текстурным блоком и кешем Tex L1. С 8 TPC у G80 было 128 ядер, генерирующих полную мощность 345.6 Гфлопс.

GeForce 8800 GTX в то время пользовалась огромной популярностью.

С архитектурой Tesla NVIDIA также представила язык программирования CUDA (Compute Unified Devide Architecture) в C, расширенный набор C99, который был долгожданным облегчением для энтузиастов GPGPU, которые приветствовали альтернативу обмануть GPU с помощью шейдеров и текстур GLSL.

Хотя в этом разделе основное внимание уделялось SM, они составляли лишь половину системы.

SM должны получать инструкции и данные, находящиеся в графической памяти графического процессора, поэтому, чтобы избежать застоя, графические процессоры не избегают «отключений» памяти с большим количеством кэш-памяти, как это делают процессоры (ЦП), а, скорее, они загромождают память bus для запросов ввода / вывода от тысяч управляемых потоков. Для этого в микросхеме G80 была реализована высокая производительность памяти через шесть двунаправленных каналов памяти DRAM.

Архитектура Ферми

Тесла был очень рискованным шагом, но он оказался очень хорошим, и он был настолько успешным, что стал основой архитектур NVIDIA на следующие два десятилетия. В 2010 году NVIDIA выпустила GF100 умереть на основе нового ферми архитектура, с многочисленными новыми функциями внутри.

Модель исполнения все еще вращается вокруг 32-проводных деформаций, запрограммированных в SM, и это было только благодаря 40нм литография (по сравнению с 90-нм техпроцессом Tesla), что NVIDIA почти все увеличила в четыре раза.

SM теперь может программировать два медиа-деформирования (16 потоков) одновременно благодаря двум наборам из 16 ядер CUDA.

Когда каждое ядро выполняло одну инструкцию за такт, один SM смог выполнить одну инструкцию деформации за цикл (а это в 4 раза больше, чем у Tesla SM).

Счетчик SFU также был усилен, хотя и в меньшей степени, потому что он «просто» умножился на два, всего четыре единицы.

Также была добавлена аппаратная поддержка вычислений float64, которых не хватало Tesla, выполняемых двумя объединенными ядрами CUDA.

GF100 может выполнять целочисленное умножение за один такт благодаря 32-битному ALU (по сравнению с 24-битным в Tesla) и имеет более высокую точность с плавающей запятой.

https://www.youtube.com/watch?v=YE1vuELvX7g\u0026t=70s

С точки зрения программирования, унифицированная система памяти Fermi позволила расширить CUDA C такими функциями C ++, как виртуальные объекты и методы.

Поскольку текстурные блоки теперь являются частью SM, понятие TPC исчезло, его заменили кластеры графического процессора (GPC), которые имеют четыре SM.

И наконец, что не менее важно, был добавлен механизм «Полиморфа» для обработки вершин объектов, преобразования представления и тесселяции.

Графическим флагманом этого поколения была GTX 480 с 512 ядрами, полная мощность которой составляла 1,345 Гфлопс.

Архитектура NVIDIA Kepler

В 2012 году появилась архитектура NVIDIA Kepler, с которой энергоэффективность его матрицы была значительно улучшена за счет понижения тактовой частоты и объединения центральных тактовых импульсов с тактовой частотой карты (они раньше имели частоту в два раза), что позволило решить проблему GTX 480 предыдущего поколения (они сильно нагревались и имели очень высокий расход).

Эти изменения должны были привести к снижению производительности, но благодаря внедрению 28-нм литографического процесса и исключению аппаратного программиста в пользу программного, NVIDIA смогла не только добавить больше SM, но и улучшить свой дизайн. ,

«Мультипроцессор потоковой передачи следующего поколения», известный как SMX, оказался монстром, в котором почти все увеличилось вдвое или даже втрое.

Благодаря четырем программаторам деформации, способным обрабатывать всю деформацию за один такт (по сравнению с конструкцией Ферми, разделенной на две половины), SMX теперь содержит 196 ядер.

У каждого программиста есть двойная диспетчеризация для выполнения второй инструкции в деформации, если она не зависит от выполняемой в данный момент инструкции, хотя это двойное программирование не всегда было возможным.

Такой подход усложнил логику программирования, но при наличии до шести команд деформации за такт SMX Kepler обеспечивает удвоенную производительность Fermi SM.

Флагманской графикой этого поколения была GeForce GTX 680 с кристаллами GK104 и 8 SMX, которые содержали невероятное количество 1536 ядер и обеспечивали до 3,250 GFLOP полной мощности.

Архитектура NVIDIA Maxwell

В 2014 году появилась архитектура NVIDIA Maxwell, ее графические процессоры 10-го поколения.

Как они объяснили в своей технической документации, в основе этих графических процессоров лежит «исключительная энергоэффективность и исключительная производительность на потребляемый ватт», и именно поэтому NVIDIA ориентировала это поколение на системы с ограниченным энергопотреблением, такие как мини-ПК и ноутбуки.

Главное решение состояло в том, чтобы отказаться от подхода Кеплера к SMX и вернуться к философии Ферми о работе с искаженными носителями.

Таким образом, впервые в своей истории SMM имеет меньше ядер, чем его предшественник с «всего» 128 ядрами.

Регулировка количества сердечников к размеру основы улучшила структуру матрицы, что привело к значительной экономии занимаемого пространства и потребляемой энергии.

Второе поколение Maxwell значительно улучшило производительность, сохранив энергоэффективность первого поколения.

Поскольку литографический процесс по-прежнему застопорился до 28 нм, инженеры NVIDIA не могли полагаться на миниатюризацию транзисторов для повышения производительности, но меньшее количество ядер на SMM уменьшало их размер, позволяя им соответствовать большему количеству SMM. в тот же день. Maxwell Gen 2 содержит вдвое больше СММ, чем Кеплер, и его площадь увеличена только на 25%.

В списке улучшений мы также должны упомянуть более упрощенную логику планирования, которая уменьшила избыточный пересчет решений по планированию, что уменьшило задержку вычислений, чтобы обеспечить лучшую занятость деформаций. Часы памяти также были увеличены в среднем на 15%.

Диаграмма чипа GM200, которую мы видим выше, почти начинает ранить глаз, верно? Это матрица, которая включала в себя GTX 980 Ti с 3072 ядрами в 24 SMM и обеспечивала общую мощность 6,060 GLOP.

Архитектура NVIDIA Pascal

В 2016 году появилось следующее поколение NVIDIA Pascal, и техническая документация выглядела почти как точная копия SMM Максвелла. Но то, что в SM нет никаких изменений, не означает, что не было никаких улучшений, и на самом деле 16-нм техпроцесс, используемый в этих микросхемах, существенно улучшил производительность за счет возможности разместить больше SM на одном чипе.

Другими важными усовершенствованиями стали система памяти GDDR5X, новинка, обеспечивающая скорость передачи до 10 Гбит / с благодаря восьми контроллерам памяти, а ее 256-битный интерфейс обеспечивает пропускную способность на 43% больше, чем в предыдущем поколении.

Графическим флагманом поколения Pascal был GTX 1080 Ti, с его кристаллом GP102, который вы можете видеть на изображении выше, и его 28 SM, упакованными в общей сложности 3584 ядра для общей мощности 11,340 11.3 GLOP (у нас уже XNUMX TFLOP ).

Архитектура NVIDIA Turing

Запущенная в 2018 году архитектура Turing была «самым большим архитектурным скачком за более чем десятилетие» (по словам NVIDIA).

Был добавлен не только SM Turing, но и впервые было представлено специализированное оборудование для трассировки лучей с тензорными ядрами и ядрами RayTracing.

Этот дизайн предполагает, что матрица снова была «фрагментирована» в стиле слоев до эпохи Тесла, о которой мы рассказывали вам в начале.

В дополнение к новым ядрам, Тьюринг добавил три основные функции: во-первых, ядро CUDA теперь масштабируемое и способно выполнять как целочисленные, так и инструкции с плавающей запятой параллельно (это напомнит многим «революционным» Intel Архитектура Pentium там сзади. 1996 год). Во-вторых, новая система памяти GDDR6, поддерживаемая 16 контроллерами, которые теперь могут достигать 14 Гбит / с, и, наконец, потоками, которые больше не разделяют указатели инструкций в деформациях.

Благодаря независимому программированию потоков, введенному в Volta (который мы здесь не включаем, поскольку это не ориентированная на пользователя архитектура), каждый поток имеет свой собственный IP, и, как результат, SM могут свободно программировать потоки в деформации без необходимость ждать их сближения как можно скорее.

Вершина линейки графиков этого поколения — RTX 2080 Ti с его кристаллом TU102 и 68 TSM, содержащим 4352 ядра, с общей мощностью 13.45 TFLOP. Мы не помещаем его полную блок-диаграмму, как в предыдущих, потому что для того, чтобы она поместилась на экране, ее пришлось бы сжать настолько сильно, что это стало бы размытием.

И что будет дальше?

Как вы хорошо знаете, следующая архитектура NVIDIA называется Ampere, и она наверняка прибудет с производственным узлом на расстоянии 7 нм от TSMC. Мы обновим эту статью, как только все данные будут доступны.