Из Википедии, бесплатной энциклопедии
Перейти к навигации Перейти к поиску

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

Fermi - это кодовое название микроархитектуры графического процессора (GPU), разработанной Nvidia и впервые выпущенной в розничную продажу в апреле 2010 года в качестве преемника микроархитектуры Tesla . Это основная микроархитектура, используемая в сериях GeForce 400 и GeForce 500 . За ним последовал Kepler , и он использовался вместе с Kepler в сериях GeForce 600 , GeForce 700 и GeForce 800 , в последних двух только в мобильных графических процессорах. На рынке рабочих станций 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

Графические процессоры ( GPU ) Fermi содержат 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) для арифметических операций с одинарной и двойной точностью. До 16 операций умножения-сложения с двойной точностью может быть выполнено на SM за такт. [1]

Polymorph-Engine [ править ]

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 Engine [ править ]

Механизм 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. [2] Следовательно, невозможно использовать SFU для выполнения более 2 операций на ядро ​​CUDA за цикл.

Теоретическая мощность обработки с двойной точностью графического процессора Fermi составляет 1/2 от производительности с одинарной точностью на GF100 / 110. Однако на практике эта мощность с двойной точностью доступна только на профессиональных картах Quadro и Tesla , в то время как потребительские карты 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

См. Также [ править ]

  • Список графических процессоров Nvidia
  • CUDA
  • Масштабируемый интерфейс связи (SLI)
  • Qualcomm Adreno

Ссылки [ править ]

  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 месяцев позже, стоило ли ждать?» . AnandTech . п. 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 года .

Общие [ править ]

  • Н. Бруквуд, "NVIDIA решает загадку вычислений на GPU".
  • П. Н. Гласковский, "NVIDIA Fermi: первая полная вычислительная архитектура на основе графических процессоров".
  • Н. Уайтхед, А. Фит-Флореа, «Точность и производительность: с плавающей запятой и соответствие стандарту IEEE 754 для графических процессоров NVIDIA». , 2011.
  • С.Ф. Оберман, М. Сиу, "Высокопроизводительный многофункциональный интерполятор с эффективным использованием площади", Proc. от 17 IEEE Symposium по компьютерной арифметике , Кап - Коде, штат Массачусетс, США, июль. 27-29, 2005, стр. 272-279.
  • Р. Фарбер, «Проектирование и разработка приложений CUDA», Морган Кауфманн, 2011 г.
  • Примечание по применению NVIDIA «Настройка приложений CUDA для Fermi».

Внешние ссылки [ править ]

  • Архитектура NVIDIA Fermi в решениях Orange Owl