Fermi (мікроархітектура)

Fermi GPU мікроархітектура, розроблена Nvidia як наступник мікроархітектури Tesla. Ця мікроархітектура використовується в серії GeForce GeForce 400 і GeForce 500 серії. На ринку Fermi знайшов застосування в X000 серії Quadro, моделі Quadro NVS, а також в Nvidia Tesla обчислювальних модулів. Fermi одна з найстаріших мікроархітектур NVIDIA, що може підтримувати DirectX 12.

Огляд

Рис.1. Архітектура Fermi. Конвенція в цифрах: помаранчевий — планування та диспетчеризація; зелений — виконання; блакитні — регістри і кеші.

Fermi Graphic Processing Units має 3 мільярди транзисторів і схематично зображена на рис.1.

• Потоковий мулитипроцесор складається із 32 ядер CUDA .
• Глобальний потоковий планувальник розподіляє блоки потоків SM планувальників і управляє перемиканням контексту між потоками під час виконання.
• Хост інтерфейс: з'єднує GPU до CPU через PCI-Express V2 шини (пікова швидкость передачі даних 8GB / с).
• DRAM: підтримується до 6 ГБ GDDR5 DRAM пам'яті завдяки в 64-бітної адресації.
• Тактова частота: 1,5 ГГц (Не випущені NVIDIA, але оцінюються за Insight 64).
• Пікова продуктивність: 1,5 Тфлопс.
• Глобальна частота пам'яті: 2 ГГц.
• Пропускна здатність DRAM: 192Gb / s

Потоковість мультипроцесора

Кожен мультипроцесор містить 32 одинарної точності CUDA ядер, 16 навантаження / збереження блоків, чотири спеціальні функціональні блоки (SFUs), і 64К високошвидкісної пам'яті на кристалі  і інтерфейс до кеш-пам'яті другого рівня. Завантаження / Збереження блоку:  адреси джерела і призначення повинні бути розраховані для 16 потоків за такт. Завантаження та зберігання даних з або в кеш або DRAM.

Спеціальні функції блоків (SFUs): виконують трансцендентні операції, такі як sin, cos, квадратний корінь. Кожна функція виконує одну операцію на одному потоці за такт. SFUs відокремлені від диспетчерського пристрою, що дозволяє диспетчерському управлінню видати виконання іншої функції, поки SFUs зайнято.

Ядро CUDA

Арифметико-логічний пристрій (АЛП): Підтримує повну 32-бітну точність для всіх операцій, відповідно до вимог стандарту мови програмування. АЛП також оптимізований для ефективної підтримки 64-бітних і підвищеної точності операцій.

Модуль операцій з плаваючою комою – реалізований на новому IEEE 754-2008 з плаваючою комою стандарту, забезпечуючий множення-додавання операцій з плаваючою комою, для одинарної і подвійної точності.

До 16 подвійної точності множення-додавання операції можуть бути виконані у відповідності з мультипроцесором за такт.

Плаваюче множення-додавання

Плаваюче множення-додавання виконується (тобто, А * B + C) з одного кінцевого кроку округлення без втрати точності в результаті. плаваюче множення-додавання є більш точним, ніж виконання операцій окремо.

Планування варпів

Архітектура фермі має два рівня планування потоків. Кожен мультипроцесор може використовувати інструкції будь-які дві з чотирьох зелених колонок виконання, показаних на схематичному рис. 1. Наприклад, мультипроцесор можете перемішати 16 операцій з 16 ядрами перших стовпців з 16 операцій з другого ядра 16 стовпців, або 16 операцій із завантаження / збереження з чотирма з SFUs, або будь-яких інших комбінацій, що визначає програма.

Зверніть увагу, що 64-бітові операції з плаваючою точкою використовує як перші дві колонки виконання. Це означає, що мультипроцесор може видавати до 32 одинарної точності (32-біт) операцій з плаваючою комою або 16 з подвійною точністю (64-розрядне) і плаваючою комою в той час.

