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

Nvidia Fermi
Дата выходаАпрель 2010 г.
Процесс изготовления40 нм и 28 нм
История
ПредшественникТесла 2.0
ПреемникКеплер

Ферми это кодовое имя для графический процессор (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]

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

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

Слитное умножение-сложение (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 + общая память

Встроенная память, которая может использоваться либо для кэширования данных для отдельных потоков (переполнение регистров / кэш 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

Рекомендации

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

Общий

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