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

Нвидиа Ферми
	NVIDIA GeForce GTX 590 из линейки видеокарт GeForce 500 была последней крупной итерацией с микроархитектурой Fermi (GF110-351-A1).
Дата выпуска	апрель 2010 г.
Производитель	ТСМК
Разработано	Нвидиа
Процесс изготовления	40 нм и 28 нм [ нужна ссылка ]
История
Предшественник	Тесла
Преемник	Кеплер
Статус поддержки
	Не поддерживается

Fermi — кодовое название графического процессора (GPU), микроархитектуры разработанной Nvidia и впервые выпущенной в розничную продажу в апреле 2010 года в качестве преемника микроархитектуры Tesla . Это была основная микроархитектура, используемая в сериях GeForce 400 и 500 . Все настольные графические процессоры Fermi производились по 40-нм техпроцессу, мобильные графические процессоры Fermi - по 40-нм и 28-нм техпроцессу. ^{[ нужна ссылка ]}. Fermi — старейшая микроархитектура от Nvidia, которая получает поддержку Microsoft API рендеринга Direct3D 12 Feature_level 11.

За Ферми последовал Kepler , и он использовался вместе с Kepler в сериях GeForce 600 , GeForce 700 и GeForce 800 , в последних двух только в мобильных графических процессорах.

На рынке рабочих станций Fermi нашел применение в серии Quadro x000, моделях Quadro NVS и в вычислительных модулях Nvidia Tesla .

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

Обзор

Рис. 1. Архитектура NVIDIA Fermi
Конвенция в цифрах: оранжевый — планирование и отправка; зеленый – исполнение; голубой — регистры и кэши.

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

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

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

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

Единицы загрузки/сохранения

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

Подразделения специальных функций (SFU)

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

Ядро CUDA

Целочисленно-арифметико-логическое устройство (АЛУ)

Поддерживает полную 32-битную точность для всех инструкций, что соответствует стандартным требованиям языка программирования. ^{[ который? ]} Он также оптимизирован для эффективной поддержки 64-разрядной версии в моделях рабочих станций и серверов, но искусственно урезан для потребительских версий.

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

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

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

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

Варп-планирование

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

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

Для 64-битных операций с плавающей запятой требуются оба первых двух столбца выполнения, поэтому они выполняются на половине скорости 32-битных операций.

Планировщик двойного варпа

На уровне SM каждый планировщик варпов распределяет варпы из 32 потоков по своим исполнительным блокам. Каждый SM имеет два планировщика варпов и два блока диспетчеризации инструкций, что позволяет одновременно запускать и выполнять два варпа. Планировщик двойного варпа выбирает два варпа и выдает по одной инструкции из каждого варпа группе из 16 ядер, 16 блоков загрузки/хранения или 4 SFU. Большинство инструкций могут быть выданы дважды; две целочисленные инструкции, две инструкции с плавающей запятой или сочетание целочисленных инструкций, инструкций с плавающей запятой, загрузки, сохранения и SFU могут выполняться одновременно. Инструкции двойной точности не поддерживают двойную отправку с любой другой операцией. ^{[ нужна ссылка ]}

Производительность

Теоретическая с одинарной точностью вычислительная мощность графического процессора Fermi в GFLOPS рассчитывается как 2 (операций на инструкцию FMA на ядро CUDA за цикл) × количество ядер CUDA × тактовая частота шейдера (в ГГц). предыдущего поколения Обратите внимание, что Tesla могла параллельно выполнять двойную выдачу MAD+MUL для ядер CUDA и SFU, но Fermi потерял эту возможность, поскольку он может выдавать только 32 инструкции за цикл на SM, что позволяет полностью использовать только 32 ядра CUDA. ^[3] Таким образом, невозможно использовать SFU для выполнения более двух операций на ядро CUDA за цикл.

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

Память

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

Регистры

Каждый SM имеет 32 КБ 32-битных регистров. Каждый поток имеет доступ к своим собственным регистрам, а не к регистрам других потоков. Максимальное количество регистров, которое может использовать ядро CUDA, — 63. Количество доступных регистров плавно уменьшается с 63 до 21 по мере увеличения рабочей нагрузки (и, следовательно, требований к ресурсам) с увеличением количества потоков. Регистры имеют очень высокую пропускную способность: около 8000 ГБ/с.

L1+общая память

Встроенная память, которую можно использовать либо для кэширования данных для отдельных потоков (распределение регистров/кэш L1), либо для совместного использования данных между несколькими потоками (общая память). Эту память объемом 64 КБ можно настроить либо как 48 КБ общей памяти с 16 КБ кэша L1, либо как 16 КБ общей памяти с 48 КБ кэша L1. Общая память позволяет потокам в одном блоке взаимодействовать, облегчает повторное использование данных на кристалле и значительно снижает трафик за пределами кристалла. Общая память доступна потокам в одном и том же блоке потоков. Он обеспечивает доступ с малой задержкой (10–20 циклов) и очень высокую пропускную способность (1600 ГБ/с) для умеренных объемов данных (например, промежуточных результатов в серии вычислений, одной строки или столбца данных для матричных операций, строки видео и т. д.). Дэвид Паттерсон говорит, что эта общая память использует идею локального блокнота. ^[5]

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

