NVIDIA планирует представить недорогие чипы Fermi для массового рынка

Последние выходные я потратил на освоение программирования CUDA и SIMT. Это плодотворно проведённое время закончилось почти 700-кратным ускорением моего «рейтрейсера на визитке» [1] — с 101 секунд до 150 мс. Такой приятный опыт стал хорошим предлогом для дальнейшего изучения темы и эволюции архитектуры Nvidia. Благодаря огромному объёму документации, опубликованному за долгие годы «зелёной» командой, мне удалось вернуться назад во времени и вкратце пройтись по удивительной эволюции её потоковых мультипроцессоров. В этой статье мы рассмотрим: Год Поколение Серия Кристалл Техпроцесс Самая мощная карта
===========================================================================
2006 Tesla GeForce 8 G80 90 nm 8800 GTX
2010 Fermi GeForce 400 GF100 40 nm GTX 480
2012 Kepler GeForce 600 GK104 28 nm GTX 680
2014 Maxwell GeForce 900 GM204 28 nm GTX 980 Ti
2016 Pascal GeForce 10 GP102 16 nm GTX 1080 Ti
2018 Turing GeForce 20 TU102 12 nm RTX 2080 Ti

Тупик

Вплоть до 2006 года архитектура GPU компании NVidia коррелировала с логическими этапами API рендеринга[2].

GeForce 7900 GTX, управлявшаяся кристаллом G71, состояла из трёх частей, занимавшихся обработкой вершин (8 блоков), генерацией фрагментов (24 блоков), и объединением фрагментов (16 блоков). Кристалл G71. Обратите внимание на оптимизацию Z-Cull, отбрасывающую фрагмент, не прошедший бы Z-тест. Эта корреляция заставила проектировщиков угадывать расположение «узких места» конвейера для правильной балансировки каждого из слоёв. С появлением в DirectX 10 ещё одного этапа — геометрического шейдера, инженеры Nvidia столкнулись со сложной задачей балансировки кристалла без знания того, насколько активно будет использоваться этот этап. Настало время для перемен.

Tesla

Nvidia решила проблему роста сложности при помощи «объединённой» архитектуры Tesla, выпущенной в 2006 году. В кристалле G80 больше не было различий между слоями. Благодаря возможности выполнения вершинного, фрагментного и геометрического «ядра», потоковый мультипроцессор (Stream Multiprocessor, SM) заменил все существовавшие ранее блоки. Уравновешивание нагрузки выполнялось автоматически, благодаря замене выполняемого каждым SM «ядра» в зависимости от требований конвейера. «Фактически, мы выбросили всю шейдерную архитектуру NV30/NV40 и с нуля создали новую, с новой общей архитектурой универсальных процессоров (SIMT), в которой также были введены новые методологии проектирования процессоров».

Джона Албен (интервью extremetech.com)

Больше не имеющие возможности выполнять инструкции SIMD «блоки шейдеров» превратились в «ядра», способные выполнять по одной целочисленной инструкции или по одной инструкции с float32 за такт. SM получает потоки в группах по 32 потока, называемых warp. В идеале все потоки одного warp выполняют одновременно одну и ту же инструкцию, только для разных данных (отсюда и название SIMT). Многопотоковый блок инструкций (Multi-threaded Instruction Unit, MT) занимается включением/отключением потоков в warp-е в случае, если их указатель инструкций (Instruction Pointer, IP) сходится/отклоняется. Два блока SFU помогают выполнять сложные математические вычисления, например, обратный квадратный корень, sin, cos, exp и rcp. Эти блоки также способны выполнять по одной инструкции за такт, но поскольку их только два, скорость выполнения warp-а делится на четыре. Аппаратная поддержка float64 отсутствует, вычисления выполняются программно, что сильно влияет на скорость выполнения. SM реализует свой максимальный потенциал, когда способен скрывать задержки памяти благодаря постоянному наличию диспетчеризируемых warp-ов, но также когда поток в warp-е не отклоняется (управляющая логика удерживает его на одном пути выполнения инструкций). Состояния потоков хранятся в 4-килобайтных файлах регистров (Register File, RF). Потоки, занимающие слишком большое пространство в стеке, снижают количество возможных потоков, которые могут выполняться одновременно, понижая при этом производительность.

