Fermi — кодовое название микроархитектуры графического процессора (GPU), разработанной Nvidia , впервые выпущенной в продажу в апреле 2010 года в качестве преемника микроархитектуры Tesla . Это была основная микроархитектура, используемая в сериях GeForce 400 и 500. Все настольные графические процессоры Fermi производились по 40 нм, мобильные графические процессоры Fermi — по 40 нм и 28 нм [ требуется ссылка ] . Fermi — старейшая микроархитектура от Nvidia, которая получила поддержку API рендеринга Microsoft Direct3D 12 feature_level 11.
За Fermi последовал 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+Shared Memory») и интерфейсом к кэшу L2 (см. подраздел «Кэш L2»).
Разрешить вычислять исходные и целевые адреса для 16 потоков за такт. Загружать и сохранять данные из/в кэш или DRAM .
Выполнять трансцендентные инструкции, такие как синус, косинус, обратная величина и квадратный корень. Каждый SFU выполняет одну инструкцию на поток за такт; варп выполняется за восемь тактов. Конвейер SFU отделен от диспетчерского блока, что позволяет диспетчерскому блоку выдавать данные другим исполнительным блокам, пока SFU занят.
Поддерживает полную 32-битную точность для всех инструкций, что соответствует стандартным требованиям языка программирования. [ какой? ] Он также оптимизирован для эффективной поддержки 64-битной архитектуры в моделях рабочих станций и серверов, но искусственно урезан для потребительских версий.
Реализует новый стандарт IEEE 754-2008 с плавающей точкой, предоставляя инструкцию fusion multiply-add (FMA) для арифметики как одинарной, так и двойной точности. До 16 операций fusion multiple-add двойной точности могут быть выполнены на SM за такт. [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 для достижения более 2 операций на ядро CUDA за цикл.
Теоретическая мощность обработки двойной точности графического процессора Fermi составляет 1/2 от производительности одинарной точности на GF100/110. Однако на практике эта мощность двойной точности доступна только на профессиональных картах Quadro и Tesla , в то время как потребительские карты GeForce ограничены 1/8. [4]
Кэш L1 на SM и унифицированный кэш L2, обслуживающий все операции (загрузку, сохранение и текстурирование).
Каждый SM имеет 32K 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.
768 КБ унифицированного кэша L2, общего для 16 SM, который обслуживает все загрузки и сохранения из/в глобальную память, включая копии на/из хоста ЦП, а также запросы текстур. Подсистема кэша L2 также реализует атомарные операции, используемые для управления доступом к данным, которые должны быть общими для блоков потоков или даже ядер.
Глобальная память (VRAM) доступна всем потокам напрямую, а также хостовой системе через шину PCIe. Она имеет высокую задержку в 400-800 циклов. [ необходима цитата ]
См. Nvidia NVDEC (ранее называвшийся NVCUVID), а также Nvidia PureVideo .
Технология Nvidia NVENC тогда еще не была доступна, но была представлена в преемнике — Kepler .
Всего 32 инструкции из одного или двух варпов могут быть отправлены в каждом цикле любым двум из четырех исполнительных блоков в Fermi SM.
Производительность FP64 серии GTX 400 ограничена 1/8 (12,5%) производительности FP32, в отличие от того, что изначально может сделать оборудование — 1/2 (50%) FP32.