Под локальной памятью понимается ячейка памяти, используемая для хранения «разлитых» регистров. Распространение регистров происходит, когда блоку потока требуется больше памяти регистров, чем доступно на SM. Локальная память используется только для некоторых автоматических переменных (которые объявлены в коде устройства без квалификаторов __device__, __shared__ или __constant__). Обычно автоматическая переменная находится в регистре, за исключением следующих случаев: (1) Массивы, которые компилятор не может определить, индексируются постоянными величинами; (2) Большие структуры или массивы, которые будут занимать слишком много регистрового пространства; Любая переменная, которую компилятор решает записать в локальную память, когда ядро использует больше регистров, чем доступно в SM.

Кэш L2

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

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

Глобальная память (VRAM) доступна всем потокам напрямую, а также хост-системе через шину PCIe. Он имеет высокую задержку 400-800 циклов. ^{[ нужна ссылка ]}

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

См. Nvidia NVDEC (ранее называвшаяся NVCUVID), а также Nvidia PureVideo .

Технология Nvidia NVENC еще не была доступна, но была представлена в преемнике Kepler .

Ферми-чипы

ГФ100
ГФ104
ГФ106
ГФ108
ГФ110
ГФ114
ГФ116
ГФ117
ГФ119

См. также

Ссылки

^ «Информация об устройстве-декодере NVIDIA GPU» .
^ «Вычислительная архитектура CUDA нового поколения NVIDIA: Fermi» (PDF) . 2009 . Проверено 7 декабря 2015 г.
^ Глазковски, Питер Н. (сентябрь 2009 г.). «Fermi от NVIDIA: первая полная вычислительная архитектура на графическом процессоре» (PDF) . п. 22 . Проверено 6 декабря 2015 г. В каждом цикле в любые два из четырех исполнительных блоков ферми-СМ может быть отправлено в общей сложности 32 инструкции из одного или двух варпов.
^ Смит, Райан (26 марта 2010 г.). «NVIDIA GeForce GTX 480 и GTX 470: опоздание на 6 месяцев, стоило ли ждать?» . АнандТех . п. 6 . Проверено 6 декабря 2015 г. Производительность FP64 серии GTX 400 ограничена 1/8 (12,5%) от производительности FP32, в отличие от того, что аппаратное обеспечение изначально может делать 1/2 (50%) FP32.
^ Паттерсон, Дэвид (30 сентября 2009 г.). «10 главных инноваций в новой архитектуре NVIDIA Fermi и 3 главных будущих задачи» (PDF) . Лаборатория исследований параллельных вычислений и NVIDIA . Проверено 3 октября 2013 г.

Общий

Н. Бруквуд, «NVIDIA решает проблему вычислений на графических процессорах».
П. Н. Глазковски, «Fermi от NVIDIA: первая полная вычислительная архитектура на графическом процессоре».
Н. Уайтхед, А. Фит-Флореа, «Точность и производительность: числа с плавающей запятой и соответствие стандарту IEEE 754 для графических процессоров NVIDIA». , 2011.
Оберман, Сан-Франциско; Сиу, МЮ (2005). «Высокопроизводительный многофункциональный интерполятор с эффективной площадью». 17-й симпозиум IEEE по компьютерной арифметике (ARITH'05) . стр. 272–279. дои : 10.1109/arith.2005.7 . ISBN 0-7695-2366-8 . S2CID 14975421 .
Р. Фарбер, «Проектирование и разработка приложений CUDA», Морган Кауфманн, 2011 г.
Примечание по применению NVIDIA «Настройка приложений CUDA для Fermi».

Внешние ссылки

Архитектура NVIDIA Fermi в решениях Orange Owl. Архивировано 4 января 2022 г. на Wayback Machine.

[1] «Информация об устройстве-декодере NVIDIA GPU» .

[fermi_whitepaper-2] «Вычислительная архитектура CUDA нового поколения NVIDIA: Fermi» (PDF) . 2009 . Проверено 7 декабря 2015 г.

[3] Глазковски, Питер Н. (сентябрь 2009 г.). «Fermi от NVIDIA: первая полная вычислительная архитектура на графическом процессоре» (PDF) . п. 22 . Проверено 6 декабря 2015 г. В каждом цикле в любые два из четырех исполнительных блоков ферми-СМ может быть отправлено в общей сложности 32 инструкции из одного или двух варпов.

[4] Смит, Райан (26 марта 2010 г.). «NVIDIA GeForce GTX 480 и GTX 470: опоздание на 6 месяцев, стоило ли ждать?» . АнандТех . п. 6 . Проверено 6 декабря 2015 г. Производительность FP64 серии GTX 400 ограничена 1/8 (12,5%) от производительности FP32, в отличие от того, что аппаратное обеспечение изначально может делать 1/2 (50%) FP32.

[5] Паттерсон, Дэвид (30 сентября 2009 г.). «10 главных инноваций в новой архитектуре NVIDIA Fermi и 3 главных будущих задачи» (PDF) . Лаборатория исследований параллельных вычислений и NVIDIA . Проверено 3 октября 2013 г.

[1]

[2]

[3]

[4]

[5]