Кристаллом-флагманом поколения Tesla был 90-нанометровый G80, представленный в GeForce 8800 GTX. Два SM объединены в кластер обработки текстур (Texture Processor Cluster, TPC) вместе с текстурным блоком (Texture Unit) и кешем Tex L1.

Обещалось, что G80 с 8 TPC и 128 ядрами генерирует 345,6 гигафлопс[3]. Карта 8800 GTX была в своё время чрезвычайно популярна, она получила замечательные отзывы и полюбилась тем, кто мог себе её позволить.

Она оказалась таким превосходным продуктом, что спустя тринадцать месяцев после выпуска оставалась одним из самых быстрых GPU на рынке.

G80, установленный в 8800 GTX. Render Output Units (ROP) занимаются выполнением сглаживания. Вместе с Tesla компания Nvidia представила язык программирования C для Compute Unified Device Architecture (CUDA) — надмножество языка C99. Это понравилось энтузиастам GPGPU, приветствовавшим альтернативу обмана GPU при помощи текстур и шейдеров GLSL. Хотя в этом разделе я в основном рассказываю о SM, это была только одна половина системы. В SM необходимо передавать инструкции и данные, хранящиеся в памяти GPU. Чтобы избежать простоев, GPU не пытаются минимизировать переходы в память при помощи больших кешей и прогнозирования, как это делают CPU. GPU пользуются задержкой, насыщая шину памяти для удовлетворения потребностей ввода-вывода тысяч потоков. Для этого кристалл (например, G80) реализует высокую пропускную способность памяти при помощи шести двусторонних шин памяти DRAM. GPU пользуются задержками памяти, в то время как CPU скрывают их при помощи огромного кеша и логике прогнозирования.

Fermi

Tesla была рискованным ходом, оказавшимся очень успешным. Она была настолько успешной, что стала фундаментом для GPU компании NVidia на следующие два десятка лет. «Хотя с тех пор мы, конечно же, внесли серьёзные архитектурные изменения (Fermi была серьёзным изменением архитектуры системы, а Maxwell стал ещё одним крупным изменением в проектировании процессоров), фундаментальная архитектура, представленная нами в G80, и сегодня осталась такой же [Pascal]».

Джона Албен (интервью extremetech.com)

В 2010 году Nvidia выпустила GF100, основанный на совершенно новой архитектуре Fermi. Внутренности её последнего чипа подробно описаны в технической документации Fermi[4]. Модель выполнения по-прежнему основана на warp-ах из 32 потоков, диспетчеризируемых в SM. NVidia удалось удвоить/учетверить все показатели только благодаря 40-нанометровому техпроцессу. Благодаря двум массивам из 16 ядер CUDA, SM теперь мог одновременно диспетчеризировать два полу-warp-а (по 16 потоков). При том, что каждое ядро выполняло по одной инструкции за такт, SM по сути был способен исключать по одной инструкции warp за такт (в четыре раза больше, чем у SM архитектуры Tesla). Количество SFU также увеличилось, однако не так сильно — мощность всего лишь удвоилась. Можно прийти к выводу, что инструкции такого типа использовались не очень активно. Присутствует полуаппаратная поддержка float64, при которой комбинируются операции, выполняемые двумя ядрами CUDA. Благодаря 32-битном АЛУ (в Tesla оно было 24-битным) GF100 может выполнять целочисленное умножение за один такт, а из-за перехода от IEEE 754-1985 к IEEE 754-2008 имеет повышенную точность при работе с конвейером float32 при помощи Fused Multiply-Add (FMA) (более точного, чем используемое в Tesla MAD). С точки зрения программирования, объединённая система памяти Fermi позволила дополнить CUDA C такими возможностями C++, как объект, виртуальные методы и исключения.

Благодаря тому, что текстурные блоки стали теперь SM, от концепции TPC отказались. Она была заменена кластерами Graphics Processor Clusters (GPC), имеющими по четыре SM.

И последнее — SM теперь одарён Polymorph Engine, занимающимся получением вершин, преобразованием окна обзора и тесселяцией.

Карта-флагман GeForce GTX 480 на основе GF100 рекламировалась, как содержащая 512 ядер и способная обеспечить 1 345 гигафлопс[5].

GF100, установленный в GeForce GTX 480. Обратите внимание на шесть контроллеров памяти, обслуживающих GPC.

