Войти
| Дата выпуска | Апрель 2010 г. |
|---|---|
| Процесс изготовления | 40 нм и 28 нм |
| История | |
| Предшественник | Tesla 2.0 |
| Преемник | Кеплер |
Ферми - это кодовое имя графического процессора (GPU) микроархитектура, разработанная Nvidia, впервые выпущенная в розничную продажу в апреле 2010 года в качестве преемника микроархитектуры Tesla. Это основная микроархитектура, используемая в сериях GeForce 400 и GeForce 500 серии. За ним последовал Kepler, и он использовался вместе с Kepler в GeForce 600 series, GeForce 700 series и GeForce 800 series, в последние два только в мобильных графических процессорах. На рынке рабочих станций Fermi нашел применение в серии Quadro x000, моделях Quadro NVS, а также в вычислительных модулях Nvidia Tesla. Все графические процессоры Fermi для настольных ПК производились по 40 нм, мобильные графические процессоры Fermi - по 40 и 28 нм. Fermi - старейшая микроархитектура от NVIDIA, получившая поддержку API рендеринга Microsoft Direct3D 12, уровень возможности 11.
Архитектура названа в честь Энрико Ферми, итальянского физика.
Рис. 1. Архитектура NVIDIA Fermi. Условные обозначения в цифрах: оранжевый - планирование и отправка; зеленый - исполнение; светло-голубые - регистры и кэши. Графические процессоры Fermi (графические процессоры ) содержат 3,0 миллиарда транзисторов, а схема схематически показана на рис. 1.
Каждый SM имеет 32 CUDA одинарной точности ядра, 16 единиц загрузки / хранения, четыре специальных функциональных блока (SFU), блок высокоскоростной встроенной памяти объемом 64 КБ (см. подраздел L1 + Shared Memory) и интерфейс к кэш-памяти L2 (см. подраздел кэш-память L2).
Разрешить вычисление адресов источника и назначения для 16 потоков за такт. Загружать и сохранять данные из / в кэш или DRAM.
Выполнять трансцендентные инструкции, такие как синус, косинус, обратная величина и квадратный корень. Каждый SFU выполняет одну инструкцию на поток за такт; деформация выполняется более восьми часов. Трубопровод SFU отделен от блока диспетчеризации, позволяя блоку диспетчеризации передавать другим исполнительным блокам, пока SFU занят.
Целочисленное арифметическое логическое устройство (ALU): поддерживает полную 32-битную точность для всех инструкций, в соответствии с требованиями стандартного языка программирования. Он также оптимизирован для эффективной поддержки 64-битных операций и операций повышенной точности.
Реализует новый стандарт IEEE 754-2008 с плавающей запятой, предоставляя команду объединенного умножения-сложения (FMA) как для одиночной, так и для арифметика двойной точности. На SM за такт может выполняться до 16 операций слитного умножения-сложения двойной точности.
Fused multiply-add (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. полностью используется. Следовательно, невозможно использовать SFU для выполнения более 2 операций на ядро CUDA за цикл.
Теоретическая мощность обработки с двойной точностью графического процессора Fermi составляет 1/2 от производительности с одинарной точностью на GF100 / 110. Однако на практике эта мощность двойной точности доступна только на профессиональных картах Quadro и Tesla, тогда как потребительские карты GeForce ограничены до 1/8.
Кэш L1 на SM и унифицированный кэш L2, который обслуживает все операции (загрузка, сохранение и текстура).
Каждый SM имеет 32 КБ 32-битных регистров. Каждый поток имеет доступ к своим регистрам, а не к регистрам других потоков. Максимальное количество регистров, которое может использоваться ядром CUDA, составляет 63. Число доступных регистров постепенно уменьшается с 63 до 21 по мере увеличения рабочей нагрузки (и, следовательно, требований к ресурсам) с увеличением количества потоков. Регистры имеют очень высокую пропускную способность: около 8000 ГБ / с.
Встроенная память, которая может использоваться либо для кэширования данных для отдельных потоков (переполнение регистров / L1-кеш) и / или для совместного использования данных между несколькими потоками (общая память). Эти 64 КБ памяти могут быть сконфигурированы либо как 48 КБ общей памяти с 16 КБ кеш-памяти L1, либо как 16 КБ общей памяти с 48 КБ кеш-памяти L1. Совместно используемая память позволяет потокам внутри одного и того же блока потоков взаимодействовать, способствует многократному использованию данных на кристалле и значительно снижает объем трафика вне кристалла. Общая память доступна потокам в том же блоке потока. Он обеспечивает доступ с малой задержкой (10-20 циклов) и очень высокую пропускную способность (1600 ГБ / с) для умеренных объемов данных (например, промежуточных результатов в серии вычислений, одной строки или столбца данных для матричных операций, строки видео и т. д.). Дэвид Паттерсон говорит, что эта общая память использует идею локальной блокнота
Локальная память - это область памяти, используемая для хранения «разлитых» регистров. Распространение регистров происходит, когда блоку потока требуется больше регистрового хранилища, чем доступно на SM. Локальная память используется только для некоторых автоматических переменных (которые объявлены в коде устройства без каких-либо квалификаторов __device__, __shared__ или __constant__). Как правило, автоматическая переменная находится в регистре, за исключением следующего: (1) массивы, которые компилятор не может определить, индексируются с постоянными величинами; (2) большие структуры или массивы, которые занимали бы слишком много места в регистрах; Любая переменная, которую компилятор решает передать в локальную память, когда ядро использует больше регистров, чем доступно на SM.
768 KB унифицированный L2-кеш, совместно используемый 16 SM, который обслуживает всю загрузку и сохранение из / в глобальную память, включая копии в / из хоста ЦП, а также запросы текстур. Подсистема кэша L2 также реализует атомарные операции, используемые для управления доступом к данным, которые должны совместно использоваться блоками потоков или даже ядрами.
Доступна для всех потоков, а также для хоста (ЦП). Высокая задержка (400-800 циклов).
См. Nvidia NVDEC (ранее назывался NVCUVID), а также Nvidia PureVideo и Nvidia NVENC.
