Azərbaycanca
Беларускі
Dansk
Deutsch
Española
Français
Indonesia
Italiana
日本語
Қазақ
Lietuvos
Nederlands
Português
Русский
සිංහල
แบบไทย
Türkçe
Українська
中國人
United State
Afrikaans
Fermi GPU мікроархітектура, розроблена Nvidia як наступник мікроархітектури Tesla. Ця мікроархітектура використовується в серії GeForce GeForce 400 і GeForce 500 серії. На ринку Fermi знайшов застосування в X000 серії Quadro, моделі Quadro NVS, а також в Nvidia Tesla обчислювальних модулів. Fermi одна з найстаріших мікроархітектур NVIDIA, що може підтримувати DirectX 12.
Огляд
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".
Вікіпедія, Українська, Україна, книга, книги, бібліотека, стаття, читати, завантажити, безкоштовно, безкоштовно завантажити, mp3, відео, mp4, 3gp, jpg, jpeg, gif, png, малюнок, музика, пісня, фільм, книга, гра, ігри, мобільний, телефон, android, ios, apple, мобільний телефон, samsung, iphone, xiomi, xiaomi, redmi, honor, oppo, nokia, sonya, mi, ПК, web, Інтернет
Fermi GPU mikroarhitektura rozroblena Nvidia yak nastupnik mikroarhitekturi Tesla Cya mikroarhitektura vikoristovuyetsya v seriyi GeForce GeForce 400 i GeForce 500 seriyi Na rinku Fermi znajshov zastosuvannya v X000 seriyi Quadro modeli Quadro NVS a takozh v Nvidia Tesla obchislyuvalnih moduliv Fermi odna z najstarishih mikroarhitektur NVIDIA sho mozhe pidtrimuvati DirectX 12 OglyadRis 1 Arhitektura Fermi Konvenciya v cifrah pomaranchevij planuvannya ta dispetcherizaciya zelenij vikonannya blakitni registri i keshi Fermi Graphic Processing Units maye 3 milyardi tranzistoriv i shematichno zobrazhena na ris 1 Potokovij mulitiprocesor skladayetsya iz 32 yader CUDA Globalnij potokovij planuvalnik rozpodilyaye bloki potokiv SM planuvalnikiv i upravlyaye peremikannyam kontekstu mizh potokami pid chas vikonannya Host interfejs z yednuye GPU do CPU cherez PCI Express V2 shini pikova shvidkost peredachi danih 8GB s DRAM pidtrimuyetsya do 6 GB GDDR5 DRAM pam yati zavdyaki v 64 bitnoyi adresaciyi Taktova chastota 1 5 GGc Ne vipusheni NVIDIA ale ocinyuyutsya za Insight 64 Pikova produktivnist 1 5 Tflops Globalna chastota pam yati 2 GGc Propuskna zdatnist DRAM 192Gb sPotokovist multiprocesoraKozhen multiprocesor mistit 32 odinarnoyi tochnosti CUDA yader 16 navantazhennya zberezhennya blokiv chotiri specialni funkcionalni bloki SFUs i 64K visokoshvidkisnoyi pam yati na kristali i interfejs do kesh pam yati drugogo rivnya Zavantazhennya Zberezhennya bloku adresi dzherela i priznachennya povinni buti rozrahovani dlya 16 potokiv za takt Zavantazhennya ta zberigannya danih z abo v kesh abo DRAM Specialni funkciyi blokiv SFUs vikonuyut transcendentni operaciyi taki yak sin cos kvadratnij korin Kozhna funkciya vikonuye odnu operaciyu na odnomu potoci za takt SFUs vidokremleni vid dispetcherskogo pristroyu sho dozvolyaye dispetcherskomu upravlinnyu vidati vikonannya inshoyi funkciyi poki SFUs zajnyato Yadro CUDAArifmetiko logichnij pristrij ALP Pidtrimuye povnu 32 bitnu tochnist dlya vsih operacij vidpovidno do vimog standartu movi programuvannya ALP takozh optimizovanij dlya efektivnoyi pidtrimki 64 bitnih i pidvishenoyi tochnosti operacij Modul operacij z plavayuchoyu komoyu realizovanij na novomu IEEE 754 2008 z plavayuchoyu komoyu standartu zabezpechuyuchij mnozhennya dodavannya operacij z plavayuchoyu komoyu dlya odinarnoyi i podvijnoyi tochnosti Do 16 podvijnoyi tochnosti mnozhennya dodavannya operaciyi mozhut buti vikonani u vidpovidnosti z multiprocesorom za takt Plavayuche mnozhennya dodavannyaPlavayuche mnozhennya dodavannya vikonuyetsya tobto A B C z odnogo kincevogo kroku okruglennya bez vtrati tochnosti v rezultati plavayuche mnozhennya dodavannya ye bilsh tochnim nizh vikonannya operacij okremo Planuvannya varpivArhitektura fermi maye dva rivnya planuvannya potokiv Kozhen multiprocesor mozhe vikoristovuvati instrukciyi bud yaki dvi z chotiroh zelenih kolonok vikonannya pokazanih na shematichnomu ris 1 Napriklad multiprocesor mozhete peremishati 16 operacij z 16 yadrami pershih stovpciv z 16 operacij z drugogo yadra 16 stovpciv abo 16 operacij iz zavantazhennya zberezhennya z chotirma z SFUs abo bud yakih inshih kombinacij sho viznachaye programa Zvernit uvagu sho 64 bitovi operaciyi z plavayuchoyu tochkoyu vikoristovuye yak pershi dvi kolonki vikonannya Ce oznachaye sho multiprocesor mozhe vidavati do 32 odinarnoyi tochnosti 32 bit operacij z plavayuchoyu komoyu abo 16 z podvijnoyu tochnistyu 64 rozryadne i plavayuchoyu komoyu v toj chas Podvijne planuvannya varpiv Na rivni multiprocesora kozhen varp maye 32 potokiv dlya svoyih vikonavchih blokiv Potoki skladayut v grupi po 32 potoki tak zvanih varpom Kozhen multiprocesor maye dva riznih rivnya planuvannya potokiv ta dva instrukciya vidpravki bloku sho dozvolyaye dvom varpam buti zapusheno i vikonano odnochasno Podvijne planuvannya varpiv vibiraye dva varpi i vikonuye odnu komandu z kozhnogo varpu do grupi iz 16 yadrami 16 zavantazhennya zberezhennya blokiv abo 4 SFUs Bilshist instrukcij mozhe buti podvijno vikonana dva cilih instrukciyi dvi plavayuchih instrukciyi abo sumish ciloyi i z plavayuchoyu zavantazhennya zberezhennya ta instrukciyi SFU mozhe buti vikonano odnochasno Podvijna tochnist instrukcij ne pidtrimuye podvijnu vidpravku z bud yakoyu inshoyu operaciyeyu Pam yatKesh pershogo rivnya na multiprocesori i yedinij kesh drugogo rivnya yakij obslugovuye vsi operaciyi zavantazhennya zberezhennya i teksturi Registri Kozhen multiprocesor maye 32K 32 rozryadnih registriv Kozhen potik maye dostup do vlasnih registriv a ne do inshih potokiv Maksimalne chislo registriv yaki mozhut buti vikoristani v yadra CUDA ye 63 Kilkist dostupnih registriv pogirshuyetsya plavno vid 63 do 21 Registri mayut duzhe visoku propusknu zdatnist blizko 8000 Gbajt s L1 Shared Memory Na chipi pam yati yaki mozhut buti vikoristani abo dlya keshuvannya danih dlya okremih potokiv abo dlya obminu danimi mizh dekilkoma potokami podilennya pam yati Ci 64 Kb pam yati mozhut buti nalashtovani yak 48 KB spilnoyi pam yati z 16 KB kesh pam yati L1 abo 16 KB spilnoyi pam yati z 48 KB kesh pam yati L1 Zagalna pam yat dozvolyaye potoku v odnomu bloci spivpracyuvati povtornomu vikoristanyu na chipi danih i znachno zmenshuye vid chipa trafik Spilna pam yat dostupna potokam v odnomu bloci potoku Ce zabezpechuye nizkoyu zatrimkoyu dostupu 10 20 cikliv i duzhe visoku propusknu zdatnist 1600 Gb s shob zmenshiti obsyagi danih napriklad promizhnih rezultativ v seriyi rozrahunkiv odin ryadok abo stovpec danih dlya matrichnih operacij tosho Devid Patterson kazhe sho ce Shared Memory vikoristovuye ideyu lokalnogo bloknotu Lokalna pam yat Lokalna pam yat priznachena yak misce pam yati vikoristovuvanoyi dlya zberigannya prolittya registriv Reyestraciya prolittya vidbuvayetsya koli potik bloku potribno bilshe pam yati nizh dostupnoyi na multiprocesori Lokalna pam yat vikoristovuyetsya tilki dlya deyakih avtomatichnih zminnih yaki ogolosheni v kodi bez bud yakih device shared abo constant Yak pravilo avtomatichna zminna znahoditsya v registri za vinyatkom takih 1 Masivi sho kompilyator ne mozhe viznachiti chi indeksuyutsya z postijnimi indeksami 2 Veliki strukturi abo masivi yaki spozhivayut zanadto bagato registriv prostir Bud yaku zminnu kompilyator virishuye dodati v lokalnu pam yat koli yadro vikoristovuye bilshe registriv nizh ye na multiprocesori Kesh drugogo rivnya 768 Kb rozpodilyayutsya sered 16 OM sho opracovuyut vsi navantazhennya i zberezhennya z v globalnij pam yati v tomu chisli kopij z centralnogo procesora a takozh teksturnih zapitiv Kesh pam yati L2 takozh realizuye atomarni operaciyi vikoristovuvani dlya upravlinnya dostupom do danih yaki povinni buti zagalnimi potokami blokiv abo navit yader Globalna pam yat Zruchnij dostup dlya vsih potokiv a takozh CPU Visoka latentnist 400 800 cikliv Versiyi FermiGF100 GF104 GF106 GF110 GF114 GF116 GF118 GF119 GF117Posilannya NVIDIA s Next Generation CUDA Compute Architecture Fermi Arhivovano 9 kvitnya 2012 u WebCite N Brookwood NVIDIA Solves the GPU Computing Puzzle 3 bereznya 2016 u Wayback Machine P N Glaskowsky NVIDIA s Fermi The First Complete GPU Computing Architecture 21 lipnya 2021 u Wayback Machine N Whitehead A Fit Florea Precision amp Performance Floating Point and IEEE 754 Compliance for NVIDIA GPUs 20 kvitnya 2015 u 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