Kepler

В 2012 году Nvidia выпустила архитектуру Kepler, названную в честь астролога, наиболее известного открытием законов движения планет. Как обычно, взглянуть внутрь нам позволила техническая документация GK104[6]. В Kepler компания Nvidia значительно улучшила энергоэффективность кристалла, снизив тактовую частоту и объединив частоту ядер с частотой карты (ранее их частота различалась вдвое). Такие изменения должны были привести к снижению производительности. Однако благодаря вдвое уменьшившемуся техпроцессу (28 нанометров) и замене аппаратного диспетчера на программный, Nvidia смогла не только разместить на чипе больше SM, но и улучшить их конструкцию. Next Generation Streaming Multiprocessor (SMX) — это монстр, почти все показатели которого были удвоены или утроены. Благодаря четырём диспетчерам warp-ов, способным на обработку целого warp-а за один такт (Fermi мог обрабатывать только половину warp-а), SMX теперь содержал 196 ядер. Каждый диспетчер имел двойную диспетчеризацию, позволявшую выполнять вторую инструкцию в warp-е, если она была независима от текущей исполняемой инструкции. Двойная диспетчеризация была не всегда возможна, потому что один столбец из 32 ядер был общим для двух операций диспетчеризации. Такая схема усложнила логику диспетчеризации (к этому мы ещё вернёмся), но благодаря выполнению до шести инструкций warp-ов за такт SMX обеспечивал удвоенную производительность по сравнению с SM архитектуры Fermi.

Заявлялось, что флагманская NVIDIA GeForce GTX 680 с кристаллом GK104 и восемью SMX имеет 1536 ядер, достигающими 3 250 гигафлопс[7]. Элементы кристалла стали настолько запутанными, что мне пришлось убрать со схемы все подписи.

GK104, установленный в GeForce GTX 680. Обратите внимание на полностью переделанные подсистемы памяти, работающие с захватывающей дух частотой 6 ГГц. Они позволили снизить количество контроллеров памяти с шести до четырёх.

Maxwell

В 2014 году Nvidia выпустила GPU десятого поколения под названием Maxwell. Как говорится в технической документации GM107[8], девизом первого поколения архитектуры стали «Максимальная энергоэффективность и чрезвычайная производительность на каждый потреблённый ватт».

Карты позиционировались для «ограниченных в мощности сред, таких как ноутбуки и PC с малым форм-фактором (small form factor, SFF)». Важнейшим решением стал отказ от структуры Kepler с количеством ядер CUDA в SM, не являющимся степенью двойки: некоторые ядра стали общими и вернулись в работе в режиме половины warp-ов.

Впервые за всю историю архитектуры SMM имел меньше ядер, чем его предшественник: «всего» 128 ядер. Согласование количества ядер и размера warp-ов улучшило сегментацию кристалла, что привело к экономии площади и энергии. Один SMM 2014 года имел столько же ядер (128), сколько вся карта GTX 8800 в 2006 году.

Второе поколение Maxwell (описанное в технической документации GM200[9]) значительно повысило производительность, сохранив при этом энергоэффективность первого поколения.

Техпроцесс оставался на уровне 28 нанометров, поэтому инженеры Nvidia не могли для повышения производительности прибегнуть к простой миниатюризации. Однако уменьшение количества ядер SMM снизило их размер, благодаря чему на кристалле удалось разместить больше SMM. По сравнению с Kepler, второе поколение Maxwell удвоило количество SMM, при этом всего на 25% увеличив площадь кристалла. В списке усовершенствований также можно найти упрощённую логику диспетчеризации, позволившую снизить количество избыточных повторных вычислений диспетчеризации и задержку вычислений, что обеспечило повышение оптимальности использования warp-ов. Также на 15% была увеличена частота памяти.

Изучение структурной схемы Maxwell GM200 уже начинает напрягать глаза. Но мы всё равно внимательно его исследуем. Флагманская карта NVIDIA GeForce GTX 980 Ti с кристаллом GM200 и 24 SMM обещала 3072 ядер и 6 060 гигафлопс[10].

GM200, установленный в GeForce GTX 980 Ti.

Pascal

