Fermi — кодовое название микроархитектуры графического процессора (GPU) , разработанной Nvidia и впервые выпущенной в розничную продажу в апреле 2010 года в качестве преемника микроархитектуры Tesla . Это была основная микроархитектура, используемая в сериях GeForce 400 и GeForce 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 .
Архитектура названа в честь Энрико Ферми , итальянского физика.
Графические процессоры Fermi ( GPU ) содержат 3,0 миллиарда транзисторов, их схема показана на рис. 1.
Каждый SM имеет 32 ядра CUDA одинарной точности, 16 блоков загрузки/сохранения, четыре блока специальных функций (SFU), блок высокоскоростной встроенной памяти объемом 64 КБ (см. подраздел «L1 + Общая память») и интерфейс к кэшу L2 ( см. подраздел «Кэш L2»).
Разрешить вычисление адресов источника и назначения для 16 потоков за такт. Загружайте и сохраняйте данные из/в кэш или DRAM .
Выполняйте трансцендентные инструкции, такие как синус, косинус, обратное число и квадратный корень. Каждый SFU выполняет одну инструкцию на поток за такт; варп выполняется в течение восьми тактов. Конвейер SFU отделен от модуля диспетчеризации, что позволяет блоку диспетчеризации передавать запросы другим исполнительным блокам, пока SFU занят.
Целочисленно-арифметико-логическое устройство (АЛУ) : поддерживает полную 32-битную точность для всех инструкций в соответствии со стандартными требованиями языка программирования. Он также оптимизирован для эффективной поддержки 64-битных операций и операций повышенной точности.
Реализует новый стандарт IEEE 754-2008 для операций с плавающей запятой, предоставляя команду объединенного умножения-сложения (FMA) для арифметических операций с одинарной и двойной точностью. На SM за такт может выполняться до 16 операций умножения-сложения с двойной точностью. [2]
Слитое умножение-сложение (FMA) выполняет умножение и сложение (т. е. A*B+C) с одним заключительным шагом округления без потери точности при сложении. FMA более точен, чем выполнение операций по отдельности.
В архитектуре Fermi используется двухуровневый распределенный планировщик потоков .
Каждый SM может выдавать инструкции, потребляющие любые два из четырех зеленых столбцов выполнения, показанных на схеме рис. 1. Например, SM может смешивать 16 операций из 16 ядер первого столбца с 16 операциями из 16 ядер второго столбца или 16 операций. из модулей загрузки/сохранения с четырьмя из SFU или любых других комбинаций, заданных программой.
Обратите внимание, что 64-битные операции с плавающей запятой используют оба первых двух столбца выполнения. Это означает, что SM может одновременно выполнять до 32 операций с плавающей запятой одинарной точности (32 бита) или 16 операций с плавающей запятой двойной точности (64 бита).
Движок GigaThread планирует блокировку потоков для различных SM.
На уровне SM каждый планировщик варпов распределяет варпы из 32 потоков по своим исполнительным блокам. Потоки планируются группами по 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 для выполнения более 2 операций на ядро CUDA за цикл.
Теоретическая вычислительная мощность двойной точности графического процессора Fermi составляет половину производительности одинарной точности на GF100/110. Однако на практике эта мощность двойной точности доступна только на профессиональных картах Quadro и Tesla , в то время как потребительские карты GeForce ограничены значением 1/8. [4]
Кэш L1 для каждого SM и унифицированный кеш L2, который обслуживает все операции (загрузка, сохранение и текстурирование).
Каждый SM имеет 32 КБ 32-битных регистров. Каждый поток имеет доступ к своим собственным регистрам, а не к регистрам других потоков. Максимальное количество регистров, которое может использовать ядро CUDA, — 63. Количество доступных регистров плавно уменьшается с 63 до 21 по мере увеличения рабочей нагрузки (и, следовательно, требований к ресурсам) с увеличением количества потоков. Регистры имеют очень высокую пропускную способность: около 8000 ГБ/с.
Встроенная память, которую можно использовать либо для кэширования данных для отдельных потоков (распределение регистров/кэш L1), либо для совместного использования данных между несколькими потоками (общая память). Эту память объемом 64 КБ можно настроить либо как 48 КБ общей памяти с 16 КБ кэша L1, либо как 16 КБ общей памяти с 48 КБ кэша L1. Общая память позволяет потокам в одном и том же блоке взаимодействовать, облегчает повторное использование данных на кристалле и значительно снижает трафик за пределами кристалла. Общая память доступна потокам в одном и том же блоке потоков. Он обеспечивает доступ с малой задержкой (10–20 циклов) и очень высокую пропускную способность (1600 ГБ/с) для умеренных объемов данных (например, промежуточных результатов в серии вычислений, одной строки или столбца данных для матричных операций, строки видео и т. д.). Дэвид Паттерсон говорит, что эта общая память использует идею локального блокнота [5]
Под локальной памятью понимается ячейка памяти, используемая для хранения «разлитых» регистров. Распространение регистров происходит, когда блоку потока требуется больше памяти регистров, чем доступно на SM. Локальная память используется только для некоторых автоматических переменных (которые объявлены в коде устройства без квалификаторов __device__, __shared__ или __constant__). Обычно автоматическая переменная находится в регистре, за исключением следующих случаев: (1) Массивы, которые компилятор не может определить, индексируются постоянными величинами; (2) Большие структуры или массивы, которые будут занимать слишком много регистрового пространства; Любая переменная, которую компилятор решает перенести в локальную память, когда ядро использует больше регистров, чем доступно в SM.
Единый кеш L2 объемом 768 КБ, общий для 16 SM, который обслуживает всю загрузку и сохранение из/в глобальную память, включая копирование в/из хоста ЦП, а также запросы текстур. Подсистема кэша L2 также реализует атомарные операции, используемые для управления доступом к данным, которые должны совместно использоваться блоками потоков или даже ядрами.
Доступен всем потокам, а также хосту (ЦП). Высокая задержка (400-800 циклов).
См. Nvidia NVDEC (ранее называвшаяся NVCUVID), а также Nvidia PureVideo .
Технология Nvidia NVENC еще не была доступна, но была представлена в преемнике Kepler .
В каждом цикле в любые два из четырех исполнительных блоков ферми-СМ может быть отправлено в общей сложности 32 инструкции из одной или двух деформаций.
Производительность FP64 серии GTX 400 ограничена 1/8 (12,5%) производительности FP32, в отличие от того, что аппаратное обеспечение изначально может делать 1/2 (50%) FP32.