Подвійне планування варпів: На рівні мультипроцесора, кожен варп має 32 потоків для своїх виконавчих блоків. Потоки складають в групи по 32 потоки, так званих варпом. Кожен мультипроцесор має два різних рівня планування потоків та два інструкція відправки блоку, що дозволяє двом варпам бути запущено і виконано одночасно. Подвійне планування варпів вибирає два варпи і виконує одну команду з кожного варпу до групи із 16 ядрами,16 завантаження / збереження блоків, або 4 SFUs. Більшість інструкцій може бути подвійно виконана; два цілих інструкції, дві плаваючих інструкції, або суміш цілої і з плаваючою, завантаження, збереження та інструкції SFU може бути виконано одночасно. Подвійна точність інструкцій не підтримує подвійну відправку з будь-якою іншою операцією.

Пам’ять

Кеш першого рівня на мультипроцесорі і єдиний кеш другого рівня, який обслуговує всі операції (завантаження, збереження і текстури). Регістри: Кожен мультипроцесор має 32К 32-розрядних регістрів. Кожен потік має доступ до власних регістрів а не до інших потоків. Максимальне число регістрів, які можуть бути використані в ядра CUDA є 63. Кількість доступних регістрів погіршується плавно від 63 до 21, Регістри мають дуже високу пропускну здатність: близько 8000 Гбайт / с.

L1 + Shared Memory: На чипі пам'яті, які можуть бути використані або для кешування даних для окремих потоків або для обміну даними між декількома потоками (поділення пам'яті). Ці 64 Кб пам'яті можуть бути налаштовані як 48 КБ спільної пам'яті з 16 КБ кеш-пам'яті L1, або 16 КБ спільної пам'яті з 48 КБ кеш-пам'яті L1. Загальна пам'ять дозволяє потоку в одному блоці співпрацювати повторному використаню на чипі даних,і значно зменшує від чипа трафік. Спільна пам'ять доступна потокам в одному блоці потоку. Це забезпечує низькою затримкою доступу (10-20 циклів) і дуже високу пропускну здатність (1600 Гб / с), щоб зменшити обсяги даних (наприклад, проміжних результатів в серії розрахунків, один рядок або стовпець даних для матричних операцій, тощо). Девід Паттерсон каже, що це Shared Memory використовує ідею локального блокноту.

Локальна пам'ять: Локальна пам'ять призначена як місце пам'яті, використовуваної для зберігання "пролиття" регістрів. Реєстрація пролиття відбувається, коли потік блоку потрібно більше пам'яті, ніж доступної на мультипроцесорі. Локальна пам'ять використовується тільки для деяких автоматичних змінних (які оголошені в коді без будь-яких __device__, __shared__ або __constant__). Як правило, автоматична змінна знаходиться в регістрі, за винятком таких: (1) Масиви, що компілятор не може визначити чи індексуються з постійними індексами; (2) Великі структури або масиви, які споживають занадто багато регістрів простір; Будь-яку змінну компілятор вирішує додати в локальну пам'ять, коли ядро використовує більше регістрів, ніж є на мультипроцесорі.

Кеш другого рівня: 768 Кб , розподіляються серед 16 ОМ, що опрацьовують всі навантаження і збереження з / в глобальній пам'яті, в тому числі копій з центрального процесора, а також текстурних запитів. Кеш-пам'яті L2, також реалізує атомарні операції, використовувані для управління доступом до даних, які повинні бути загальними потоками блоків або навіть ядер.

Глобальна пам'ять: Зручний доступ для всіх потоків, а також CPU. Висока латентність (400-800 циклів)

Версії Фермі

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

Посилання

• "NVIDIA’s Next Generation CUDA Compute Architecture: Fermi." [Архівовано 9 квітня 2012 у WebCite]
• N. Brookwood, "NVIDIA Solves the GPU Computing Puzzle." [Архівовано 3 березня 2016 у Wayback Machine.]
• P.N. Glaskowsky, "NVIDIA’s Fermi: The First Complete GPU Computing Architecture." [Архівовано 21 липня 2021 у Wayback Machine.]
• N. Whitehead, A. Fit-Florea, "Precision & Performance: Floating Point and IEEE 754 Compliance for NVIDIA GPUs." [Архівовано 20 квітня 2015 у Wayback Machine.], 2011.
• S.F. Oberman, M. Siu, "A high-performance area-efficient multifunction interpolator," Proc. of the 17th IEEE Symposium on Computer Arithmetic, Cap Cod, MA, USA, Jul. 27-29, 2005, pp. 272–279.
• R. Farber, "CUDA Application Design and Development," Morgan Kaufmann, 2011.
• NVIDIA Application Note "Tuning CUDA applications for Fermi".