В 2016 году Nvidia представила Pascal. Техническая документация GP104[11] оставляет ощущение дежавю, потому что Pascal SM выглядит точно так же, как Maxwell SMM. Отсутствие изменений SM не привело к стагнации производительности, потому что 16-нанометровый техпроцесс позволил разместить больше SM и снова удвоить количество гигафлопс.

Среди других серьёзных улучшений была система памяти, основанная на совершенно новой GDDR5X. 256-битный интерфейс памяти благодаря восьми контроллерам памяти обеспечивал скорости передачи в 10 гигафлопс, увеличив на 43% пропускную способность памяти и снизив время простоя warp-ов.

Флагман NVIDIA GeForce GTX 1080 Ti с кристаллом GP102 и 28 TSM обещал 3584 ядер и 11 340 гигафлопс[12].

GP104, установленный в GeForce GTX 1080.

Turing

Выпуском в 2018 году Turing компания Nvidia произвела свой «крупнейший за десять лет архитектурный шаг вперёд»[13]. В «Turing SM» появились не только специализированные ядра Tensor с искусственным интеллектом, но и ядра для трассировки лучей (rautracing, RT).

Такая фрагментированная структура напоминает мне многослойную архитектуру, существовавшую до Tesla, и это ещё раз доказывает, что история любит повторения. Кроме новых ядер, в Turing появилось три важные особенности.

Во-первых, ядро CUDA теперь стало суперскалярным, что позволяет параллельно выполнять инструкции с целыми числами и с числами с плавающей запятой. Если вы застали 1996 год, то это может напомнить вам об «инновационной» архитектуре Pentium компании Intel.

Во-вторых, новая подсистема памяти на GDDR6X, поддерживаемая 16 контроллерами, способна теперь обеспечивать 14 гигафлопс. В-третьих, потоки теперь не имеют общих указателей инструкций (IP) в warp-е. Благодаря появившейся в Volta диспетчеризации Independent Thread Scheduling, каждый поток имеет собственный IP.

В результате этого SM способны гибче настраивать диспетчеризацию потоков в warp-е без необходимости как можно более быстрого их схождения.

Флагманская карта NVIDIA GeForce GTX 2080 Ti с кристаллом TU102 и 68 TSM имеет 4352 и достигает 13 45 гигафлопс[14]. Я не стал рисовать структурную схему, потому что она выглядела бы как размытое зелёное пятно.

Что ждёт нас дальше

По слухам, следующая архитектура под кодовым названием Ampere будет объявлена в 2020 году.

Так как Intel доказала на примере Ice Lake, что по-прежнему существует потенциал миниатюризации при помощи 7-нанометрового техпроцесса, почти нет сомнения в том, что Nvidia использует его для дальнейшего уменьшения SM и удвоения производительности.

Терафлопс/с для каждого кристалла/карты Nvidia (источник данных: techpowerup.com). Интересно будет посмотреть, как Nvidia продолжит эволюцию идеи кристаллов, имеющих три типа ядер, выполняющих разные задачи. Увидим ли мы кристаллы, целиком состояние из Tensor-ядер или RT-ядер? Любопытно.

Справочные материалы

[ 1] Источник: Revisiting the Business Card Raytracer [ 2] Источник: Fermi: The First Complete GPU Computing Architecture [ 3] Источник: NVIDIA GeForce 8800 GTX (techpowerup.com) [ 4] Источник: Fermi (GF100) whitepaper [ 5] Источник: NVIDIA GeForce GTX 480 [ 6] Источник: Kepler (GK104) whitepaper [ 7] Источник: NVIDIA GeForce GTX 680 [ 8] Источник: Maxwell Gen1 (GM107) whitepaper [ 9] Источник: Maxwell Gen2 (GM200) whitepaper [10] Источник: NVIDIA GeForce GTX 980 Ti [11] Источник: Pascal (GP102) whitepaper [12] Источник: NVIDIA GeForce GTX 1080 Ti [13] Источник: Turing (TU102) whitepaper [14] Источник: NVIDIA GeForce GTX 2080 Ti

Криптовалюта не майнится

Майнеры в России массово начали продавать видеокарты на Avito. Причем цены на них в разы ниже тех, что были у ритейлеров в марте. Например, одну из востребованных видеокарт Nvidia можно купить за 165 тыс. руб.

, в то время как в марте ее стоимость в магазине достигала почти 540 тыс руб. Как сообщает Forbes, рынок добычи криптовалют в июне рухнул примерно в пять раз по сравнению с началом весны.

Эксперты связывают это, в частности, с падением курса биткойна и других цифровых валют, а также с укреплением рубля.

“Ъ FM” спросил авторов объявлений на Avito, почему они продают видеокарты и верят ли в будущее майнинга в России:

«Я собрался с друзьями в Крым, но ехать было не на что. Продавать было жалко, но я выложил объявление. Просмотры были — звонков нет. Я уже обрадовался, думал, что подожду до конца месяца — если что, сниму объявления. У меня их два.

Это либо одна карточка 3090, либо две 3060. Все, что добываю майнингом, я коплю, пока не будет хорошего курса. Другого выхода я не вижу. Прибыль с двух ферм у меня всего 23 тыс. руб. До этого, буквально три месяца назад, было 120 тыс. руб.

»

«Мы работаем под заказ — возим из Китая. Я сам в нем. Как мы можем продавать оборудование для майнинга и не быть в нем? С майнингом все в порядке. Он как был прибыльным, так прибыльным и остается. Просто сейчас точка входа дешевле, если смотреть на перспективу, а не в данный момент. Этот пузырь уже лет 12 пытается лопнуть. Если бы у меня были свободные деньги, я бы точно еще докупил».

«Не так давно, когда были совершенно другие времена, купишь карту — 10–15 тыс. руб. максимум скидываешь, и ее расхватывают. Сейчас упали курсы, еще жарко, и я не готов вкладывать 25–30 тыс. руб. Я снял три карты — остальные пять чувствуют себя более или менее нормально.

Я вообще на ASIC майню — просто решил попробовать на видюхах. Мне легче продать три карты и купить крипту. Потом курс поднимется, я ее продам и возьму эти три карты. Я лично верю, что через 10–15 лет бумагу на руках никто держать не будет — все деньги будут электронные».

Одной из причин массовой продажи видеокарт собеседники “Ъ FM” называют сезонное повышение стоимости электричества. С учетом падения курса криптовалют многие розничные майнинговые фермы становятся убыточными. Поэтому предприниматели вынуждены продавать видеокарты, пока это можно сделать по более или менее выгодной цене. В противном случае устройства могут дешеветь до 30% ежемесячно.

По данным аналитической компании JPR, за 2021 год во всем мире производители поставили на рынок почти 50 млн видеокарт на сумму около $52 млрд.

Теперь на фоне проблем с майнингом они стали невостребованы, считает независимый IT-эксперт, специалист по компьютерным играм Александр Кузьменко.

Кроме того, дефицита видеокарт, несмотря на трудности с логистикой и санкции, по его словам, нет: «Видеокарты из золотого стандарта превращаются в дорогие железки, которые очень выгодно будет продать, если ты успеешь за паническим рынком.

При этом нет никакого понимания, кто может платить за подобные устройства.

С поставкой оборудования в России ситуация запутанная, тем не менее появляются всевозможные серые схемы, и через тот же самый AliExpress можно купить видеокарты по цене, которая не слишком превышает общую и международную.

Если вдруг вам очень сильно захотелось приобрести последнюю видеокарту Nvidia, например, 3080 или 4800, это можно сделать, переплатив буквально 10-20% от ее номинальной стоимости в каком-либо магазине в Европе или в Соединенных Штатах Америки».

Тем временем в Минэнерго обсуждают введение налога за использование электричества для майнеров. Предполагается, что они будут платить за потребление так же, как бизнес.

Потенциальные поступления в госбюджет могут составить до 1,5 трлн руб.

Регулирование криптовалют и сопровождающее их падение сделают рынок майнинга все более промышленным и недоступным для энтузиастов, полагает ведущий аналитик 8848 Invest Виктор Першиков:

«Буквально на днях стало известно, что «Газпромнефть» подписала меморандум о сотрудничестве с крупнейшей майнинговой фермой в России BitRiver.

Теперь компания может предоставлять ей территории технологически незапущенных месторождений с излишком электроэнергии для коммерческой деятельности. Таких факторов в последние месяцы наблюдается довольно много.

Когда майнинг в России будет регулироваться, определенная доля игроков останется в серой зоне, не захочет обеляться и быть в так называемом реестре майнеров.

Но, полагаю, у правоохранительных органов довольно большой объем информации о том, кто и где занимается добычей криптовалют на территории страны. Прежде всего вопрос в том, как они будут с этим работать.

Последние года два-три розничный майнинг постепенно отмирает. Здесь нужно учитывать, что происходят технологические изменения, например, на видеокартах добываются в основном Ethereum и форки Ethereum.

Эта криптовалюта перейдет на новую версию: на смену протоколу Proof-of-Work придет Proof-of-Stake, там майнинга не будет, поэтому так или иначе, скорее всего, в конце этого — начале следующего года майнить не получится».

По данным оператора дата-центров BitRiver, к началу этого года в России насчитывалось около 300 тыс. физических и юридических лиц, которые занимаются добычей криптовалют. Аналитики компании оценивают совокупную усредненную выручку от майнинга примерно в $1,5 млрд в год.

Новости в вашем ритме — Telegram-канал «Ъ FM».

Станислав Корягин

Nvidia заявляет, что новый процессор Fermi будет запускать суперкомпьютеры 2022

В среду Nvidia продемонстрировала новую архитектуру графического процессора, которая, как она надеется, позволит ей выйти за рамки игр, чтобы играть большую роль на рынке суперкомпьютеров.

Новая архитектура Fermi должна обеспечить более реалистичную графику для геймеров, но он также включает в себя технологии, которые делают его хорошо подходящим для высокопараллельных вычислительных сред, сказал генеральный директор Nvidia Джен-Сюнь Хуан в своем выступлении на конференции компании GPU Technology, которая была трансляцией из Сан-Хосе, Калифорния.

Он показал видеокарту на сцене с прототипом нового чипа внутри нее, но он не объявил никаких конкретных планов продукта в среду или сказал, когда Ферми попадет на рынок.

Национальная лаборатория Oak Ridge будет использовать фишки Fermi в суперкомпьютере для sci — сказал Джеффри Николс (Jeffrey Nichols), помощник директора лаборатории, который присоединился к Хуангу на сцене. В настоящее время в лаборатории находится суперкомпьютер Jaguar, разработанный Cray, который основан на процессорах Advanced Micro Devices Opteron.

Микросхема Fermi имеет 3 миллиарда транзисторов и 512 процессорных ядер и включает в себя новую технологию GigaThread 3.0, которая может управлять тысячами потоков Параллельно, сказал Хуан.

Fermi преуспеет в архитектуре NVIDIA G80 GPU, которая была представлена ​​в прошлом году и используется в графических картах GeForce GT200 компании, которые включают до 240 процессорных ядер.

Fermi также содержит память ECC для защиты битов данных от электромагнитных волн или частиц от окружающей среды, которые могут повлиять на выполнение нитей. Huang сказал, что GPU удвоит пропускную способность памяти существующих архитектур с поддержкой памяти DRAM GDDR5 и поддерживает до 1 Тбайт памяти графического процессора.

Архитектура также поддерживает плавную точку с двойной точностью, что может обеспечить до восьми- Nvidia сказал, что повышение производительности для некоторых научных и математических задач. И он будет поддерживать программирование на C ++ в дополнение к программированию на C современных моделей.

Чипы Fermi также будут совместимы с средой разработки CUDA Nvidia, которая помогает разработчикам писать параллельный код. Кроме того, Nvidia и Microsoft совместно объявили Nexus, среду разработки для архитектуры Fermi, которая была интегрирована в среду разработки Microsoft Visual Studio.

Fermi (microarchitecture) — Wikipedia

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 за такт.

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

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.

Следовательно, невозможно использовать 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.

Кэш L2

Унифицированный кэш L2 размером 768 КБ, совместно используемый 16 модулями SM, обслуживает всю загрузку и хранение из / в глобальную память, включая копии в / из хоста ЦП, а также запросы текстур. Подсистема кэша L2 также реализует атомарные операции, используемые для управления доступом к данным, которые должны совместно использоваться блоками потоков или даже ядрами.

Глобальная память

Доступен для всех потоков, а также для хоста (ЦП). Высокая задержка (400-800 циклов).

Распаковка / сжатие видео

См. Nvidia NVDEC (ранее назывался NVCUVID), а также Nvidia PureVideo .

Технология Nvidia NVENC еще не была доступна, но была представлена ​​в преемнике Kepler .

Чипы Fermi

• GF100
• GF104
• GF106
• GF108
• GF110
• GF114
• GF116
• GF118
• GF119
• GF117

Смотрите также

использованная литература

Общий

• Н. Бруквуд, «NVIDIA решает загадку вычислений на GPU».
• П. Н. Гласковский, «NVIDIA Fermi: первая полная вычислительная архитектура на основе графических процессоров».
• Н. Уайтхед, А. Фит-Флореа, «Точность и производительность: с плавающей точкой и соответствие стандарту IEEE 754 для графических процессоров NVIDIA». , 2011.
• Оберман, СФ; Сиу, MY (2005). «Высокоэффективный многофункциональный интерполятор». 17-й симпозиум IEEE по компьютерной арифметике (ARITH'05) . С. 272–279. DOI : 10.1109 / arith.2005.7 . ISBN 0-7695-2366-8. S2CID  14975421 .
• Р. Фарбер, «Проектирование и разработка приложений CUDA», Морган Кауфманн, 2011 г.
• Примечание по применению NVIDIA «Настройка приложений CUDA для Fermi».

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

NVIDIA Fermi – процессор из трех миллиардов транзисторов

Осень 2009 года принесла некоторое оживление на рынок графических адаптеров. В сентябре компания AMD презентовала видеокарты ATI Radeon HD 5870 и ATI Radeon HD 5850 на основе процессоров RV870.

Сразу же стало понятно, что до появления нового флагмана компании NVIDIA именно эти видеокарты являются самими производительными из всех однопроцессорных графических адаптеров.

По сравнению со своими предшественниками — ATI Radeon HD 4870 (которые сами являются весьма удачными продуктами компании AMD) – видеокарты ATI Radeon HD 5870 обладают вдвое более высокой производительностью.

Причина столь впечатляющего результата кроется в двукратном увеличении числа основных вычислительных блоков графического процессора: потоковых процессоров, текстурных модулей, блоков растеризации и пр.

Модернизация видеочипов именно в этом направлении привела к вполне ожидаемому резкому увеличению и количества транзисторов, из которых состоят интегральные схемы – если конструкция процессора ATI Radeon HD 4870 предусматривала использование 0,956 млрд транзисторов, то в случае ATI Radeon HD 5870 их количество составило уже 2,15 млрд. Но даже эта астрономическая цифра меркнет перед новым продуктом NVIDIA, анонсированного представителями компании в ходе конференции GPU Technologies Conference – графическим процессором NVIDIA Fermi.

Разработчики из Калифорнии основательно подошли к созданию графического процессора нового поколения – микрочипы Fermi (кстати, их «старое» кодовое обозначение – GT300) состоят из более чем трех миллиардов транзисторов.

Эта цифра сразу на 40% выше количества транзисторов у процессора RV870, при том, что они являются продуктами одного поколения и изготавливаются по одному технологическому процессу – 40-нм на мощностях тайваньской компании TSMC.

Если сравнивать процессоры Fermi с решениями предыдущего поколения GT200, то преимущество в количестве транзисторов и вовсе двукратное – 1,4 млрд против 3 млрд.

Увеличение количества транзисторов вполне предсказуемо сказалось на характеристиках процессоров: по сравнению с GT200 увеличено до 512 количество вычислительных блоков, увеличена до 384 бит разрядность интерфейса графической памяти (шесть 64-разрядных блоков), реализована поддержка памяти стандарта GDDR5 максимальным объемом аж до 6 Гб. Надеемся, что все читатели помнят о поддержке процессорами GT200 стандарта GDDR3, а значит, переход на более скоростную графическую память позволит заметно увеличить возможности соответствующей подсистемы видеокарт. Ожидалась и аппаратная поддержка DirectX 11, которая была реализована разработчиками. Но на этом сюрпризы не заканчиваются, ведь дополнительные транзисторы «потрачены» не только на простое увеличение количества исполнительных блоков, как сделали инженеры AMD/ATI. В отличие от их продукта, графический процессор имеет заметно переработанную архитектуру, в которой реализовано значительное количество интересных и передовых (для графических процессоров как класса) нововведений.

На «верхнем» уровне архитектуры графических процессоров существенных качественных отличий не наблюдается. С этой позиции Fermi можно рассматривать всего лишь как масштабированную версию графических процессоров GT200.

Но как только мы обращаем свое внимание на «нижние» уровни архитектуры, на ее фундамент, то сразу же появляются существенные нововведения инженеров NVIDIA. Первое на что необходимо обращать внимание – на графические ядра, которые ранее обозначались как потоковые процессоры (SP — Streaming Processor).

Сейчас разработчики вместо привычного термина перешли на использования термина CUDA-ядер (CUDA Core). В случае графических интегральных микросхем G80 и GT200 разработчики объединяли по восемь потоковых процессоров в единые группы – потоковые мультипроцессорные SM-блоки (SM — Streaming Multiprocessor).

Похожая организация сохранена и для процессоров Fermi, за тем лишь исключением, что теперь в единый блок объединяются не восемь, а тридцать два ядра.

В зависимости от конкретной реализации, в основу графических процессоров NVIDIA Fermi будут входить до шестнадцати SM-блоков, состоящих из 32 CUDA-ядер, оптимизированных для работы с вычислениями общего назначения. В результате и получаются 512 вычислительных ядра, которые и являются базой для высочайшей производительности процессоров.

В дополнение к упомянутым CUDA-ядрам, в состав мультипотоковых блоков входят и дополнительные вычислительные элементы. Речь идет о блоках Special Function Unit (SFU), основная область применения которых — трансцендентальная математика и интерполяция.

Впрочем, блоки SFU сложно назвать главными исполнительными элементами графического процессора. На это указывает и тот факт, что разработчики решили лишь удвоить количество этих компонентов для каждого из SM-блоков графического процессора – до четырех штук.

К сожалению, более подробную информацию об основных исполнительных компонентах архитектуры NVIDIA Fermi разработчики пока предпочли не разглашать.

В данном случае за кадром остались такие элементы, как блоки растеризации, текстурной фильтрации, и многое другое, что непосредственно касается возможностей процессоров по обработке трехмерной графики.

Поэтому пока практически невозможно сказать, насколько интереснее архитектура Fermi применительно к компьютерным играм по сравнению со своими предшественниками в лице GT200.

Помимо исполнительных блоков, отвечающих непосредственно за обработку информации, каждый из SM-блоков графического процессора на основе архитектуры NVIDIA Fermi оснащается еще и блоками временного хранения данных – кэш-память.

В случае видеочипов предыдущих поколений разработчики также оснащали SM-блоки кэш-памятью первого и второго уровней фиксированным объемом 24 Кб и 256 Кб соответственно. Но использовалась она лишь для хранения «текстурных» данных, к тому же, блок кэш-памяти уровня L1 распределялся между тремя потоковыми мультипроцессорными блоками.

Графические процессоры следующего поколения оснащены уже универсальной кэш-памятью, причем каждому SM-блоку соответствует собственный блок кэш-памяти первого уровня. Но что самое важное, эта память является конфигурируемой.

Что это означает? Каждый SM-блок имеет доступ к 64 Кб памяти, которая может быть разделена непосредственно на кэш-память первого уровня и разделяемую память, причем разделена двумя способами: 48 Кб/16 Кб, либо 16 Кб/48 Кб. Такой подход позволяет решить сразу несколько задач.

Во-первых, возможность различной конфигурации кэш и разделяемой памяти означает полную совместимость с приложениями, оптимизированными для работы с графическими процессорами GT200 с их памятью объемом 16 Кб.

Во-вторых, при выполнении вычислений общего назначения, специализированная традиционная «текстурная» кэш-память оказывается неэффективной – эта проблема решена в случае Fermi с ее универсальной кэш-памятью первого и второго уровней.

В-третьих, возможность различной конфигурации кэш-памяти позволяет с максимальной эффективностью организовать работу программного обеспечения, оптимизированного для работы с большим объемом кэш-памяти – до 48 Кб. И последнее нововведение – увеличение объема универсальной кэш-памяти второго уровня до 768 Кб, которое позволяет существенно повысить производительность при работе с так называемыми атомарными операциями, часто используемыми при вычислениях общего назначения. Согласно заверениям разработчиков, эффективность работы повышается по сравнению с GT200 в 4 – 20 раз (!).