Įrenginiai, skirti asmeniniams kompiuteriams paversti mažais superkompiuteriais, buvo žinomi gana ilgą laiką. Dar praėjusio amžiaus 80-aisiais rinkoje buvo siūlomi vadinamieji transputeriai, kurie buvo įkišti į tuo metu įprastus ISA išplėtimo lizdus. Iš pradžių jų atlikimas atitinkamose užduotyse buvo įspūdingas, tačiau paskui paspartėjo universalių procesorių spartos augimas, jie sustiprino savo pozicijas lygiagrečiame skaičiavime, o transuteriams nebuvo prasmės. Nors panašūs įrenginiai egzistuoja ir šiandien – tai įvairūs specializuoti greitintuvai. Tačiau dažnai jų taikymo sritis yra siaura ir tokie greitintuvai nėra ypač paplitę.

Tačiau pastaruoju metu lygiagretaus skaičiavimo lazdelė perėjo į masinę rinką, vienaip ar kitaip susijusią su 3D žaidimais. Bendrosios paskirties įrenginiai su kelių branduolių procesoriais lygiagrečiam vektorių skaičiavimui, naudojamiems 3D grafikoje, pasiekia didelį našumą, kurio negali prilygti bendrosios paskirties procesoriai. Žinoma, maksimalus greitis pasiekiamas tik atliekant daugybę patogių užduočių ir turi tam tikrų apribojimų, tačiau tokie įrenginiai jau pradėti gana plačiai naudoti tose srityse, kurioms jie iš pradžių nebuvo skirti. Puikus tokio lygiagretaus procesoriaus pavyzdys yra Sony-Toshiba-IBM aljanso sukurtas procesorius Cell, naudojamas Sony PlayStation 3 žaidimų konsolėje, taip pat visos šiuolaikinės vaizdo plokštės iš rinkos lyderių – Nvidia ir AMD.

Šiandien mes neliesime Cell, nors jis pasirodė anksčiau ir yra universalus procesorius su papildomomis vektorinėmis galimybėmis, šiandien apie tai nekalbame. 3D vaizdo greitintuvams prieš keletą metų pasirodė pirmosios negrafinių bendrosios paskirties skaičiavimų technologijos GPGPU (General-Purpose computation on GPUs). Galų gale, šiuolaikiniuose vaizdo lustuose yra šimtai matematinių vykdymo vienetų, o ši galia gali būti naudojama norint žymiai pagreitinti daugelį daug skaičiavimo reikalaujančių programų. O dabartinės GPU kartos turi gana lanksčią architektūrą, kuri kartu su aukšto lygio programavimo kalbomis ir aparatinės-programinės įrangos architektūromis, kaip aptarta šiame straipsnyje, atskleidžia šias galimybes ir daro jas daug labiau prieinamas.

Sukurti GPCPU paskatino gana greitų ir lanksčių šešėlių programų atsiradimas, kurias gali vykdyti šiuolaikiniai vaizdo lustai. Kūrėjai nusprendė, kad GPU skaičiuotų ne tik vaizdus 3D programose, bet ir būtų naudojami atliekant kitus lygiagrečius skaičiavimus. GPGPU tam buvo naudojamos grafinės API: OpenGL ir Direct3D, kai duomenys į vaizdo lustą buvo perduodami tekstūrų pavidalu, o skaičiavimo programos buvo įkeliamos šešėlių pavidalu. Šio metodo trūkumai yra gana didelis programavimo sudėtingumas, mažas duomenų mainų tarp procesoriaus ir GPU greitis ir kiti apribojimai, kuriuos aptarsime vėliau.

GPU kompiuterija vystėsi ir vystosi labai greitai. Vėliau du pagrindiniai vaizdo lustų gamintojai, Nvidia ir AMD, sukūrė ir paskelbė apie atitinkamas platformas, atitinkamai vadinamas CUDA (Compute Unified Device Architecture) ir CTM (Close To Metal arba AMD Stream Computing). Skirtingai nuo ankstesnių GPU programavimo modelių, jie buvo sukurti su tiesiogine prieiga prie vaizdo plokščių techninės įrangos galimybių. Platformos nesuderinamos viena su kita, CUDA yra C programavimo kalbos plėtinys, o CTM yra virtuali mašina, kuri vykdo surinkimo kodą. Tačiau abi platformos pašalino kai kuriuos svarbius ankstesnių GPGPU modelių, kuriuose buvo naudojamas tradicinis grafikos vamzdynas ir atitinkamos Direct3D arba OpenGL sąsajos, apribojimus.

Žinoma, atvirieji standartai, naudojantys OpenGL, atrodo patys nešiojamiausi ir universaliausi, leidžiantys tą patį kodą naudoti skirtingų gamintojų vaizdo lustuose. Tačiau tokie metodai turi daug trūkumų, jie yra daug mažiau lankstūs ir nėra tokie patogūs naudoti. Be to, jie neleidžia naudoti specifinių tam tikrų vaizdo plokščių galimybių, pavyzdžiui, sparčiai bendrinamos (bendros) atminties, esančios šiuolaikiniuose skaičiavimo procesoriuose.

Štai kodėl „Nvidia“ išleido CUDA platformą – į C panašią programavimo kalbą su savo kompiliatoriumi ir bibliotekomis, skirtomis GPU skaičiavimui. Žinoma, parašyti optimalų kodą vaizdo lustams nėra taip paprasta ir ši užduotis reikalauja daug rankų darbo, tačiau CUDA atskleidžia visas galimybes ir suteikia programuotojui daugiau galimybių valdyti GPU aparatines galimybes. Svarbu, kad „Nvidia CUDA“ palaikymas yra G8x, G9x ir GT2xx lustuose, naudojamuose „Geforce 8“, 9 ir 200 serijų vaizdo plokštėse, kurios yra labai paplitusios. Dabar buvo išleista galutinė CUDA 2.0 versija, kurioje yra keletas naujų funkcijų, pavyzdžiui, dvigubo tikslumo skaičiavimų palaikymas. CUDA galima naudoti 32 bitų ir 64 bitų Linux, Windows ir MacOS X operacinėse sistemose.

Skirtumas tarp procesoriaus ir GPU lygiagrečių skaičiavimų metu

Universalių procesorių dažnių padidėjimas susidūrė su fiziniais apribojimais ir dideliu energijos suvartojimu, o jų našumas vis labiau didinamas į vieną lustą dedant kelis branduolius. Šiuo metu parduodamuose procesoriuose yra tik iki keturių branduolių (tolimesnis augimas nebus greitas) ir jie skirti bendroms programoms, naudojant MIMD kelių komandų ir duomenų srautą. Kiekvienas branduolys veikia atskirai nuo kitų, vykdydamas skirtingas instrukcijas skirtingiems procesams.

Specializuotos vektorinės galimybės (SSE2 ir SSE3) keturių komponentų (vieno tikslumo slankiojo kablelio) ir dviejų komponentų (dvigubo tikslumo) vektoriams atsirado bendrosios paskirties procesoriuose, visų pirma dėl padidėjusių grafikos programų poreikių. Štai kodėl tam tikroms užduotims atlikti GPU naudojimas yra pelningesnis, nes iš pradžių jie buvo sukurti jiems.

Pavyzdžiui, „Nvidia“ vaizdo lustuose pagrindinis įrenginys yra daugiaprocesorius, turintis aštuonių iki dešimties branduolių ir iš viso šimtus ALU, kelis tūkstančius registrų ir nedaug bendros atminties. Be to, vaizdo plokštėje yra greita visuotinė atmintis, kurią gali pasiekti visi daugiaprocesoriai, vietinė atmintis kiekviename daugiaprocesoriuje ir speciali atmintis konstantoms.

Svarbiausia, kad šie keli kelių procesorių GPU branduoliai yra SIMD (vieno instrukcijų srauto, kelių duomenų srauto) branduoliai. Ir šie branduoliai vienu metu vykdo tas pačias instrukcijas – toks programavimo stilius, kuris yra įprastas grafiniams algoritmams ir daugeliui mokslinių užduočių, tačiau reikalauja specifinio programavimo. Tačiau šis metodas leidžia padidinti vykdymo vienetų skaičių dėl jų supaprastinimo.

Taigi, išvardinkime pagrindinius CPU ir GPU architektūros skirtumus. CPU branduoliai skirti vykdyti vieną nuoseklių komandų srautą esant maksimaliam našumui, o GPU sukurti taip, kad greitai vykdytų daug lygiagrečių komandų srautų. Bendrosios paskirties procesoriai yra optimizuoti, kad būtų pasiektas didelis našumas naudojant vieną komandų srautą, apdorojant sveikuosius ir slankiojo kablelio skaičius. Šiuo atveju prieiga prie atminties yra atsitiktinė.

CPU kūrėjai stengiasi lygiagrečiai vykdyti kuo daugiau instrukcijų, kad padidintų našumą. Tam, pradedant Intel Pentium procesoriais, atsirado superskaliarinis vykdymas, užtikrinantis dviejų instrukcijų vykdymą per laikrodžio ciklą, o Pentium Pro pasižymėjo netvarkingu instrukcijų vykdymu. Tačiau lygiagretus nuoseklaus komandų srauto vykdymas turi tam tikrų pagrindinių apribojimų, o vykdymo vienetų skaičiaus padidinimas negali pasiekti daugkartinio greičio padidėjimo.

Vaizdo lustai nuo pat pradžių veikia paprastai ir lygiagrečiai. Vaizdo lustas paima daugiakampių grupę kaip įvestį, atlieka visas būtinas operacijas ir sukuria pikselius kaip išvestį. Daugiakampių ir pikselių apdorojimas yra nepriklausomas, jie gali būti apdorojami lygiagrečiai, atskirai vienas nuo kito. Todėl dėl iš esmės lygiagrečio darbo organizavimo GPU naudoja daug vykdymo vienetų, kuriuos lengva įkelti, priešingai nei nuoseklus procesoriaus instrukcijų srautas. Be to, šiuolaikiniai GPU taip pat gali vykdyti daugiau nei vieną komandą per laikrodžio ciklą (dviguba problema). Taigi, Tesla architektūra, esant tam tikroms sąlygoms, vienu metu paleidžia MAD+MUL arba MAD+SFU operacijas.

GPU taip pat skiriasi nuo procesoriaus atminties prieigos principais. GPU jis yra nuoseklus ir lengvai nuspėjamas - jei tekstūros tekselis skaitomas iš atminties, tada po kurio laiko ateis laikas kaimyniniams teksams. O įrašant nutinka tas pats – į kadrų buferį įrašomas pikselis, o po kelių laikrodžio ciklų bus įrašomas šalia esantis. Todėl atminties organizavimas skiriasi nuo naudojamo procesoriaus. O vaizdo lustui, skirtingai nei universaliems procesoriams, tiesiog nereikia didelės talpyklos atminties, o tekstūroms reikia vos kelių (dabartiniuose GPU iki 128-256) kilobaitų.

O pats darbas su atmintimi GPU ir CPU kiek skiriasi. Taigi, ne visi centriniai procesoriai turi įmontuotus atminties valdiklius, o visi GPU paprastai turi kelis valdiklius, iki aštuonių 64 bitų kanalų Nvidia GT200 mikroschemoje. Be to, vaizdo plokštės naudoja greitesnę atmintį, todėl vaizdo lustai turi prieigą prie žymiai didesnio atminties pralaidumo, o tai taip pat labai svarbu atliekant lygiagrečius skaičiavimus, kurie veikia didžiuliais duomenų srautais.

Universaliuose procesoriuose daug tranzistorių ir lusto ploto naudojami instrukcijų buferiams, aparatinės įrangos šakų numatymui ir dideliam lustinės talpyklos atminties kiekiui. Visi šie aparatūros blokai reikalingi norint pagreitinti kelių komandų srautų vykdymą. Vaizdo lustai išleidžia tranzistorius vykdymo blokų, srauto valdymo blokų, mažos bendros atminties ir kelių kanalų atminties valdiklių masyvuose. Tai, kas išdėstyta aukščiau, nepagreitina atskirų gijų vykdymo, tai leidžia lustui apdoroti kelis tūkstančius gijų, kurias vienu metu vykdo lustas ir kurioms reikalingas didelis atminties pralaidumas.

Apie talpyklos skirtumus. Bendrosios paskirties procesoriai naudoja talpyklą, kad padidintų našumą, sumažindami prieigos prie atminties delsą, o GPU naudoja talpyklą arba bendrinamą atmintį, kad padidintų pralaidumą. CPU sumažina prieigos prie atminties delsą naudodami dideles talpyklas ir kodo šakų numatymą. Šios techninės įrangos dalys užima didžiąją dalį lusto ploto ir sunaudoja daug energijos. Vaizdo lustai apeina atminties prieigos vėlavimo problemą, vienu metu vykdydami tūkstančius gijų – kol viena iš gijų laukia duomenų iš atminties, vaizdo lustas gali atlikti skaičiavimus pagal kitą giją, nelaukdamas ir nedelsdamas.

Taip pat yra daug skirtumų, susijusių su kelių sriegių palaikymu. Centrinis procesorius atlieka 1–2 skaičiavimų gijas viename procesoriaus branduolyje, o vaizdo lustai gali palaikyti iki 1024 gijų viename daugiaprocesoriuje, kurių mikroschemoje yra keletas. Ir jei perjungimas iš vienos gijos į kitą CPU kainuoja šimtus laikrodžio ciklų, tai GPU perjungia keletą gijų per vieną laikrodžio ciklą.

Be to, centriniai procesoriai naudoja SIMD (viena instrukcija per kelis duomenis) vektoriaus skaičiavimams, o GPU naudoja SIMT (vienos instrukcijos kelių gijų) skaliarinių gijų apdorojimui. SIMT nereikalauja, kad kūrėjas konvertuotų duomenis į vektorius ir leidžia savavališkai išsišakoti srautuose.

Trumpai tariant, galime pasakyti, kad, skirtingai nuo šiuolaikinių universalių procesorių, vaizdo lustai yra skirti lygiagrečiam skaičiavimui su daugybe aritmetinių operacijų. O žymiai didesnis skaičius GPU tranzistorių dirba pagal numatytą paskirtį – apdoroja duomenų masyvus, ir nekontroliuoja kelių nuoseklių skaičiavimo gijų vykdymo (srauto valdymo). Tai diagrama, kiek vietos CPU ir GPU užima įvairios logikos:

Todėl efektyvaus GPU galios panaudojimo moksliniams ir kitiems negrafiniams skaičiavimams pagrindas yra algoritmų lygiagretinimas šimtuose vaizdo lustuose esančių vykdymo vienetų. Pavyzdžiui, daugelis molekulinio modeliavimo programų puikiai tinka skaičiavimams vaizdo lustuose, joms reikia didelės skaičiavimo galios, todėl jos yra patogios lygiagrečiam skaičiavimui. O kelių GPU naudojimas suteikia dar daugiau skaičiavimo galios tokioms problemoms išspręsti.

Skaičiavimai naudojant GPU rodo puikius algoritmų, kurie naudoja lygiagretų duomenų apdorojimą, rezultatus. Tai yra, kai ta pati matematinių operacijų seka taikoma dideliam duomenų kiekiui. Šiuo atveju geriausi rezultatai pasiekiami, jei aritmetinių komandų skaičiaus ir atminties prieigos skaičiaus santykis yra pakankamai didelis. Tai kelia mažiau reikalavimų srauto valdymui, o didelis matematikos tankis ir didelis duomenų kiekis pašalina didelių talpyklų poreikį, kaip ir CPU.

Dėl visų aukščiau aprašytų skirtumų teorinis vaizdo lustų našumas gerokai viršija procesoriaus našumą. „Nvidia“ pateikia tokią procesoriaus ir GPU našumo augimo per pastaruosius kelerius metus grafiką:

Natūralu, kad šie duomenys nėra be klastos. Juk ant procesoriaus praktiškai daug lengviau pasiekti teorinius skaičius, o skaičiai pateikiami dėl vienkartinio tikslumo, jei yra GPU, ir dėl dvigubo tikslumo, jei tai yra procesorius. Bet kokiu atveju kai kurioms lygiagrečioms užduotims pakanka vieno tikslumo, o greičio skirtumas tarp universalių ir grafinių procesorių yra labai didelis, todėl žaidimas yra vertas žvakės.

Pirmieji bandymai naudoti GPU skaičiavimus

Jie gana ilgą laiką bandė naudoti vaizdo lustus lygiagretiesiems matematiniams skaičiavimams. Pirmieji tokių programų bandymai buvo labai primityvūs ir apsiribojo tam tikrų aparatinės įrangos funkcijų, tokių kaip rastravimas ir Z-buferizavimas, naudojimu. Tačiau dabartiniame amžiuje, atsiradus šešėliams, matricos skaičiavimai pradėjo spartėti. 2003 m. SIGGRAPH GPU skaičiavimui buvo skirta atskira sekcija, kuri vadinosi GPGPU (General-Purpose computation on GPU).

Geriausiai žinomas yra BrookGPU Brook srauto programavimo kalbos kompiliatorius, skirtas atlikti negrafinius GPU skaičiavimus. Prieš pasirodant, kūrėjai, naudojantys vaizdo lustų galimybes skaičiavimams, pasirinko vieną iš dviejų įprastų API: Direct3D arba OpenGL. Tai rimtai apribojo GPU naudojimą, nes 3D grafikoje naudojami šešėliai ir tekstūros, apie kurias lygiagretaus programavimo specialistai neturi žinoti, jie naudoja gijas ir branduolius. Brook galėjo padėti palengvinti jų užduotį. Šie C kalbos transliacijos plėtiniai, sukurti Stanfordo universitete, paslėpė 3D API nuo programuotojų, o vaizdo lustą pristatė kaip lygiagretų koprocesorių. Kompiliatorius apdorojo .br failą su C++ kodu ir plėtiniais, sukurdamas kodą, susietą su DirectX, OpenGL arba x86 biblioteka.

Natūralu, kad Brook turėjo daug trūkumų, prie kurių mes apsistojome ir kuriuos išsamiau aptarsime vėliau. Tačiau net tik jo išvaizda sukėlė didelį tos pačios Nvidia ir ATI dėmesį į skaičiavimo ant GPU iniciatyvą, nes šių galimybių plėtra ateityje rimtai pakeitė rinką, atverdama visiškai naują jos sektorių - lygiagrečius kompiuterius. remiantis vaizdo lustais.

Vėliau kai kurie „Brook“ projekto tyrėjai prisijungė prie „Nvidia“ kūrimo komandos, siekdami pristatyti aparatinės ir programinės įrangos lygiagrečiojo skaičiavimo strategiją, atverdami naują rinkos dalį. O pagrindinis šios Nvidia iniciatyvos privalumas yra tas, kad kūrėjai iki smulkmenų žino visas savo GPU galimybes ir nereikia naudoti grafikos API, o su aparatine įranga galite dirbti tiesiogiai naudodami tvarkyklę. Šios komandos pastangų rezultatas buvo Nvidia CUDA (Compute Unified Device Architecture) – nauja programinės ir aparatinės įrangos architektūra, skirta lygiagrečiam Nvidia GPU skaičiavimui, kuri yra šio straipsnio tema.

Lygiagrečių skaičiavimų GPU taikymo sritys

Norėdami suprasti skaičiavimų perkėlimo į vaizdo lustus naudą, pateikiame vidutinius tyrėjų visame pasaulyje gautus skaičius. Vidutiniškai perkeliant skaičiavimus į GPU, daugelis užduočių pagreitina 5–30 kartų, palyginti su sparčiais bendrosios paskirties procesoriais. Didžiausi skaičiai (maždaug 100 kartų ar net daugiau!) pasiekiami su kodu, kuris nėra labai tinkamas skaičiavimams naudojant SSE blokus, tačiau yra gana patogus GPU.

Tai tik keli sintetinio kodo pagreitinimo GPU ir SSE vektorizuoto procesoriaus kodo pavyzdžiai (pagal Nvidia):

• Fluorescencinė mikroskopija: 12x;
• Molekulinė dinamika (nesurištos jėgos skaičiuok.): 8-16x;
• Elektrostatika (tiesioginis ir daugiapakopis Kulono sumavimas): 40-120x ir 7x.

Ir tai yra ženklas, kurį „Nvidia“ labai myli, rodydamas jį visuose pristatymuose, apie kuriuos plačiau kalbėsime antroje straipsnio dalyje, skirtoje konkretiems praktinio CUDA skaičiavimo taikymo pavyzdžiams:

Kaip matote, skaičiai labai patrauklūs, ypač įspūdingi 100-150 kartų padidėjimai. Kitame straipsnyje apie CUDA išsamiai apžvelgsime kai kuriuos iš šių skaičių. Dabar išvardinkime pagrindines programas, kuriose šiuo metu naudojamas GPU skaičiavimas: vaizdų ir signalų analizė ir apdorojimas, fizinis modeliavimas, skaičiavimo matematika, skaičiavimo biologija, finansiniai skaičiavimai, duomenų bazės, dujų ir skysčių dinamika, kriptografija, adaptyvioji spindulinė terapija, astronomija, garso apdorojimas, bioinformatika, biologinis modeliavimas, kompiuterinis matymas, duomenų gavyba, skaitmeninis kinas ir televizija, elektromagnetinis modeliavimas, geografinės informacijos sistemos, karinės programos, minų planavimas, molekulinė dinamika, magnetinio rezonanso tomografija (MRT), neuroniniai tinklai, okeanografiniai tyrimai, dalelių fizika , baltymų lankstymo modeliavimas, kvantinė chemija, spindulių sekimas, vizualizacija, radaras, rezervuarų modeliavimas, dirbtinis intelektas, palydovinių duomenų analizė, seisminis tyrinėjimas, chirurgija, ultragarsas, vaizdo konferencijos.

Išsamią informaciją apie daugelį programų galite rasti Nvidia svetainės skyriuje. Kaip matote, sąrašas yra gana didelis, bet tai dar ne viskas! Galima tęsti ir tikrai galime manyti, kad ateityje bus rasta ir kitų lygiagrečių skaičiavimų vaizdo lustuose taikymo sričių, kurių dar nežinome.

Nvidia CUDA funkcijos

CUDA technologija yra „Nvidia“ aparatinės ir programinės įrangos skaičiavimo architektūra, pagrįsta C kalbos plėtiniu, kuri leidžia organizuoti prieigą prie grafikos greitintuvo instrukcijų rinkinio ir valdyti jo atmintį organizuojant lygiagretųjį skaičiavimą. CUDA padeda įgyvendinti algoritmus, vykdomus aštuntos kartos ir senesniuose GeForce vaizdo greitintuvuose (Geforce 8, Geforce 9, Geforce 200 serijos), taip pat Quadro ir Tesla.

Nors GPU programavimo sudėtingumas naudojant CUDA yra gana didelis, jis yra mažesnis nei naudojant ankstesnius GPGPU sprendimus. Tokios programos reikalauja padalinti programą tarp kelių procesorių, panašiai kaip MPI programavimas, tačiau nesidalinant duomenimis, saugomais bendroje vaizdo atmintyje. Ir kadangi CUDA programavimas kiekvienam kelių procesorių yra panašus į OpenMP programavimą, tam reikia gerai suprasti atminties organizavimą. Tačiau, žinoma, kūrimo ir perkėlimo į CUDA sudėtingumas labai priklauso nuo programos.

Kūrėjo rinkinyje yra daug kodų pavyzdžių ir jis yra gerai dokumentuotas. Mokymosi procesas užtruks apie dvi ar keturias savaites tiems, kurie jau yra susipažinę su OpenMP ir MPI. API yra pagrįsta išplėstine C kalba, o norint išversti kodą iš šios kalbos, CUDA SDK apima nvcc komandinės eilutės kompiliatorių, pagrįstą atvirojo kodo Open64 kompiliatoriumi.

Išvardinkime pagrindines CUDA savybes:

• vieningas programinės ir techninės įrangos sprendimas lygiagrečiam skaičiavimui Nvidia vaizdo lustuose;
• platus palaikomų sprendimų asortimentas – nuo ​​mobiliųjų iki kelių lustų
• standartinė C programavimo kalba;
• standartinės skaitmeninės analizės bibliotekos FFT (Fast Furjė transformacija) ir BLAS (tiesinė algebra);
• optimizuotas keitimasis duomenimis tarp procesoriaus ir GPU;
• sąveika su grafinėmis API OpenGL ir DirectX;
• 32 ir 64 bitų operacinių sistemų palaikymas: Windows XP, Windows Vista, Linux ir MacOS X;
• Galimybė vystytis žemame lygyje.

Kalbant apie operacinės sistemos palaikymą, reikia pridurti, kad visi pagrindiniai Linux platinimai yra oficialiai palaikomi (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), tačiau, sprendžiant iš entuziastų gautų duomenų, CUDA puikiai veikia kitose versijose: Fedora Core, Ubuntu, Gentoo ir kt.

CUDA kūrimo aplinka (CUDA Toolkit) apima:

• nvcc kompiliatorius;
• FFT ir BLAS bibliotekos;
• profiliuotojas;
• gdb derinimo priemonė GPU;
• CUDA vykdymo tvarkyklė įtraukta į standartines Nvidia tvarkykles
• programavimo vadovas;
• CUDA Developer SDK (šaltinio kodas, komunalinės paslaugos ir dokumentacija).

Šaltinio kodo pavyzdžiuose: lygiagretus bitoninis rūšiavimas, matricos perkėlimas, lygiagretus didelių masyvų priešdėlių sumavimas, vaizdo konvoliucija, diskrečiųjų bangelių transformacija, sąveikos su OpenGL ir Direct3D pavyzdys, CUBLAS ir Cufft bibliotekų naudojimas, pasirinkimo kainos apskaičiavimas (juodas formulė Scholes, binominis modelis, Monte Karlo metodas), lygiagretusis atsitiktinių skaičių generatorius Mersenne Twister, didelio masyvo histogramos skaičiavimas, triukšmo mažinimas, Sobelio filtras (ribų radimas).

CUDA privalumai ir apribojimai

Programuotojo požiūriu grafikos vamzdynas yra apdorojimo etapų rinkinys. Geometrijos blokas generuoja trikampius, o rastravimo blokas – monitoriuje rodomus pikselius. Tradicinis GPGPU programavimo modelis atrodo taip:

Norint perkelti skaičiavimus į šio modelio GPU, reikalingas specialus požiūris. Netgi pridedant du vektorius pagal elementus reikės nupiešti figūrą ekrane arba ne ekrano buferyje. Figūra rastrizuota, kiekvieno pikselio spalva apskaičiuojama naudojant tam tikrą programą (pixel shader). Programa nuskaito įvesties duomenis iš kiekvieno pikselio tekstūrų, juos prideda ir įrašo į išvesties buferį. Ir visos šios daugybės operacijų reikalingos tam, kas parašyta vienu operatoriumi įprasta programavimo kalba!

Todėl GPGPU naudojimas bendrosios paskirties kompiuteriams yra ribotas, nes jį per sunku mokyti kūrėjus. Taip, ir yra pakankamai kitų apribojimų, nes pikselių atspalvis yra tik galutinės pikselio spalvos priklausomybės nuo jo koordinatės formulė, o pikselių šešėlių kalba yra šių formulių rašymo kalba su C tipo sintaksė. . Ankstyvieji GPGPU metodai yra puikus triukas, leidžiantis išnaudoti GPU galią, bet be jokio patogumo. Ten esantys duomenys atvaizduojami vaizdais (tekstūros), o algoritmas – rastrizacijos procesu. Ypač atkreiptinas dėmesys į labai specifinį atminties ir vykdymo modelį.

„Nvidia“ programinės ir techninės įrangos architektūra, skirta GPU skaičiavimui, skiriasi nuo ankstesnių GPGPU modelių tuo, kad leidžia GPU programas rašyti tikra C kalba su standartine sintaksė, rodyklėmis ir būtinybe turėti minimalų plėtinį, kad būtų galima pasiekti vaizdo lustų skaičiavimo išteklius. CUDA nepriklauso nuo grafinių API ir turi kai kurias funkcijas, specialiai sukurtas bendrosios paskirties kompiuteriams.

CUDA pranašumai, palyginti su tradiciniu GPGPU skaičiavimo metodu:

• CUDA taikomųjų programų programavimo sąsaja yra pagrįsta standartine C programavimo kalba su plėtiniais, kas supaprastina CUDA architektūros mokymosi ir diegimo procesą;
• CUDA suteikia prieigą prie 16 KB gijų bendrinamos atminties vienam daugiaprocesoriui, kurią galima naudoti talpyklai sutvarkyti didesniu pralaidumu nei tekstūros gavimas;
• efektyvesnis duomenų perdavimas tarp sistemos ir vaizdo atminties
• nereikia grafinių API su pertekliumi ir pridėtinėmis išlaidomis;
• linijinis atminties adresavimas, surinkimas ir išsklaidymas, galimybė rašyti savavališkais adresais;
• aparatinės įrangos palaikymas sveikųjų ir bitų operacijoms.

Pagrindiniai CUDA apribojimai:

• vykdomųjų funkcijų rekursijos palaikymo trūkumas;
• minimalus bloko plotis 32 sriegiai;
• uždara CUDA architektūra, priklausanti Nvidia.

Ankstesnių GPGPU metodų programavimo trūkumai yra tai, kad šie metodai nenaudoja viršūnių šešėlių vykdymo vienetų ankstesnėse nesuvienodintose architektūrose, duomenys saugomi tekstūrose ir išvedami į ne ekrano buferį, o kelių žingsnių algoritmuose naudojami pikselių šešėlių vienetai. GPGPU apribojimai gali apimti: nepakankamą techninės įrangos galimybių naudojimą, atminties pralaidumo apribojimus, sklaidos veikimo trūkumą (tik rinkti), privalomą grafikos API naudojimą.

Pagrindiniai CUDA pranašumai, palyginti su ankstesniais GPGPU metodais, kyla iš to, kad architektūra sukurta efektyviai panaudoti ne grafinį skaičiavimą GPU ir naudoja C programavimo kalbą, nereikalaujant, kad algoritmai būtų perkeliami į grafinio konvejerio koncepcijai palankią formą. . CUDA siūlo naują kelią į GPU skaičiavimą, kuris nenaudoja grafikos API, siūlydamas atsitiktinę prieigą prie atminties (išsklaidyti arba surinkti). Ši architektūra neturi GPGPU trūkumų ir naudoja visus vykdymo blokus, taip pat išplečia galimybes dėl sveikųjų skaičių matematikos ir bitų poslinkio operacijų.

Be to, CUDA atveria kai kurias aparatinės įrangos galimybes, kurių nėra grafikos API, pvz., bendrinamą atmintį. Tai maža atmintis (16 kilobaitų vienam procesoriui), prie kurios gijų blokai turi prieigą. Tai leidžia talpykloje išsaugoti dažniausiai pasiekiamus duomenis ir gali užtikrinti didesnį greitį nei naudojant tekstūros gavimą šiai užduočiai atlikti. O tai savo ruožtu sumažina lygiagrečių algoritmų pralaidumo jautrumą daugelyje programų. Pavyzdžiui, tai naudinga tiesinei algebrai, greitajai Furjė transformacijai ir vaizdo apdorojimo filtrams.

Prieiga prie atminties taip pat patogesnė naudojant CUDA. Grafikos API kodas išveda duomenis kaip 32 vieno tikslumo slankiojo kablelio reikšmes (RGBA reikšmės vienu metu į aštuonis atvaizdavimo objektus) į iš anksto nustatytas sritis, o CUDA palaiko išsklaidytą rašymą – neribotą įrašų skaičių bet kuriuo adresu. Tokie pranašumai leidžia vykdyti kai kuriuos GPU algoritmus, kurių negalima efektyviai įgyvendinti naudojant GPGPU metodus, pagrįstus grafinėmis API.

Taip pat grafinės API būtinai saugo duomenis tekstūrose, o tam reikia iš anksto supakuoti didelius masyvus į tekstūras, o tai apsunkina algoritmą ir verčia naudoti specialų adresavimą. O CUDA leidžia skaityti duomenis bet kuriuo adresu. Kitas CUDA privalumas yra optimizuotas keitimasis duomenimis tarp procesoriaus ir GPU. O kūrėjams, norintiems žemo lygio prieigos (pavyzdžiui, rašant kitą programavimo kalbą), CUDA siūlo žemo lygio asamblėjos kalbos programavimo galimybes.

CUDA vystymosi istorija

Apie CUDA kūrimą kartu su G80 lustu buvo pranešta 2006 m. lapkritį, o vieša CUDA SDK beta versija buvo išleista 2007 m. vasario mėn. 1.0 versija buvo išleista 2007 m. birželio mėn., kad būtų pristatyti Tesla sprendimai, pagrįsti G80 lustu ir skirti didelio našumo kompiuterių rinkai. Tada metų pabaigoje buvo išleista CUDA 1.1 beta versija, kuri, nepaisant nežymaus versijos skaičiaus padidėjimo, pristatė gana daug naujų dalykų.

CUDA 1.1 naujovė yra CUDA funkcijos įtraukimas į įprastas Nvidia vaizdo tvarkykles. Tai reiškė, kad bet kurios CUDA programos reikalavimuose užteko nurodyti Geforce 8 ir naujesnės serijos vaizdo plokštę, taip pat minimalią tvarkyklės versiją 169.xx. Tai labai svarbu kūrėjams, jei šios sąlygos bus įvykdytos, CUDA programos veiks bet kuriam vartotojui. Taip pat buvo pridėtas asinchroninis vykdymas kartu su duomenų kopijavimu (tik G84, G86, G92 ir aukštesniems lustams), asinchroninis duomenų perkėlimas į vaizdo atmintį, prieigos prie atominės atminties operacijos, 64 bitų Windows versijų palaikymas ir galimybė valdyti kelis CUDA lustas SLI režimu.

Šiuo metu dabartinė versija skirta GT200 CUDA 2.0 pagrindu, išleista kartu su Geforce GTX 200 linija. Beta versija buvo išleista dar 2008 m. pavasarį. Pridėta antroji versija: dvigubo tikslumo skaičiavimų palaikymas (aparatinės įrangos palaikymas tik GT200), galiausiai palaiko Windows Vista (32 ir 64 bitų versijos) ir Mac OS X, pridedami derinimo ir profiliavimo įrankiai, palaikomos 3D tekstūros, optimizuotas duomenų perdavimas. .

Kalbant apie dvigubo tikslumo skaičiavimus, jų greitis dabartinėje aparatūros kartoje yra kelis kartus mažesnis nei vieno tikslumo. Priežastys aptariamos mūsų. Šios paramos įgyvendinimas GT200 yra tas, kad FP32 blokai nenaudojami norint gauti rezultatus keturis kartus mažesniu greičiu, kad palaikytų FP64 skaičiavimus, „Nvidia“ nusprendė sukurti specialius skaičiavimo įrenginius. O GT200 jų yra dešimt kartų mažiau nei FP32 blokų (po vieną dvigubo tikslumo bloką kiekvienam daugiaprocesoriui).

Realybėje našumas gali būti dar mažesnis, nes architektūra optimizuota 32 bitų skaitymui iš atminties ir registrų, be to, grafinėse programose nereikia dvigubo tikslumo, o GT200 jis yra labiau tikėtinas. Ir šiuolaikiniai keturių branduolių procesoriai rodo ne ką mažiau tikrą našumą. Tačiau būdama net 10 kartų lėtesnė už vienkartinį tikslumą, tokia parama yra naudinga mišriam tikslumo dizainui. Vienas iš įprastų metodų yra iš pradžių gauti apytikslius rezultatus vienu tikslumu, o tada juos patobulinti dvigubu tikslumu. Dabar tai galima padaryti tiesiai vaizdo plokštėje, nesiunčiant tarpinių duomenų į centrinį procesorių.

Kita naudinga CUDA 2.0 funkcija, kaip bebūtų keista, neturi nieko bendra su GPU. Tiesiog dabar galite sukompiliuoti CUDA kodą į labai efektyvų kelių gijų SSE kodą, kad būtų galima greitai vykdyti CPU. Tai yra, dabar ši funkcija tinka ne tik derinimui, bet ir realiam naudojimui sistemose be Nvidia vaizdo plokštės. Juk CUDA naudojimą įprastame kode apsunkina tai, kad Nvidia vaizdo plokštės, nors ir populiariausios tarp dedikuotų vaizdo sprendimų, yra prieinamos ne visose sistemose. O iki 2.0 versijos tokiais atvejais būtų reikėję sukurti du skirtingus kodus: CUDA ir atskirai CPU. Ir dabar galite paleisti bet kurią CUDA programą didelio efektyvumo CPU, nors ir mažesniu greičiu nei vaizdo lustuose.

Nvidia CUDA palaikantys sprendimai

Visos CUDA palaikančios vaizdo plokštės gali padėti pagreitinti sudėtingiausias užduotis – nuo ​​garso ir vaizdo apdorojimo iki medicinos ir mokslinių tyrimų. Vienintelis tikras apribojimas yra tas, kad daugeliui CUDA programų reikia mažiausiai 256 megabaitų vaizdo atminties, ir tai yra viena iš svarbiausių techninių CUDA programų charakteristikų.

Naujausią CUDA palaikančių produktų sąrašą galite rasti adresu. Rašymo metu CUDA skaičiavimai palaikė visus „GeForce 200“, „GeForce 9“ ir „GeForce 8“ serijų produktus, įskaitant mobiliuosius produktus, pradedant nuo „GeForce 8400M“, taip pat „GeForce 8100“, „8200“ ir „8300“ lustų rinkinius bei „Modern Quadro“ ir 8300 gaminius palaiko CUDA visas Teslas: S1070, C1060, C870, D870 ir S870.

Ypatingai pažymime, kad kartu su naujomis Geforce GTX 260 ir 280 vaizdo plokštėmis buvo paskelbti ir atitinkami didelio našumo skaičiavimo sprendimai: Tesla C1060 ir S1070 (parodyta aukščiau esančioje nuotraukoje), kuriuos bus galima įsigyti jau šį rudenį. Jie naudoja tą patį GPU - GT200, C1060 yra vienas, S1070 - keturi. Tačiau, skirtingai nei žaidimų sprendimai, jie naudoja keturis gigabaitus atminties vienai lustai. Vienintelis trūkumas yra tas, kad atminties dažnis ir pralaidumas yra mažesni nei žaidimų kortelių, užtikrinant 102 GB/s vienam lustui.

Nvidia CUDA sudėtis

CUDA apima dvi API: aukšto lygio (CUDA Runtime API) ir žemo lygio (CUDA Driver API), nors neįmanoma naudoti abiejų vienu metu, turite naudoti vieną ar kitą. Aukšto lygio veikia „viršuje“ žemo lygio, visi vykdymo iškvietimai paverčiami paprastomis instrukcijomis, kurias apdoroja žemo lygio tvarkyklės API. Tačiau net „aukšto lygio“ API reikalauja žinių apie „Nvidia“ vaizdo lustų dizainą ir veikimą, ten nėra per didelio abstrakcijos lygio.

Yra dar vienas lygis, dar aukštesnis – dvi bibliotekos:

CUBLAS BLAS (Basic Linear Algebra Subprograms) CUDA versija, skirta skaičiuoti tiesinės algebros problemas ir naudojant tiesioginę prieigą prie GPU išteklių;

RANKELIAI Greitosios Furjė transformacijos bibliotekos CUDA versija, skirta greitajai Furjė transformacijai apskaičiuoti, plačiai naudojama signalų apdorojimui. Palaikomi šie transformacijų tipai: kompleksinis kompleksas (C2C), realus kompleksas (R2C) ir kompleksinis-realus (C2R).

Pažvelkime į šias bibliotekas atidžiau. CUBLAS yra standartiniai tiesinės algebros algoritmai, išversti į CUDA kalbą, šiuo metu palaikomas tik tam tikras pagrindinių CUBLAS funkcijų rinkinys. Biblioteka labai paprasta naudotis: vaizdo plokštės atmintyje reikia sukurti matricą ir vektorinius objektus, užpildyti juos duomenimis, iškviesti reikiamas CUBLAS funkcijas, o rezultatus iš vaizdo atminties įkelti atgal į sistemos atmintį. CUBLAS yra specialios funkcijos, skirtos GPU atmintyje esantiems objektams kurti ir naikinti, taip pat duomenims nuskaityti ir įrašyti į šią atmintį. Palaikomos BLAS funkcijos: 1, 2 ir 3 lygiai tikriems skaičiams, CGEMM 1 lygis kompleksiniams skaičiams. 1 lygis yra vektoriaus-vektoriaus operacijos, 2 lygis yra vektoriaus-matricos operacijos, 3 lygis yra matricos-matricos operacijos.

Cufft CUDA greito Furjė transformacijos funkcijos versija, plačiai naudojama ir labai svarbi signalų analizėje, filtravime ir kt. CUFFT suteikia paprastą sąsają, leidžiančią efektyviai apskaičiuoti FFT naudojant Nvidia GPU, nereikia kurti savo GPU FFT. FFT CUDA variantas palaiko sudėtingų ir realių duomenų 1D, 2D ir 3D transformacijas, kelių 1D transformacijų paketinį vykdymą lygiagrečiai, 2D ir 3D transformacijų dydžiai gali būti , 1D atveju iki 8 milijonų elementų. palaikoma.

Programų kūrimo CUDA pagrindu pagrindai

Norėdami suprasti tolesnį tekstą, turėtumėte suprasti pagrindines Nvidia vaizdo lustų architektūrines ypatybes. GPU susideda iš kelių tekstūros vienetų grupių (Texture Processing Cluster). Kiekvieną klasterį sudaro didelis tekstūros gavimo blokas ir du arba trys srautiniai daugiaprocesoriai, kurių kiekvienas susideda iš aštuonių skaičiavimo įrenginių ir dviejų superfunkcinių blokų. Visos instrukcijos vykdomos naudojant SIMD principą, kai viena instrukcija taikoma visoms metmenų gijomis (tekstilės pramonės terminas, CUDA tai yra 32 gijų grupė, kuri yra minimalus kelių procesorių apdorojamas duomenų kiekis). Šis vykdymo būdas buvo vadinamas SIMT (viena instrukcija, kelios gijos viena instrukcija ir daug gijų).

Kiekvienas iš kelių procesorių turi tam tikrus išteklius. Taigi, yra speciali bendra 16 kilobaitų atmintis vienam kelių procesoriui. Bet tai nėra talpykla, nes programuotojas gali ją naudoti bet kokiam poreikiui, kaip vietinė parduotuvė ląstelių procesorių SPU. Ši bendra atmintis leidžia keistis informacija tarp to paties bloko gijų. Svarbu, kad visos vieno bloko gijos visada būtų vykdomos to paties daugiaprocesoriaus. Tačiau gijos iš skirtingų blokų negali keistis duomenimis, todėl reikia atsiminti šį apribojimą. Bendrinama atmintis dažnai yra naudinga, išskyrus atvejus, kai kelios gijos pasiekia tą patį atminties banką. Daugiaprocesoriai taip pat gali pasiekti vaizdo atmintį, tačiau su didesniu delsimu ir prastesniu pralaidumu. Norėdami pagreitinti prieigą ir sumažinti prieigos prie vaizdo atminties dažnį, daugiaprocesoriai turi 8 kilobaitus talpyklos konstantoms ir tekstūros duomenims.

Daugiaprocesorius naudoja 8192-16384 (atitinkamai G8x/G9x ir GT2xx) registrus, bendrus visoms visų jame vykdomų blokų gijomis. Maksimalus blokų skaičius viename daugiaprocesoriuje G8x / G9x yra aštuoni, o metmenų skaičius yra 24 (768 gijos viename daugiaprocesoriuje). Iš viso aukščiausios klasės Geforce 8 ir 9 serijos vaizdo plokštės vienu metu gali apdoroti iki 12288 gijų. „GeForce GTX 280“, pagrįstas GT200, siūlo iki 1024 gijų viename daugiaprocesoriuje, jame yra 10 trijų kelių procesorių grupių, apdorojančių iki 30720 gijų. Žinodami šiuos apribojimus galite optimizuoti turimų išteklių algoritmus.

Pirmas žingsnis perkeliant esamą programą į CUDA yra jos profiliavimas ir kodo sričių, kurios yra kliūtys, lėtinančios darbą, nustatymas. Jei tarp tokių skyrių yra tinkamų greitam lygiagrečiam vykdymui, šios funkcijos perkeliamos į CUDA C plėtinius, kad būtų galima vykdyti GPU. Programa kompiliuojama naudojant „Nvidia“ tiekiamą kompiliatorių, kuris generuoja CPU ir GPU kodą. Kai programa vykdoma, centrinis procesorius vykdo savo kodo dalis, o GPU vykdo CUDA kodą su labiausiai lygiagrečiais skaičiavimais. Ši dalis, skirta GPU, vadinama branduoliu. Branduolys apibrėžia operacijas, kurios bus atliekamos su duomenimis.

Vaizdo lustas gauna branduolį ir sukuria kiekvieno duomenų elemento kopijas. Šios kopijos vadinamos gijomis. Sraute yra skaitiklis, registrai ir būsena. Dideliam duomenų kiekiui, pvz., vaizdo apdorojimui, paleidžiama milijonai gijų. Siūlai vykdomi grupėmis po 32, vadinami metimais, kuriems priskiriamas vykdymas konkrečiuose gijų daugiaprocesoriuose. Kiekvienas daugiaprocesorius susideda iš aštuonių branduolių srauto procesorių, kurie vykdo vieną MAD komandą per laikrodžio ciklą. Norint atlikti vieną 32 gijų deformaciją, reikalingi keturi kelių procesorių laikrodžio ciklai (kalbame apie šešėlio srities dažnį, kuris yra 1,5 GHz ir didesnis).

Daugiaprocesorius nėra tradicinis kelių branduolių procesorius, kuris vienu metu palaiko iki 32 deformacijų. laikrodžio ciklai. Jei nubrėžtume analogiją su centriniu procesoriumi, tai panašu į 32 programų vykdymą vienu metu ir perjungimą tarp jų kiekvieną laikrodžio ciklą neprarandant konteksto perjungimo. Tiesą sakant, procesoriaus branduoliai palaiko vienos programos vykdymą vienu metu ir perjungia į kitas su šimtų ciklų vėlavimu.

CUDA programavimo modelis

Pakartokime, kad CUDA naudoja lygiagretųjį skaičiavimo modelį, kai kiekvienas SIMD procesorius lygiagrečiai vykdo tą pačią komandą skirtingiems duomenų elementams. GPU yra skaičiavimo įrenginys, centrinio procesoriaus (host) koprocesorius (įrenginys), kuris turi savo atmintį ir lygiagrečiai apdoroja daugybę gijų. Branduolys yra GPU funkcija, vykdoma gijomis (analogija iš 3D grafikos – atspalvis).

Aukščiau minėjome, kad vaizdo lustas skiriasi nuo procesoriaus tuo, kad vienu metu gali apdoroti dešimtis tūkstančių gijų, o tai būdinga gerai lygiagrečiai grafikai. Kiekvienas srautas yra skaliarinis ir nereikalauja duomenų pakavimo į 4 komponentų vektorius, o tai yra patogiau daugeliui užduočių. Loginių gijų ir gijų blokų skaičius viršija fizinių vykdymo įrenginių skaičių, o tai užtikrina gerą mastelio keitimą visam įmonės sprendimų spektrui.

CUDA programavimo modelis apima gijų grupavimą. Gijos yra suskirstytos į gijų blokus – vienmačius arba dvimačius gijų tinklelius, kurie bendrauja tarpusavyje naudodami bendrą atmintį ir sinchronizavimo taškus. Programa (branduolis) veikia per gijų blokų tinklelį, žr. toliau pateiktą paveikslėlį. Vienu metu vykdomas vienas tinklelis. Kiekvienas blokas gali būti vienos, dviejų arba trijų matmenų formos ir gali būti sudarytas iš 512 dabartinės aparatinės įrangos gijų.

Siūlų blokai vykdomi mažose grupėse, vadinamose metmenyse, kurių dydis yra 32 gijos. Tai yra mažiausias duomenų kiekis, kurį galima apdoroti keliuose procesoriuose. Ir kadangi tai ne visada patogu, CUDA leidžia dirbti su blokais, kuriuose yra nuo 64 iki 512 gijų.

Blokų grupavimas į tinklelius leidžia išvengti apribojimų ir vienu skambučiu pritaikyti branduolį daugiau gijų. Tai taip pat padeda didinti mastelį. Jei GPU neturi pakankamai išteklių, jis blokus vykdys nuosekliai. Priešingu atveju blokai gali būti vykdomi lygiagrečiai, o tai svarbu norint optimaliai paskirstyti darbą skirtingų lygių vaizdo lustuose, pradedant mobiliaisiais ir integruotais.

CUDA atminties modelis

CUDA atminties modelis išsiskiria baitų adresavimo galimybe, tiek rinkimo, tiek išsklaidymo palaikymu. Kiekvienam srauto procesoriui yra gana daug registrų, iki 1024 vienetų. Prieiga prie jų yra labai greita ir juose galite saugoti 32 bitų sveikuosius skaičius arba slankiojo kablelio skaičius.

Kiekviena gija turi prieigą prie šių tipų atminties:

Pasaulinė atmintis didžiausias visų vaizdo lusto kelių procesorių atminties kiekis, kurio dydis dabartiniuose sprendimuose svyruoja nuo 256 megabaitų iki 1,5 gigabaitų (ir iki 4 GB „Tesla“). Jis turi didelį pralaidumą, daugiau nei 100 gigabaitų per sekundę geriausiems Nvidia sprendimams, tačiau labai didelis kelių šimtų ciklų delsos laikas. Neįkeliama talpykloje, palaiko bendrąsias įkėlimo ir saugojimo instrukcijas bei įprastas atminties nuorodas.

Vietinė atmintis tai nedidelis atminties kiekis, prie kurio gali prieiti tik vienas srauto procesorius. Jis gana lėtas – toks pat kaip pasaulinis.

Bendra atmintis tai 16 kilobaitų (dabartinės architektūros vaizdo lustuose) atminties blokas su bendra prieiga visiems srauto procesoriams daugiaprocesoriuje. Ši atmintis yra labai greita, tokia pati kaip registrai. Tai leidžia gijomis bendrauti, yra tiesiogiai valdoma kūrėjo ir turi mažą delsą. Bendrosios atminties privalumai: naudoti kaip programuotojo valdomą pirmojo lygio talpyklą, sumažinta delsa, kai pasiekiami duomenų vykdymo vienetai (ALU), sumažėja prieigų prie pasaulinės atminties skaičius.

Nuolatinė atmintis- 64 kilobaitų atminties plotas (tas pats dabartiniams GPU), skaitomas visų kelių procesorių. Viename kelių procesorių talpykloje yra 8 kilobaitai. Gana lėtas - kelių šimtų ciklų vėlavimas, kai talpykloje nėra reikiamų duomenų.

Tekstūros atmintis atminties blokas, kurį gali nuskaityti visi daugiaprocesoriai. Duomenų atranka atliekama naudojant vaizdo lusto tekstūros blokus, todėl linijinė duomenų interpoliacija galima be papildomų išlaidų. Viename kelių procesorių talpykloje saugomi 8 kilobaitai. Lėtas kaip pasaulinis šimtai delsos ciklų, kai talpykloje nėra duomenų.

Natūralu, kad globali, lokali, tekstūrinė ir pastovioji atmintis fiziškai yra ta pati atmintis, žinoma kaip vietinė vaizdo plokštės vaizdo atmintis. Jų skirtumai slypi skirtinguose talpyklos algoritmuose ir prieigos modeliuose. CPU gali atnaujinti ir prašyti tik išorinės atminties: visuotinės, pastovios ir tekstūros.

Iš to, kas buvo parašyta aukščiau, aišku, kad CUDA prisiima ypatingą požiūrį į kūrimą, ne visai tą patį, kaip priimtas CPU programose. Turite atsiminti apie skirtingus atminties tipus, kad vietinė ir globali atmintis nėra talpykloje, o prieigos prie jos delsa yra daug didesnė nei registrinės atminties, nes ji fiziškai yra atskiruose lustuose.

Tipiškas, bet neprivalomas problemų sprendimo modelis:

• užduotis suskirstyta į dalines užduotis;
• įvesties duomenys yra suskirstyti į blokus, kurie telpa į bendrą atmintį;
• kiekvienas blokas apdorojamas gijų bloku;
• subblokas įkeliamas į bendrą atmintį iš pasaulinės atminties;
• atliekami atitinkami bendrosios atminties duomenų skaičiavimai;
• rezultatai nukopijuojami iš bendros atminties atgal į pasaulinę atmintį.

Programavimo aplinka

CUDA apima vykdymo bibliotekas:

• bendroji dalis, teikianti integruotus vektorinius tipus ir RTL skambučių pogrupius, palaikomus CPU ir GPU;
• CPU komponentas, skirtas valdyti vieną ar daugiau GPU;
• GPU komponentas, teikiantis GPU specifines funkcijas.

Pagrindinis CUDA taikymo procesas veikia universaliame procesoriuje (host), jis vykdo kelias branduolio procesų kopijas vaizdo plokštėje. CPU kodas atlieka šiuos veiksmus: inicijuoja GPU, paskiria atmintį vaizdo plokštėje ir sistemoje, nukopijuoja konstantas į vaizdo plokštės atmintį, paleidžia keletą branduolio procesų vaizdo plokštėje kopijų, nukopijuoja rezultatą iš vaizdo atminties, atlaisvina atmintis ir išėjimai.

Kaip suprasti supratimo pavyzdį, čia yra vektoriaus pridėjimo procesoriaus kodas, pateiktas CUDA:

Vaizdo lusto vykdomos funkcijos turi šiuos apribojimus: nėra rekursijos, nėra statinių kintamųjų funkcijų viduje ir nėra kintamo argumentų skaičiaus. Palaikomi du atminties valdymo tipai: linijinė atmintis su prieiga per 32 bitų rodykles ir CUDA matricos su prieiga tik per tekstūros gavimo funkcijas.

CUDA programos gali sąveikauti su grafinėmis API: pateikti programoje sugeneruotus duomenis, skaityti atvaizdavimo rezultatus ir apdoroti juos naudojant CUDA įrankius (pvz., diegiant papildomo apdorojimo filtrus). Norint tai pasiekti, grafikos API ištekliai gali būti susieti (gaunantys išteklių adresą) į CUDA pasaulinę atminties erdvę. Palaikomi šie grafikos API išteklių tipai: Buferiniai objektai (PBO/VBO) OpenGL, viršūnių buferiai ir tekstūros (2D, 3D ir kubinės schemos) Direct3D9.

CUDA programos sudarymo etapai:

CUDA C šaltinio kodo failai kompiliuojami naudojant NVCC programą, kuri yra kitų įrankių įvynioklis ir juos iškviečia: cudacc, g++, cl ir tt NVCC generuoja: CPU kodą, kuris kompiliuojamas kartu su likusia programa, parašyta grynai. C ir vaizdo lusto PTX objekto kodas. Vykdomiems failams su CUDA kodu būtinai reikalinga CUDA vykdymo biblioteka (cudart) ir CUDA pagrindinė biblioteka (cuda).

CUDA programų optimizavimas

Natūralu, kad apžvalginiame straipsnyje neįmanoma apsvarstyti rimtų CUDA programavimo optimizavimo problemų. Todėl trumpai pakalbėsime apie pagrindinius dalykus. Norėdami efektyviai išnaudoti CUDA galimybes, turite pamiršti įprastus CPU programų rašymo būdus ir naudoti tuos algoritmus, kurie yra gerai lygiagretūs tūkstančiams gijų. Taip pat svarbu rasti optimalią duomenų saugojimo vietą (registrus, bendrinamą atmintį ir kt.), sumažinti duomenų perdavimą tarp procesoriaus ir GPU bei naudoti buferį.

Apskritai, optimizuodami CUDA programą, turėtumėte stengtis pasiekti optimalų balansą tarp dydžio ir blokų skaičiaus. Daugiau gijų viename bloke sumažins atminties delsos poveikį, bet taip pat sumažins galimų registrų skaičių. Be to, 512 gijų blokas yra neefektyvus, norint pasiekti optimalų delsą ir registrų skaičių, pati Nvidia rekomenduoja naudoti 128 arba 256 gijų blokus.

Tarp pagrindinių CUDA programų optimizavimo taškų: aktyviausias bendrosios atminties naudojimas, nes jis yra daug greitesnis nei pasaulinė vaizdo plokštės vaizdo atmintis; Skaitymas ir rašymas iš pasaulinės atminties turėtų būti sujungti, kai tik įmanoma. Norėdami tai padaryti, turite naudoti specialius duomenų tipus, kad per vieną operaciją vienu metu būtų galima nuskaityti ir įrašyti 32/64/128 bitų duomenų. Jei skaitymus sunku derinti, galite pabandyti naudoti tekstūros gavimą.

išvadas

Nvidia pateikta programinės ir techninės įrangos architektūra, skirta CUDA vaizdo lustų skaičiavimams, puikiai tinka sprendžiant daugybę labai lygiagrečių užduočių. CUDA veikia su įvairiais Nvidia GPU ir pagerina GPU programavimo modelį, jį labai supaprastindama ir pridedant daugybę funkcijų, tokių kaip bendra atmintis, gijų sinchronizavimas, dvigubo tikslumo skaičiavimai ir sveikųjų skaičių operacijos.

CUDA yra technologija, prieinama kiekvienam programinės įrangos kūrėjui, ją gali naudoti bet kuris programuotojas, žinantis C kalbą. Jūs tiesiog turite priprasti prie skirtingos programavimo paradigmos, būdingos lygiagrečiam skaičiavimui. Bet jei algoritmas iš esmės yra gerai lygiagretus, studijos ir laikas, praleistas programuojant CUDA, grįš daug kartų.

Tikėtina, kad dėl plačiai paplitusių vaizdo plokščių pasaulyje, lygiagretaus skaičiavimo kūrimas GPU didelės įtakos didelio našumo kompiuterijos pramonei. Šios galimybės jau sukėlė didelį susidomėjimą mokslo sluoksniuose, ir ne tik juose. Galų gale, potencialios galimybės pagreitinti algoritmus, kuriuos galima lengvai lygiagretinti (turimoje aparatinėje įrangoje, kas ne mažiau svarbu) dešimtis kartų, nėra tokios dažnos.

Bendrosios paskirties procesoriai vystosi gana lėtai ir neturi tokių našumo šuolių. Iš esmės, nors tai skamba kaip dideli pinigai, kiekvienas, kuriam reikia greito skaičiavimo, dabar gali turėti nebrangų asmeninį superkompiuterį ant stalo, kartais net neinvestuojant papildomai, nes „Nvidia“ vaizdo plokštės yra taip plačiai prieinamos. Jau nekalbant apie padidintą GFLOPS/$ ir GFLOPS/W efektyvumą, kurį taip mėgsta GPU gamintojai.

Daugelio kompiuterijos ateitis aiškiai slypi lygiagrečiuose algoritmuose, ir beveik visi nauji sprendimai ir iniciatyvos yra nukreiptos šia kryptimi. Tačiau kol kas naujų paradigmų kūrimas yra pradiniame etape, jūs turite rankiniu būdu kurti gijas ir planuoti prieigą prie atminties, o tai apsunkina užduotis, palyginti su įprastu programavimu. Tačiau CUDA technologija žengė žingsnį teisinga linkme ir akivaizdžiai atrodo sėkmingas sprendimas, ypač jei „Nvidia“ pavyks įtikinti kuo daugiau kūrėjų savo pranašumais ir perspektyvomis.

Bet, žinoma, GPU nepakeis procesorių. Dabartiniu pavidalu jie tam nėra skirti. Dabar, kai vaizdo lustai pamažu juda link procesoriaus, tampa vis universalesni (vieno ir dvigubo tikslumo slankiojo kablelio skaičiavimai, sveikųjų skaičių skaičiavimai), CPU tampa vis „lygiagretesni“, įgyja daug branduolių, kelių gijų technologijas. , jau nekalbant apie SIMD blokų ir nevienalyčių procesorių projektų atsiradimą. Labiausiai tikėtina, kad GPU ir CPU ateityje tiesiog susijungs. Yra žinoma, kad daugelis kompanijų, įskaitant „Intel“ ir AMD, dirba su panašiais projektais. Ir nesvarbu, ar CPU absorbuoja GPU, ar atvirkščiai.

Straipsnyje daugiausia kalbėjome apie CUDA naudą. Tačiau yra ir musė. Vienas iš nedaugelio CUDA trūkumų yra prastas nešiojamumas. Ši architektūra veikia tik šios kompanijos vaizdo lustuose ir ne visose, o pradedant Geforce 8 ir 9 serijomis bei atitinkamais Quadro ir Tesla. Taip, pasaulyje yra daug tokių sprendimų, „Nvidia“ nurodo 90 milijonų su CUDA suderinamų vaizdo lustų. Tai tiesiog puiku, tačiau konkurentai siūlo savo sprendimus, kurie skiriasi nuo CUDA. Taigi, AMD turi Stream Computing, o Intel ateityje turės Ct.

Kuri technologija laimės, paplis ir gyvens ilgiau nei kitos – parodys laikas. Tačiau CUDA turi daug šansų, nes, pavyzdžiui, palyginti su srautiniu skaičiavimu, ji yra labiau išvystyta ir lengviau naudojama programavimo aplinka įprasta C kalba. Galbūt trečioji šalis gali padėti apsispręsti, pateikdama kokį nors bendrą sprendimą. Pavyzdžiui, kitame „DirectX“ atnaujinime pagal 11 versiją „Microsoft“ pažadėjo skaičiavimo atspalvius, kurie gali tapti kažkokiu vidutiniu sprendimu, tinkančiu visiems arba beveik visiems.

Sprendžiant iš preliminarių duomenų, šis naujas šešėlių tipas daug pasiskolino iš CUDA modelio. O programuodami šioje aplinkoje dabar galite gauti tiesioginės naudos ir ateičiai reikalingų įgūdžių. Žvelgiant iš didelio našumo skaičiavimo perspektyvos, „DirectX“ taip pat turi aiškų trūkumą – prastą perkeliamumą, nes API apsiriboja „Windows“ platforma. Tačiau kuriamas dar vienas standartas – atvira kelių platformų iniciatyva OpenCL, kurią palaiko dauguma kompanijų, tarp jų – Nvidia, AMD, Intel, IBM ir daugelis kitų.

Nepamirškite, kad kitame straipsnyje apie CUDA bus nagrinėjami specifiniai praktiniai mokslinio ir kitokio negrafinio skaičiavimo pritaikymai, kuriuos atlieka kūrėjai iš įvairių mūsų planetos vietų naudodami Nvidia CUDA.

Šerdys CUDA - simbolis skaliariniai skaičiavimo vienetai vaizdo lustuose NVidia, pradedant nuo G 80 (GeForce 8 xxx, Tesla C-D-S870, FX4/5600 , 360 mln). Patys lustai yra architektūros dariniai. Beje, kadangi įmonė NVidia taip noriai ėmėsi savo procesorių kūrimo Tegra serija, taip pat remiantis RISC architektūra. Turiu daug patirties dirbant su šiomis architektūromis.

CUDAšerdyje yra vienas vienas vektorius Ir vienas skaliaras vienetai, atliekantys vieną vektorinę ir vieną skaliarinę operaciją per laikrodžio ciklą, perkeldami skaičiavimus į kitą daugiaprocesorių arba į kitą tolesniam apdorojimui. Šimtų ir tūkstančių tokių branduolių masyvas reiškia didelę skaičiavimo galią ir gali atlikti įvairias užduotis, priklausomai nuo reikalavimų, jei yra tam tikra pagalbinė programinė įranga. Taikymas gali būti įvairus: vaizdo srauto dekodavimas, 2D/3D grafikos pagreitinimas, debesų kompiuterija, specializuotos matematinės analizės ir kt.

Gana dažnai, kartu NVidia Tesla profesionalios kortelės Ir NVidia Quadro, yra šiuolaikinių superkompiuterių pagrindas.

CUDA— nuo to laiko branduoliai nepasikeitė G 80, bet jų skaičius didėja (kartu su kitais blokais - ROP, Tekstūros vienetai ir tt) ir lygiagrečios sąveikos efektyvumą (moduliai tobulinami Giga gija).

Pvz.:

GeForce

GTX 460 – 336 CUDA branduoliai

GTX 580 – 512 CUDA branduoliai

8800GTX – 128 CUDA branduoliai

Iš srauto procesorių skaičiaus ( CUDA), šešėlių skaičiavimų našumas didėja beveik proporcingai (tolygiai didėjant kitų elementų skaičiui).

Pradedant nuo lusto GK110(NVidia GeForce GTX 680) – CUDAšerdys turi nebe dvigubą dažnį, o bendrą su visais kitais lustų blokais. Vietoj to jų skaičius buvo padidintas maždaug triskart palyginti su ankstesne karta G110.

Nauja technologija yra tarsi naujai atsiradusi evoliucinė rūšis. Keista būtybė, kitaip nei daugelis senbuvių. Kartais nepatogu, kartais juokinga. Ir iš pradžių atrodo, kad naujos jo savybės niekaip netinka šiam nusistovėjusiam ir stabiliam pasauliui.

Tačiau praeina šiek tiek laiko, ir pasirodo, kad pradedantysis bėga greičiau, šokinėja aukščiau ir apskritai yra stipresnis. Ir jis valgo daugiau musių nei jo retrogradiniai kaimynai. Ir tada tie patys kaimynai pradeda suprasti, kad nėra prasmės ginčytis su šiuo nerangiu buvusiu. Geriau su juo draugauti, o dar geriau organizuoti simbiozę. Pamatysi, kad musių bus daugiau.

GPGPU technologija (General-Purpose Graphics Processing Units – bendrosios paskirties grafikos procesorius) ilgą laiką egzistavo tik protingų akademikų teoriniuose skaičiavimuose. Kaip kitaip? Siūlyti radikaliai pakeisti dešimtmečiais susiklosčiusį skaičiavimo procesą, jo lygiagrečių šakų skaičiavimą patikėjus vaizdo plokštei – tai sugeba tik teoretikai.

CUDA technologijos logotipas primena, kad ji augo gilumoje
3D grafika.

Tačiau GPGPU technologija neketino ilgai rinkti dulkių universitetų žurnalų puslapiuose. Išsipūtusi geriausių savo savybių plunksnas, ji patraukė gamintojų dėmesį. Taip gimė CUDA – GPGPU įgyvendinimas „nVidia“ gaminamuose „GeForce“ grafikos procesoriuose.

CUDA dėka GPGPU technologijos tapo įprasta. O dabar tik pats trumparegiškiausias ir storu tinginystės sluoksniu padengtas programavimo sistemų kūrėjas nedeklaruoja paramos CUDA savo produktu. IT leidiniai manė, kad yra garbė pristatyti technologijos detales daugybėje apkūnių mokslo populiarinimo straipsnių, o konkurentai iškart sėdo su modeliais ir kryžminiais kompiliatoriais, kad sukurtų kažką panašaus.

Viešas pripažinimas – ne tik trokštančiųjų žvaigždžių, bet ir naujai gimusių technologijų svajonė. Ir CUDA pasisekė. Ji gerai žinoma, apie ją kalba, rašo.

Jie tiesiog rašo taip, lyg ir toliau diskutuotų apie GPGPU tirštuose mokslo žurnaluose. Jie bombarduoja skaitytoją daugybe terminų, tokių kaip „tinklelis“, „SIMD“, „metimai“, „šeimininkas“, „tekstūra ir nuolatinė atmintis“. Jie panardina jį į nVidia GPU organizavimo schemas, veda vingiuotais lygiagrečių algoritmų takais ir (stipriausias žingsnis) rodo ilgus kodų sąrašus C kalba. Dėl to paaiškėja, kad įvesdami straipsnį turime naują skaitytoją, kuris degina norą suprasti CUDA, o išvestyje turime tą patį skaitytoją, bet su ištinusia galva, užpildyta faktų, diagramų netvarka. , kodas, algoritmai ir terminai.

Tuo tarpu bet kurios technologijos tikslas yra palengvinti mūsų gyvenimą. Ir CUDA puikiai dirba su tuo. Jos darbo rezultatai įtikins bet kurį skeptiką geriau nei šimtai schemų ir algoritmų.

Ne visur

CUDA palaiko didelio našumo superkompiuteriai
„nVidia Tesla“.

Ir vis dėlto, prieš žvelgiant į CUDA darbo rezultatus lengvinant paprasto vartotojo gyvenimą, verta suprasti visus jo apribojimus. Kaip ir su džinu: bet koks noras, bet vienas. CUDA taip pat turi savo Achilo kulnus. Vienas iš jų yra platformų, kuriose jis gali veikti, apribojimai.

„nVidia“ vaizdo plokščių, palaikančių CUDA, sąrašas pateikiamas specialiame sąraše, pavadintame „CUDA Enabled Products“. Sąrašas yra gana įspūdingas, tačiau jį lengva klasifikuoti. CUDA palaikymas nėra atmestas:

„nVidia GeForce“ 8, 9, 100, 200 ir 400 serijų modeliai su mažiausiai 256 megabaitų vaizdo atmintimi. Palaikymas apima ir stalinius, ir mobiliuosius korteles.

Didžioji dauguma stalinių ir mobiliųjų vaizdo plokščių yra „nVidia Quadro“.

Visi sprendimai iš nvidia ION netbook serijos.

Didelio našumo HPC (High Performance Computing) ir nVidia Tesla superkompiuterių sprendimai, naudojami tiek asmeniniams kompiuteriams, tiek keičiamo dydžio klasterių sistemoms organizuoti.

Todėl prieš naudojant CUDA pagrindu pagamintus programinės įrangos produktus, verta patikrinti šį mėgstamiausių sąrašą.

Be pačios vaizdo plokštės, norint palaikyti CUDA, reikalinga atitinkama tvarkyklė. Tai jungiamoji grandis tarp centrinio ir grafikos procesoriaus, veikianti kaip tam tikra programinės įrangos sąsaja, skirta prieiti prie programos kodo ir duomenų prie kelių branduolių GPU lobyno. Kad nesuklystumėte, „nVidia“ rekomenduoja apsilankyti tvarkyklių puslapyje ir gauti naujausią versiją.

...bet pats procesas

Kaip veikia CUDA? Kaip paaiškinti sudėtingą lygiagretaus skaičiavimo procesą naudojant specialią GPU aparatinės įrangos architektūrą, neįtraukiant skaitytojo į konkrečių terminų bedugnę?

Galite pabandyti tai padaryti įsivaizduodami, kaip centrinis procesorius vykdo programą simbiozėje su grafikos procesoriumi.

Architektūriniu požiūriu centrinis procesorius (CPU) ir jo grafinis atitikmuo (GPU) yra suprojektuoti skirtingai. Jei pateiksime analogiją su automobilių pramonės pasauliu, tada CPU yra universalas, vienas iš tų, kurie vadinami „tvartu“. Tai atrodo kaip lengvasis automobilis, bet tuo pat metu (kūrėjų požiūriu) „tai šveicaras, javapjūtė ir grotuvas ant vamzdžio“. Vienu metu atlieka mažo sunkvežimio, autobuso ir hipertrofuoto hečbeko vaidmenį. Trumpai tariant, universalas. Jame yra nedaug cilindrų branduolių, tačiau jie susidoroja su beveik bet kokia užduotimi, o įspūdinga talpyklos atmintis gali saugoti krūvą duomenų.

Tačiau GPU yra sportinis automobilis. Yra tik viena funkcija: kuo greičiau pristatyti pilotą į finišą. Todėl nėra didelės bagažinės atminties, jokių papildomų sėdynių. Tačiau cilindrų branduolių yra šimtus kartų daugiau nei procesoriaus.

CUDA dėka GPGPU programų kūrėjams nereikia gilintis į programavimo sudėtingumą
grafikos variklių, tokių kaip „DirectX“ ir „OpenGL“, kūrimas

Skirtingai nuo centrinio procesoriaus, kuris gali išspręsti bet kokią užduotį, įskaitant grafiką, bet su vidutiniu našumu, grafikos procesorius yra pritaikytas didelės spartos vienos užduoties sprendimui: daugiakampių krūva įvestyje paverčiama pikselių krūva. išvestis. Be to, šią problemą galima išspręsti lygiagrečiai naudojant šimtus gana paprastų skaičiavimo branduolių GPU.

Taigi koks gali būti tandemas iš universalo ir sportinio automobilio? CUDA veikia maždaug taip: programa veikia CPU, kol atsiranda kodo dalis, kurią galima vykdyti lygiagrečiai. Tada, užuot lėtai vykdomas dviejuose (ar net aštuoniuose) šauniausio procesoriaus branduoliuose, jis perkeliamas į šimtus GPU branduolių. Tuo pačiu metu šios dalies vykdymo laikas žymiai sutrumpėja, o tai reiškia, kad sutrumpėja ir visos programos vykdymo laikas.

Technologiškai programuotojui niekas nesikeičia. CUDA programų kodas parašytas C kalba. Tiksliau, savo specialia tarme „C su upeliais“ (C su upeliais). Šis C kalbos plėtinys, sukurtas Stanforde, vadinamas Brook. Sąsaja, perduodanti Brook kodą į GPU, yra vaizdo plokštės, kuri palaiko CUDA, tvarkyklė. Jis organizuoja visą šios programos dalies apdorojimo procesą taip, kad programuotojui GPU atrodytų kaip procesoriaus koprocesorius. Labai panašus į matematikos koprocesoriaus naudojimą asmeninio skaičiavimo pradžioje. Atsiradus Brook, vaizdo plokštėms su CUDA palaikymu ir joms skirtomis tvarkyklėmis, bet kuris programuotojas galėjo pasiekti savo programose esantį GPU. Tačiau anksčiau šis šamanizmas priklausė siauram atrinktų žmonių ratui, kurie ilgus metus tobulino „DirectX“ arba „OpenGL“ grafikos variklių programavimo metodus.

Į šio pretenzingo medaus statinę – CUDA pagyros – verta įmesti musę, tai yra apribojimus. Ne kiekviena problema, kurią reikia užprogramuoti, gali būti išspręsta naudojant CUDA. Įprastų biuro užduočių sprendimo paspartinti nepavyks, tačiau galite patikėti, kad CUDA apskaičiuos tūkstančių to paties tipo naikintuvų elgesį „World of Warcraft“. Bet tai yra sugalvota užduotis. Pažvelkime į pavyzdžius, ką CUDA jau išsprendžia labai efektyviai.

Teisingi darbai

CUDA yra labai pragmatiška technologija. Įdiegusi savo palaikymą savo vaizdo plokštėse, nVidia visiškai pagrįstai tikėjosi, kad CUDA reklamjuostę paims daug entuziastų tiek universiteto aplinkoje, tiek komercijoje. Taip ir atsitiko. CUDA pagrįsti projektai gyvuoja ir duoda naudos.

NVIDIA PhysX

Reklamuodami savo kitą žaidimų šedevrą, gamintojai dažnai pabrėžia jo 3D tikroviškumą. Bet kad ir koks realus būtų 3D žaidimų pasaulis, jei elementarūs fizikos dėsniai, tokie kaip gravitacija, trintis ir hidrodinamika, bus įgyvendinami neteisingai, melas pasijus akimirksniu.

Viena iš NVIDIA PhysX fizinio variklio galimybių yra realus darbas su audiniais.

Pagrindinių fizinių dėsnių kompiuterinio modeliavimo algoritmų įgyvendinimas yra labai daug darbo reikalaujanti užduotis. Žymiausios šios srities įmonės yra Airijos kompanija Havok su savo kelių platformų fizine Havok Physics ir Kalifornijos Ageia – pirmojo pasaulyje fizinio procesoriaus (PPU – Physics Processing Unit) ir atitinkamo PhysX fizinio variklio pirmtakas. Pirmasis iš jų, nors ir įsigijo Intel, dabar aktyviai dirba Havok variklio optimizavimo ATI vaizdo plokštėms ir AMD procesoriams srityje. Tačiau „Ageia“ su savo „PhysX“ varikliu tapo „nVidia“ dalimi. Tuo pačiu metu „nVidia“ išsprendė gana sudėtingą „PhysX“ pritaikymo CUDA technologijai problemą.

Tai tapo įmanoma statistikos dėka. Statistiškai įrodyta, kad nesvarbu, koks sudėtingas GPU atvaizdavimas, kai kurie jo branduoliai vis tiek neveikia. Būtent ant šių branduolių veikia PhysX variklis.

CUDA dėka vaizdo plokštėje buvo pradėta atlikti liūto dalis su žaidimų pasaulio fizika susijusių skaičiavimų. Atlaisvinta centrinio procesoriaus galia buvo panaudota kitoms žaidimo problemoms spręsti. Rezultatas netruko laukti. Pasak ekspertų, žaidimo našumas, kai „PhysX“ veikia su CUDA, padidėjo bent eilės tvarka. Taip pat padidėjo tikimybė suvokti fizikinius dėsnius. CUDA rūpinasi įprastiniu daugiamačių objektų trinties, gravitacijos ir kitų mums žinomų dalykų skaičiavimu. Dabar į mums pažįstamo fizinio pasaulio dėsnius puikiai dera ne tik herojai ir jų įranga, bet ir dulkės, rūkas, sprogimo banga, liepsna ir vanduo.

NVIDIA Texture Tools 2 tekstūros glaudinimo paketo CUDA versija

Ar jums patinka realistiški objektai šiuolaikiniuose žaidimuose? Verta padėkoti tekstūros kūrėjams. Tačiau kuo daugiau tikrovės tekstūroje, tuo didesnė jos apimtis. Kuo daugiau tai užima brangią atmintį. Kad to išvengtumėte, tekstūros iš anksto suglaudinamos ir, jei reikia, dinamiškai išspaudžiamos. O suspaudimas ir dekompresija yra gryni skaičiavimai. Norėdami dirbti su tekstūromis, „nVidia“ išleido „NVIDIA Texture Tools“ paketą. Jis palaiko efektyvų „DirectX“ tekstūrų (vadinamojo HF formato) suspaudimą ir išskleidimą. Antroji šio paketo versija gali pasigirti BC4 ir BC5 glaudinimo algoritmų, įdiegtų DirectX 11 technologijoje, palaikymu. Tačiau svarbiausia, kad NVIDIA Texture Tools 2 apima CUDA palaikymą. Pasak nVidia, tai 12 kartų padidina tekstūros glaudinimo ir išskleidimo užduočių našumą. Tai reiškia, kad žaidimo kadrai bus įkeliami greičiau ir džiugins žaidėją savo tikroviškumu.

NVIDIA Texture Tools 2 paketas sukurtas dirbti su CUDA. Našumo padidėjimas suspaudžiant ir išspaudžiant tekstūras yra akivaizdus.

Naudojant CUDA galima žymiai pagerinti vaizdo stebėjimo efektyvumą.

Vaizdo srauto apdorojimas realiuoju laiku

Kad ir ką sakytume, dabartinis pasaulis šnipinėjimo požiūriu yra daug artimesnis Orwello didžiojo brolio pasauliui, nei atrodo. Vaizdo kamerų žvilgsnį jaučia ir automobilių vairuotojai, ir viešųjų vietų lankytojai.

Pilnos tekančios vaizdo informacijos upės teka į jos apdorojimo centrus ir... subėga į siaurą grandį – žmogų. Daugeliu atvejų jis yra paskutinė institucija, stebinti vaizdo pasaulį. Be to, valdžia nėra pati veiksmingiausia. Mirksi, blaškosi ir bando užmigti.

CUDA dėka tapo įmanoma įdiegti algoritmus, skirtus vienu metu sekti kelis objektus vaizdo sraute. Šiuo atveju procesas vyksta realiu laiku, o vaizdo įrašas yra pilnas 30 kadrų per sekundę. Palyginti su tokio algoritmo įgyvendinimu šiuolaikiniuose kelių branduolių procesoriuose, CUDA našumą padidina du ar tris kartus, o tai, matote, yra gana daug.

Vaizdo konvertavimas, garso filtravimas

Badaboom vaizdo konverteris yra pirmasis, kuris naudoja CUDA, kad pagreitintų konversiją.

Smagu žiūrėti naują vaizdo nuomos produktą FullHD kokybe ir dideliame ekrane. Tačiau į kelią negalite pasiimti didelio ekrano, o „FullHD“ vaizdo kodekas užklups mažos galios mobiliosios programėlės procesorių. Atsivertimas ateina į pagalbą. Tačiau dauguma tų, kurie su tuo susidūrė praktiškai, skundžiasi ilgu konversijos laiku. Tai suprantama, procesas yra įprastas, tinka lygiagretavimui, o jo vykdymas CPU nėra labai optimalus.

Tačiau CUDA su tuo susidoroja su kaupu. Pirmasis ženklas yra „Badaboom“ keitiklis iš „Elevental“. Badaboom kūrėjai pasirinko CUDA teisingai. Testai rodo, kad jis paverčia standartinį pusantros valandos filmą į iPhone/iPod Touch formatą greičiau nei per dvidešimt minučių. Ir tai nepaisant to, kad naudojant tik centrinį procesorių, šis procesas trunka ilgiau nei valandą.

Padeda CUDA ir profesionaliems audiofilams. Bet kuris iš jų duotų pusę karalystės efektyviam FIR krosoveriui – filtrų rinkiniui, dalijančiam garso spektrą į kelias juostas. Šis procesas yra labai daug darbo reikalaujantis ir, esant dideliam garso medžiagos kiekiui, verčia garso inžinierių kelias valandas eiti „rūkyti“. CUDA pagrindu sukurto FIR krosoverio įdiegimas šimtus kartų pagreitina jo veikimą.

CUDA ateitis

GPGPU technologiją pavertusi realybe, CUDA neužmigo ant laurų. Kaip ir visur, CUDA veikia atspindžio principas: dabar CUDA SDK versijų kūrimui įtakos turi ne tik nVidia vaizdo procesorių architektūra, bet pati CUDA technologija verčia nVidia persvarstyti savo lustų architektūrą. Tokio atspindžio pavyzdys yra nVidia ION platforma. Antroji jo versija yra specialiai optimizuota CUDA problemoms spręsti. Tai reiškia, kad net ir santykinai nebrangiuose aparatinės įrangos sprendimuose vartotojai gaus visą CUDA galią ir puikias galimybes.

CUDA technologija

Vladimiras Frolovas,[apsaugotas el. paštas]

anotacija

Straipsnyje kalbama apie CUDA technologiją, kuri leidžia programuotojui naudoti vaizdo plokštes kaip galingus skaičiavimo įrenginius. „Nvidia“ teikiami įrankiai leidžia rašyti grafikos apdorojimo bloko (GPU) programas C++ kalbos pogrupyje. Tai atleidžia programuotoją nuo būtinybės naudoti atspalvius ir suprasti grafikos dujotiekio veikimą. Straipsnyje pateikiami programavimo pavyzdžiai naudojant CUDA ir įvairias optimizavimo technikas.

1. Įvadas

Per pastaruosius dešimtmečius kompiuterinių technologijų raida sparčiai vystėsi. Taip greitai, kad procesorių kūrėjai jau beveik pasiekė vadinamąją „silicio aklavietę“. Nežabotas laikrodžio dažnio didinimas tapo neįmanomas dėl daugelio rimtų technologinių priežasčių.

Iš dalies dėl to visi šiuolaikinių skaičiavimo sistemų gamintojai žengia link procesorių ir branduolių skaičiaus didinimo, o ne vieno procesoriaus dažnio didinimo. Centrinio procesoriaus (CPU) branduolių skaičius pažangiose sistemose dabar jau yra 8.

Kita priežastis – palyginti mažas RAM greitis. Kad ir kaip greitai dirbtų procesorius, kliūtis, kaip rodo praktika, yra visai ne aritmetinės operacijos, o veikiau nesėkmingos atminties prieigos – talpyklos praleidimai.

Tačiau jei pažvelgtumėte į GPU (Graphics Processing Unit) grafikos procesorius, jie paralelizmo keliu pasuko daug anksčiau. Šiuolaikinėse vaizdo plokštėse, pavyzdžiui, GF8800GTX, procesorių skaičius gali siekti 128. Tokių sistemų našumas, sumaniai programuojant, gali būti gana reikšmingas (1 pav.).

Ryžiai. 1. CPU ir GPU slankiojo kablelio operacijų skaičius

Kai pirmosios vaizdo plokštės pasirodė rinkoje, jos buvo gana paprasti (palyginti su centriniu procesoriumi) labai specializuoti įrenginiai, skirti atleisti procesorių nuo dvimačių duomenų vizualizavimo naštos. Tobulėjant žaidimų industrijai ir atsiradus tokiems trimačiams žaidimams kaip Doom (2 pav.) ir Wolfenstein 3D (3 pav.), atsirado poreikis 3D vizualizacijai.

2,3 pav. Žaidimai Doom ir Wolfenstein 3D

Nuo to laiko, kai 3Dfx sukūrė pirmąsias Voodoo vaizdo plokštes (1996 m.) iki 2001 m., GPU buvo įdiegtas tik fiksuotas įvesties duomenų operacijų rinkinys.

Programuotojai neturėjo pasirinkimo atvaizdavimo algoritme, o lankstumui padidinti atsirado šešėliai – mažos programos, vykdomos vaizdo plokštės kiekvienai viršūnei arba kiekvienam pikseliui. Jų užduotys apėmė viršūnių transformacijas ir šešėliavimą – apskaičiuojant apšvietimą taške, pavyzdžiui, naudojant Phong modelį.

Nors šiais laikais šešėliai nuėjo ilgą kelią, reikia suprasti, kad jie buvo sukurti labai specializuotoms užduotims, tokioms kaip 3D transformacija ir rastravimas. Nors GPU tobulėja prie bendros paskirties kelių procesorių sistemų, šešėlių kalbos išlieka labai specializuotos.

Juos galima palyginti su FORTRAN ta prasme, kad, kaip ir FORTRAN, jie buvo pirmieji, tačiau skirti išspręsti tik vieno tipo problemas. Shader'iai mažai naudingi sprendžiant kitas problemas, išskyrus 3D transformacijas ir rastravimą, kaip ir FORTRAN nėra patogus spręsti problemas, nesusijusias su skaitiniais skaičiavimais.

Šiandien vyrauja tendencija netradiciniu būdu naudoti vaizdo plokštes sprendžiant kvantinės mechanikos, dirbtinio intelekto, fizikinių skaičiavimų, kriptografijos, fiziškai teisingos vizualizacijos, rekonstrukcijos iš nuotraukų, atpažinimo ir kt. Šias užduotis nepatogu išspręsti naudojant grafines API (DirectX, OpenGL), nes šios API buvo sukurtos visiškai skirtingoms programoms.

Bendrojo programavimo ant GPU (General Programming on GPU, GPGPU) kūrimas logiškai paskatino technologijų, skirtų platesniam užduočių spektrui nei rastravimas, atsiradimą. Dėl to „Nvidia“ sukūrė „Compute Unified Device Architecture“ (arba trumpiau CUDA), o konkuruojanti įmonė ATI sukūrė STREAM technologiją.

Pažymėtina, kad šio straipsnio rašymo metu STREAM technologija gerokai atsiliko nuo CUDA plėtojant, todėl čia ji nebus svarstoma. Daugiausia dėmesio skirsime CUDA – GPGPU technologijai, kuri leidžia rašyti programas C++ kalbos pogrupyje.

2. Esminis skirtumas tarp procesoriaus ir GPU

Greitai pažvelkime į keletą reikšmingų procesoriaus ir vaizdo plokščių programų sričių ir funkcijų skirtumų.

2.1. Galimybės

Iš pradžių CPU skirtas bendroms problemoms spręsti ir veikia su atsitiktinai adresuojama atmintimi. CPU programos gali tiesiogiai pasiekti bet kokias linijines ir vienarūšes atminties ląsteles.

Tai netaikoma GPU atveju. Kaip sužinosite perskaitę šį straipsnį, CUDA turi net 6 atminties tipus. Galite skaityti iš bet kurio langelio, kuris yra fiziškai pasiekiamas, bet galite rašyti ne į visus langelius. Priežastis ta, kad GPU bet kuriuo atveju yra konkretus įrenginys, sukurtas konkretiems tikslams. Šis apribojimas buvo įvestas siekiant padidinti tam tikrų algoritmų greitį ir sumažinti įrangos kainą.

2.2. Atminties našumas

Daugumos kompiuterių sistemų nuolatinė problema yra ta, kad atmintis yra lėtesnė nei procesorius. CPU gamintojai šią problemą išsprendžia įvesdami talpyklas. Dažniausiai naudojamos atminties sritys dedamos į atmintį arba talpyklą, veikiančią procesoriaus dažniu. Tai leidžia sutaupyti laiko pasiekiant dažniausiai naudojamus duomenis ir apkrauti procesorių su faktiniais skaičiavimais.

Atminkite, kad talpyklos programuotojui iš esmės yra skaidrios. Tiek skaitant, tiek rašant duomenys patenka ne tiesiai į RAM, o per talpyklas. Tai ypač leidžia greitai perskaityti reikšmę iškart ją parašius.

GPU (čia turime galvoje GF aštuntos serijos vaizdo plokštes) taip pat turi talpyklas, jos taip pat yra svarbios, tačiau šis mechanizmas nėra toks galingas kaip CPU. Pirma, ne visų tipų atmintis yra talpykloje, antra, talpyklos yra tik skaitymo.

GPU lėtos prieigos prie atminties paslėptos naudojant lygiagretųjį skaičiavimą. Kol vienos užduotys laukia duomenų, kitos dirba, pasiruošusios skaičiavimams. Tai vienas iš pagrindinių CUDA principų, galinčių labai pagerinti visos sistemos veikimą.

3. CUDA šerdis

3.1. Srautinio perdavimo modelis

CUDA skaičiavimo architektūra yra pagrįsta koncepcijaviena komanda daugeliui duomenų(Single Instruction Multiple Data, SIMD) ir koncepcija daugiaprocesorius.

SIMD koncepcija reiškia, kad viena instrukcija gali apdoroti kelis duomenis vienu metu. Pavyzdžiui, komanda addps Pentium 3 ir naujesniuose Pentium procesoriuose leidžia vienu metu pridėti 4 vieno tikslumo slankiojo kablelio skaičius.

Daugiaprocesorius yra kelių branduolių SIMD procesorius, leidžiantis bet kuriuo metu visuose branduoliuose vykdyti tik vieną komandą. Kiekviena kelių procesorių šerdis yra skaliarinė, t.y. jis nepalaiko grynųjų vektorinių operacijų.

Prieš tęsdami, pristatykime keletą apibrėžimų. Atkreipkite dėmesį, kad šiame straipsnyje mes turėsime omenyje įrenginį ir pagrindinį kompiuterį visiškai kitaip nei dauguma programuotojų. Mes naudosime tokias sąlygas, kad išvengtume neatitikimų su CUDA dokumentacija.

Mūsų straipsnyje, pagal įrenginį, turėsime omenyje vaizdo adapterį, palaikantį CUDA tvarkyklę, arba kitą specializuotą įrenginį, skirtą CUDA naudojančioms programoms vykdyti (pvz., NVIDIA Tesla). Savo straipsnyje mes laikysime GPU tik kaip loginį įrenginį, vengdami konkrečių diegimo detalių.

Priegloba vadinsime įprastoje kompiuterio RAM esančią programą, kuri naudoja procesorių ir atlieka darbo su įrenginiu valdymo funkcijas.

Tiesą sakant, jūsų programos dalis, kuri veikia CPUšeimininkas, ir tavo vaizdo plokštė - prietaisas. Logiškai mąstant, įrenginį galima pavaizduoti kaip kelių procesorių rinkinį (4 pav.) ir CUDA tvarkyklę.

Ryžiai. 4. Įrenginys

Tarkime, kad norime paleisti tam tikrą procedūrą savo įrenginyje N gijų (tai yra, norime lygiagretinti jos darbą). Pagal CUDA dokumentaciją šią procedūrą pavadinkime branduoliu.

CUDA architektūros ypatybė yra jos blokų tinklelio struktūra, kuri yra neįprasta kelių gijų programoms (5 pav.). Šiuo atveju CUDA tvarkyklė savarankiškai paskirsto įrenginio išteklius tarp gijų.

Ryžiai. 5. Gijų organizavimas

Fig. 5. Branduolys yra pažymėtas kaip branduolys. Visos gijos, vykdančios šį branduolį, yra sujungtos į blokus (Block), o blokai, savo ruožtu, yra sujungti į tinklelį (Grid).

Kaip matyti 5 paveiksle, siūlams identifikuoti naudojami dvimačiai indeksai. CUDA kūrėjai suteikė galimybę dirbti su trimačiais, dvimačiais arba paprastais (vienmačiais) indeksais, priklausomai nuo to, kas patogiau programuotojui.

Apskritai indeksai yra trimačiai vektoriai. Kiekvienai gijai bus žinoma: gijos indeksas bloko threadIdx viduje ir bloko indeksas tinklelio blockIdx viduje. Paleidžiant visos gijos skirsis tik šiais indeksais. Tiesą sakant, būtent per šiuos indeksus programuotojas vykdo kontrolę, nustatydamas, kuri jo duomenų dalis yra apdorojama kiekvienoje gijoje.

Atsakymas į klausimą, kodėl kūrėjai pasirinko būtent šią organizaciją, yra nebanalus. Viena iš priežasčių yra ta, kad vienas blokas garantuojamas ant vieno įrenginio kelių procesorių, tačiau vienas daugiaprocesorius gali vykdyti kelis skirtingus blokus. Likusios priežastys paaiškės vėliau straipsnyje.

Užduočių (gijų) blokas vykdomas daugiaprocesoriuje dalimis arba telkiniuose, vadinamuose metmenyse. Dabartinis metmenų dydis vaizdo plokštėse su CUDA palaikymu yra 32 gijos. Užduotys metmenų baseino viduje vykdomos SIMD stiliumi, t.y. Visos metmenų gijos vienu metu gali vykdyti tik vieną komandą.

Čia reikia padaryti vieną įspėjimą. Šiuolaikinėse architektūrose šio straipsnio rašymo metu procesorių skaičius viename daugiaprocesoriuje yra 8, o ne 32. Iš to išplaukia, kad ne visas metimas vykdomas vienu metu, jis padalintas į 4 dalis, kurios vykdomos nuosekliai (kadangi procesoriai yra skaliariniai).

Tačiau, pirma, CUDA kūrėjai griežtai nereglamentuoja metmenų dydžio. Savo darbuose jie mini metmenų dydžio parametrą, o ne skaičių 32. Antra, loginiu požiūriu metmenys yra minimali gijų sąjunga, apie kurią galima teigti, kad visos jos viduje esančios gijos vykdomos vienu metu – ir Tuo pačiu metu nėra prielaidų, kad sistema nebus atlikta.

3.1.1. Išsišakojimas

Iš karto kyla klausimas: jei tuo pačiu metu visi siūlai metmenų viduje vykdo tą pačią instrukciją, tai kaip su šakomis? Galų gale, jei programos kode yra šaka, tada instrukcijos bus kitokios. Čia naudojamas standartinis SIMD programavimo sprendimas (6 pav.).

Ryžiai. 6. Filialų organizavimas SIMD

Tarkime, kad turime šį kodą:

jei (sąlyga)B;

SISD (Single Instruction Single Data) atveju vykdome sakinį A, patikriname sąlygą, tada vykdome sakinius B ir D (jei sąlyga teisinga).

Dabar turime 10 gijų, veikiančių SIMD stiliumi. Visose 10 gijų vykdome teiginį A, tada patikriname sąlygą cond ir paaiškėja, kad 9 iš 10 gijų tai teisinga, o vienoje gijoje klaidinga.

Aišku, kad negalime paleisti 9 gijų, kad vykdytume operatorių B, o likusios – operatoriaus C, nes visose gijose vienu metu gali būti vykdoma tik viena instrukcija. Tokiu atveju turite tai padaryti: pirmiausia „nužudome“ nutrūkusią giją, kad ji nesugadintų niekieno duomenų, ir vykdome likusias 9 gijas. Tada „nužudome“ 9 gijas, kurios įvykdė operatorių B, ir pereiname per vieną giją su operatoriumi C. Po to gijos vėl susilieja ir vienu metu vykdo operatorių D.

Liūdnas rezultatas yra tai, kad procesoriaus ištekliai ne tik švaistomi tuščiam bitų šlifavimui suskaidytose gijose, bet ir daug blogiau yra tai, kad galiausiai turėsime vykdyti ABEJAS šakas.

Tačiau ne viskas taip blogai, kaip gali pasirodyti iš pirmo žvilgsnio. Labai didelis technologijos privalumas yra tas, kad šiuos triukus CUDA tvarkyklė atlieka dinamiškai ir jie yra visiškai skaidrūs programuotojui. Tuo pačiu metu, dirbdamas su šiuolaikinių procesorių SSE komandomis (ypač tuo atveju, kai vienu metu bandoma vykdyti 4 algoritmo kopijas), pats programuotojas turi pasirūpinti smulkmenomis: sujungti duomenis į keturračius, nepamiršti derinimo. , ir apskritai rašyti žemu lygiu, tiesą sakant, kaip surinkime.

Iš viso to, kas išdėstyta pirmiau, išplaukia viena labai svarbi išvada. Šakos savaime nesukelia našumo pablogėjimo. Vienintelės žalingos šakos yra tos, kuriose siūlai išsiskiria tame pačiame metmenų siūlų telkinyje. Be to, jei siūlai skiriasi viename bloke, bet skirtinguose metmenų telkiniuose arba skirtinguose blokuose, tai neturi jokios įtakos.

3.1.2. Ryšys tarp gijų

Šio straipsnio rašymo metu bet kokia sąveika tarp gijų (sinchronizavimas ir duomenų mainai) buvo įmanoma tik bloke. Tai yra, neįmanoma organizuoti sąveikos tarp skirtingų blokų gijų naudojant tik dokumentuotas galimybes.

Kalbant apie nedokumentuotas funkcijas, labai nerekomenduojama jomis naudotis. Taip yra todėl, kad jie priklauso nuo konkrečios sistemos aparatinės įrangos savybių.

Visų bloko užduočių sinchronizavimas atliekamas iškviečiant funkciją __synchtreads. Keitimasis duomenimis galimas naudojant bendrą atmintį, nes tai bendra visoms bloko užduotims.

3.2. Atmintis

CUDA išskiria šešis atminties tipus (7 pav.). Tai registrai, vietinė, globali, bendra, pastovi ir tekstūrinė atmintis.

Tokią gausą lemia vaizdo plokštės specifika ir jos pirminė paskirtis, taip pat kūrėjų noras padaryti sistemą kuo pigesnę, įvairiais atvejais aukojant arba universalumą, arba greitį.

Ryžiai. 7. CUDA atminties tipai

3.2.0. Registrai

Kai tik įmanoma, kompiliatorius stengiasi visus vietinius funkcijų kintamuosius patalpinti į registrus. Tokie kintamieji pasiekiami maksimaliu greičiu. Dabartinėje architektūroje vienam daugiaprocesoriui yra 8192 32 bitų registrai. Norėdami nustatyti, kiek registrų yra vienai gijai, turite padalyti šį skaičių (8192) iš bloko dydžio (gijų skaičiaus jame).

Įprastai padalijus 64 gijas viename bloke, rezultatas yra tik 128 registrai (yra keletas objektyvių kriterijų, tačiau 64 vidutiniškai tinka daugeliui užduočių). Iš tikrųjų nvcc niekada neskirs 128 registrų. Paprastai jis neduoda daugiau nei 40, o likę kintamieji pateks į vietinę atmintį. Taip nutinka todėl, kad viename daugiaprocesoriuje gali būti vykdomi keli blokai. Kompiliatorius bando maksimaliai padidinti vienu metu veikiančių blokų skaičių. Norėdami padidinti efektyvumą, turėtumėte pabandyti užimti mažiau nei 32 registrus. Tada teoriškai viename daugiaprocesoriuje galima paleisti 4 blokus (8 metmenys, jei viename bloke 64 gijos). Tačiau čia taip pat turėtumėte atsižvelgti į bendrinamos atminties kiekį, kurį užima gijos, nes jei vienas blokas užima visą bendrinamą atmintį, du tokie blokai negali būti vykdomi keliuose procesoriuose vienu metu.

3.2.1. Vietinė atmintis

Tais atvejais, kai vietinės procedūros duomenys yra per dideli arba kompiliatorius negali apskaičiuoti tam tikro nuolatinės prieigos žingsnio, jis gali juos įrašyti į vietinę atmintį. Tai galima palengvinti, pavyzdžiui, užmetant skirtingų dydžių tipų rodykles.

Fiziškai vietinė atmintis yra analogiška globaliai ir veikia tokiu pat greičiu. Rašymo metu nėra mechanizmų, kurie aiškiai neleistų kompiliatoriui naudoti vietinę atmintį konkretiems kintamiesiems. Kadangi vietinę atmintį valdyti gana sunku, geriau jos visai nenaudoti (žr. 4 skyrių „Optimizavimo rekomendacijos“).

3.2.2. Pasaulinė atmintis

CUDA dokumentacijoje kaip vienas iš pagrindinių laimėjimųtechnologija suteikia galimybę savavališkai kreiptis į pasaulinę atmintį. Tai reiškia, kad galite skaityti iš bet kurios atminties ląstelės, taip pat galite rašyti į savavališką langelį (dažniausiai taip nėra GPU).

Tačiau šiuo atveju už universalumą reikia mokėti greičiu. Pasaulinė atmintis nėra talpykloje. Jis veikia labai lėtai, prieigų prie pasaulinės atminties skaičius bet kokiu atveju turėtų būti sumažintas.

Visuotinė atmintis daugiausia reikalinga norint išsaugoti programos rezultatus prieš siunčiant juos pagrindiniam kompiuteriui (įprastoje DRAM). Taip yra todėl, kad globali atmintis yra vienintelė atmintis, kurioje galite rašyti bet ką.

Kintamieji, deklaruoti su __global__ kvalifikatoriumi, yra priskiriami visuotinei atminčiai. Visuotinę atmintį taip pat galima paskirstyti dinamiškai, pagrindiniame kompiuteryje iškviečiant cudaMalloc(void* mem, int size). Šios funkcijos negalima iškviesti iš įrenginio. Iš to išplaukia, kad atminties paskirstymą turėtų tvarkyti pagrindinė programa, veikianti CPU. Duomenys iš pagrindinio kompiuterio gali būti siunčiami į įrenginį iškvietus funkciją cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int dydis, cudaMemcpyHostToDevice);

Atvirkštinę procedūrą galite atlikti lygiai taip pat:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int dydis, cudaMemcpyDeviceToHost);

Šį skambutį taip pat atlieka šeimininkas.

Dirbant su globalia atmintimi, svarbu laikytis sujungimo taisyklės. Pagrindinė idėja yra ta, kad gijos turėtų pasiekti nuoseklias 4, 8 arba 16 baitų atminties ląsteles. Tokiu atveju pati pirmoji gija turi pasiekti adresą, suderintą su atitinkamai 4, 8 arba 16 baitų riba. CudaMalloc grąžinti adresai yra sulygiuoti ant mažiausiai 256 baitų ribos.

3.2.3. Bendra atmintis

Bendroji atmintis nėra talpinama, bet greita. Rekomenduojama naudoti kaip valdomą talpyklą. Vienam kelių procesorių bendrai naudojama atmintis yra tik 16 KB. Padalijus šį skaičių iš užduočių skaičiaus bloke, gauname didžiausią bendrinamos atminties kiekį vienai gijai (jei planuojate ją naudoti atskirai visose gijose).

Skiriamasis bendrosios atminties bruožas yra tas, kad ji vienodai skirta visoms užduotims bloke (7 pav.). Iš to išplaukia, kad jis gali būti naudojamas keistis duomenimis tik tarp vieno bloko gijų.

Garantuojama, kad bloko vykdymo metu daugiaprocesoriuje bendrosios atminties turinys bus išsaugotas. Tačiau po kelių procesorių bloko pakeitimo negarantuojama, kad senojo bloko turinys bus išsaugotas. Todėl neturėtumėte bandyti sinchronizuoti užduočių tarp blokų, palikdami bet kokius duomenis bendroje atmintyje ir tikėdamiesi jų saugumo.

Kintamieji, deklaruoti su __shared__ kvalifikatoriumi, yra paskirstomi bendrojoje atmintyje.

Shared__ float mem_shared;

Dar kartą reikia pabrėžti, kad blokui yra tik viena bendra atmintis. Todėl, jei jums reikia jį naudoti tiesiog kaip valdomą talpyklą, turėtumėte pasiekti skirtingus masyvo elementus, pavyzdžiui, taip:

float x = mem_shared;

Kur threadIdx.x yra bloko viduje esančios gijos x indeksas.

3.2.4. Nuolatinė atmintis

Nuolatinė atmintis yra talpykloje, kaip matyti Fig. 4. Talpykla yra vienoje kopijoje vienam kelių procesorių, o tai reiškia, kad ji yra bendra visoms bloko užduotims. Pagrindiniame kompiuteryje galite ką nors įrašyti į nuolatinę atmintį iškviesdami funkciją cudaMemcpyToSymbol. Nuolatinė įrenginio atmintis yra tik skaitoma.

Nuolatinė atmintis yra labai patogu naudoti. Į jį galite įdėti bet kokio tipo duomenis ir juos perskaityti naudodami paprastą užduotį.

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ reiškia, kad įrenginio_kernel yra branduolys, kurį galima paleisti GPU

Global__ negalioja įrenginio_kernel()

int a = gpu_buferis;

int b = gpu_buferis + gpu_buferis;

// gpu_buffer = a; KLAIDA! nuolatinė atmintis yra tik skaitymo

Kadangi nuolatinė atmintis naudoja talpyklą, prieiga prie jos paprastai yra gana greita. Vienintelis, bet labai didelis pastovios atminties trūkumas yra tai, kad jos dydis yra tik 64 KB (visam įrenginiui). Tai rodo, kad tikslinga kontekstinėje atmintyje saugoti tik nedidelį dažnai naudojamų duomenų kiekį.

3.2.5. Tekstūros atmintis

Tekstūros atmintis yra talpykloje (4 pav.). Kiekvienam daugiaprocesoriui yra tik viena talpykla, o tai reiškia, kad ši talpykla yra bendra visoms bloko užduotims.

Tekstūros atminties pavadinimas (ir, deja, funkcionalumas) yra paveldėtas iš sąvokų „tekstūra“ ir „tekstūravimas“. Tekstūravimas – tai tekstūros (tik paveikslėlio) pritaikymas daugiakampiui rastravimo proceso metu. Tekstūros atmintis yra optimizuota 2D duomenų atrankai ir turi šias galimybes:

greitas fiksuoto dydžio verčių (baitų, žodžių, dvigubų ar keturių žodžių) gavimas iš vienmačio ar dvimačio masyvo;

normalizuotas adresavimas su slankiaisiais skaičiais intervale . Tada galite juos pasirinkti naudodami normalizuotą adresą. Gauta reikšmė bus float4 tipo žodis, susietas su intervalu ;

CudaMalloc((void**) &gpu_memory, N*dydis (uint4 )); // paskirstyti atmintį GPU

// tekstūros parametrų nustatymas tekstūra

Texture.addressMode = cudaAddressModeWrap; // režimu Apvyniokite

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //artimiausią vertę

Tekstūra.normalizuota = false; // nenaudokite normalizuoto adresavimo

CudaBindTexture(0, tekstūra, gpu_memory, N ) // nuo šiol ši atmintis bus laikoma tekstūrine atmintimi


CudaMemcpy(gpu_memory, cpu_buffer, N*sizeof(uint 4), cudaMemcpyHostToDevice ); // nukopijuokite duomenis įGPU

// __global__ reiškia, kad įrenginio_kernel yra lygiagretinamas branduolys

Global__ negalioja įrenginio_kernel()

uint4 a = tex1Dfetch(tekstūra,0); // duomenis galite pasirinkti tik tokiu būdu!

uint4 b = tex1Dfetch(tekstūra,1);

int c = a.x * b.y;

...

3.3. Paprastas pavyzdys

Kaip paprastą pavyzdį apsvarstykite programą cppIntegration iš CUDA SDK. Jame demonstruojami darbo su CUDA metodai, taip pat nvcc (specialus Nvidia C++ kompiliatoriaus pogrupis) naudojimas kartu su MS Visual Studio, o tai labai supaprastina programų kūrimą CUDA.

4.1. Teisingai suskirstykite savo užduotį

Ne visos užduotys tinka SIMD architektūroms. Jei jūsų užduotis tam netinka, gali būti, kad neverta naudoti GPU. Bet jei esate pasiryžę naudoti GPU, turėtumėte pabandyti suskaidyti algoritmą į dalis, kurias galima efektyviai vykdyti SIMD stiliumi. Jei reikia, pakeiskite algoritmą, kad išspręstumėte savo problemą, sugalvokite naują – tokį, kuris puikiai tiktų su SIMD. Tinkamos GPU naudojimo srities pavyzdys yra piramidinis masyvo elementų pridėjimas.

4.2. Atminties tipo pasirinkimas

Įdėkite duomenis į tekstūrą arba nuolatinę atmintį, jei visos užduotys tame pačiame bloke pasiekia tą pačią atminties vietą arba arti esančias sritis. 2D duomenis galima efektyviai apdoroti naudojant text2Dfetch ir text2D funkcijas. Tekstūros atmintis yra specialiai optimizuota 2D mėginių ėmimui.

Naudokite visuotinę atmintį kartu su bendra atmintimi, jei visos užduotys atsitiktinai pasiekia skirtingas, plačiai atskirtas atminties vietas (su labai skirtingais adresais arba koordinatėmis, jei tai yra 2D / 3D duomenys).

pasaulinė atmintis => bendroji atmintis

Syncthreads ();

Apdorokite duomenis bendrojoje atmintyje

Syncthreads ();

pasaulinė atmintis<= разделяемая память

4.3. Įjungti atminties skaitiklius

Kompiliatoriaus vėliavėlė --ptxas-options=-v leidžia tiksliai pasakyti, kiek ir kokios atminties (registrų, bendrinamos, vietinės, pastovios) naudojate. Jei kompiliatorius naudoja vietinę atmintį, jūs tikrai apie tai žinote. Duomenų apie naudojamos atminties kiekį ir tipą analizė gali labai padėti optimizuoti programą.

4.4. Stenkitės kuo labiau sumažinti registrų ir bendrosios atminties naudojimą

Kuo daugiau registrų ar bendrosios atminties naudoja branduolys, tuo mažiau gijų (tiksliau kreipimų) gali būti vykdoma vienu metu daugiaprocesoriuje, nes Kelių procesorių ištekliai riboti. Todėl šiek tiek padidėjus registrų ar bendros atminties užimtumui, kai kuriais atvejais našumas gali sumažėti per pusę – būtent todėl, kad dabar keliuose procesoriuose vienu metu atliekama lygiai pusė tiek deformacijų.

4.5. Bendra atmintis, o ne vietinė

Jei „Nvidia“ kompiliatorius dėl kokių nors priežasčių paskirstė duomenis vietinei atminčiai (paprastai tai pastebima labai smarkiai sumažėjus našumui tose vietose, kur nėra daug išteklių), išsiaiškinkite, kokie duomenys tiksliai atsidūrė vietinėje atmintyje ir įdėkite juos į bendrinamą atmintis).

Dažnai kompiliatorius įdeda kintamąjį į vietinę atmintį, jei jis nenaudojamas dažnai. Pavyzdžiui, tai yra tam tikras kaupiklis, kuriame jūs kaupiate vertę apskaičiuodami ką nors ciklo būdu. Jei ciklas yra didelis pagal kodo tūrį (bet ne pagal vykdymo laiką!), tada kompiliatorius gali įdėti jūsų akumuliatorių į vietinę atmintį, nes ji naudojama palyginti retai, o registrų nedaug. Šiuo atveju našumo praradimas gali būti pastebimas.

Jei tikrai retai naudojate kintamąjį, geriau jį aiškiai įdėti į pasaulinę atmintį.

Nors kompiliatoriui gali atrodyti patogu tokius kintamuosius automatiškai įdėti į vietinę atmintį, iš tikrųjų taip nėra. Vėlesnių programos modifikacijų metu bus sunku rasti kliūtis, jei kintamasis bus pradėtas naudoti dažniau. Kompiliatorius gali perkelti tokį kintamąjį į registro atmintį arba ne. Jei modifikatorius __global__ yra nurodytas aiškiai, programuotojas labiau linkęs į jį atkreipti dėmesį.

4.6. Išvyniojamos kilpos

Kilpų išvyniojimas yra standartinė daugelio sistemų veikimo technika. Jo esmė – kiekvienos iteracijos metu atlikti daugiau veiksmų, taip sumažinant bendrą iteracijų skaičių, taigi ir sąlyginių šakų, kurias turės atlikti procesorius, skaičių.

Štai kaip išvynioti masyvo sumos (pavyzdžiui, sveikojo skaičiaus) suradimo kilpą:

int a[N]; int suma;

for (int i=0;i

Žinoma, kilpas galima išvynioti ir rankiniu būdu (kaip parodyta aukščiau), bet tai neproduktyvus darbas. Daug geriau naudoti C++ šablonus kartu su įterptomis funkcijomis.

šabloną

klasė ArraySumm

Įrenginys__ statinis T exec(const T* arr) ( return arr + ArraySumm (arr+1); )

šabloną

klasė ArraySumm<0,T>

Įrenginys__ statinis T exec(const T* arr) (grįžta 0; )

už (int i=0;i

suma+= ArraySumm<4,int>::exec(a);

Reikėtų pažymėti vieną įdomią nvcc kompiliatoriaus ypatybę. Kompiliatorius pagal numatytuosius nustatymus visada įtrauks __device__ tipo funkcijas (tai nepaisyti yra speciali direktyva __noinline__).

Todėl galite būti tikri, kad toks pavyzdys kaip aukščiau pateiktas išsiskleis į paprastą teiginių seką ir savo efektyvumu jokiu būdu nebus prastesnis už ranka parašytą kodą. Tačiau bendruoju atveju (ne nvcc) negalite tuo būti tikri, nes inline yra tik nuoroda kompiliatoriui, kad jis gali nepaisyti. Todėl negarantuojama, kad jūsų funkcijos bus įtrauktos.

4.7. Duomenų derinimas ir 16 baitų atranka

Sulygiuokite duomenų struktūras ant 16 baitų ribų. Tokiu atveju kompiliatorius galės naudoti jiems skirtas specialias instrukcijas, kurios vienu metu įkelia duomenis 16 baitų.

Jei struktūra yra 8 baitai ar mažiau, galite lygiuoti ją 8 baitais. Tačiau šiuo atveju galite pasirinkti du kintamuosius vienu metu, sujungdami du 8 baitų kintamuosius į struktūrą, naudodami jungtį arba žymeklį. Perdavimas turėtų būti naudojamas atsargiai, nes kompiliatorius gali įdėti duomenis į vietinę atmintį, o ne į registrus.

4.8. Bendri atminties banko konfliktai

Bendra atmintis suskirstyta į 16 (iš viso!) atminties bankų su 4 baitų žingsniu. Vykdant metmenų gijų telkinį daugiaprocesoriuje, jis yra padalintas į dvi dalis (jei metmenų dydis = 32) iš 16 gijų, kurios savo ruožtu pasiekia bendrą atmintį.

Užduotys skirtingose ​​metmenų pusėse neprieštarauja bendrojoje atmintyje. Dėl to, kad vienos pusės metmenų baseino užduotys pasieks tuos pačius atminties bankus, atsiras susidūrimų ir dėl to sumažės našumas. Užduotys, esančios vienoje deformacijos pusėje, tam tikru žingsniu gali pasiekti skirtingas bendrinamos atminties dalis.

Optimalūs žingsniai yra 4, 12, 28, ..., 2^n-4 baitai (8 pav.).



Ryžiai. 8. Optimalūs žingsniai.

Neoptimalūs žingsniai yra 1, 8, 16, 32, ..., 2^n baitų (9 pav.).



Ryžiai. 9. Neoptimalūs žingsniai

4.9. Pagrindinio kompiuterio duomenų judėjimo sumažinimas<=>Įrenginys

Stenkitės kuo mažiau perduoti tarpinius rezultatus į pagrindinį kompiuterį, kad jie būtų apdoroti naudojant CPU. Įdiekite jei ne visą algoritmą, tai bent pagrindinę jo dalį GPU, palikdami tik valdymo užduotis CPU.

5. CPU/GPU nešiojama matematikos biblioteka

Šio straipsnio autorius darbui su paprastais erdviniais objektais parašė nešiojamą biblioteką MGML_MATH, kurios kodas veikia tiek įrenginyje, tiek pagrindiniame kompiuteryje.

MGML_MATH biblioteka gali būti naudojama kaip pagrindas CPU/GPU nešiojamoms (arba hibridinėms) sistemoms skaičiuoti fizines, grafines ar kitas erdvines problemas. Pagrindinis jo privalumas yra tas, kad tą patį kodą galima naudoti ir CPU, ir GPU, o tuo pačiu greitis yra pagrindinis bibliotekos reikalavimas.

6 . Literatūra

Krisas Kasperskis. Programos optimizavimo būdai. Efektyvus atminties naudojimas. - Sankt Peterburgas: BHV-Petersburg, 2003. - 464 p.: iliustr.

CUDA programavimo vadovas 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA programavimo vadovas 1.1. 14-15 psl

CUDA programavimo vadovas 1.1. 48 puslapis

Pagal Darvino evoliucijos teoriją, pirmoji beždžionė (jei
tiksliau – homo antecessor, žmogaus pirmtakas) vėliau virto
mumyse. Daugiatonių kompiuterių centrai su tūkstančiu ar daugiau radijo lempų,
užėmę ištisus kambarius buvo pakeisti pusės kilogramo nešiojamaisiais kompiuteriais, kurie, beje,
našumu nebus prastesnis už pirmąjį. Priešpilio rašomosios mašinėlės tapo
spausdinant bet ką ir ant bet ko (net ant žmogaus kūno)
daugiafunkciniai įrenginiai. Procesorių gigantai staiga nusprendė pasitraukti
grafikos šerdis „akmenyje“. Ir vaizdo plokštės pradėjo ne tik rodyti vaizdą su
priimtina FPS ir grafikos kokybė, bet ir atlikti visokius skaičiavimus. Taip
vis tiek kaip gaminti! Bus aptarta kelių gijų skaičiavimo technologija naudojant GPU.

Kodėl GPU?

Įdomu, kodėl jie nusprendė visą skaičiavimo galią perkelti į grafiką
adapteris? Kaip matote, procesoriai vis dar madingi, ir vargu ar jie atsisakys savo šiltumo
vieta. Tačiau GPU rankovėje yra pora tūzų, juokdarys ir kai kurios rankovės
pakankamai. Šiuolaikinis centrinis procesorius sukurtas maksimaliai pasiekti
našumas apdorojant sveikųjų skaičių ir slankiojo kablelio duomenis
kableliu, ypač nesijaudindami dėl lygiagretaus informacijos apdorojimo. Tuo pačiu
Laikui bėgant, vaizdo plokštės architektūra leidžia greitai ir be problemų „sulyginti“
duomenų apdorojimas. Viena vertus, skaičiuojami daugiakampiai (dėl 3D konvejerio),
kita vertus, pikselių tekstūros apdorojimas. Akivaizdu, kad egzistuoja „harmoninga
apkrovos gedimas kortelės šerdyje. Be to, atminties ir vaizdo procesoriaus našumas
optimalesnis nei „RAM-cache-processor“ derinys. Šiuo metu duomenų vienetas
vaizdo plokštėje pradeda apdoroti vienas GPU srauto procesorius, kitas
vienetas lygiagrečiai įkeliamas į kitą, ir iš esmės tai lengva pasiekti
GPU apkrova, panaši į magistralės pralaidumą,
tačiau kad tai įvyktų, konvejeriai turi būti pakraunami tolygiai, be
bet kokie sąlyginiai perėjimai ir šakos. Centrinis procesorius dėl savo
universalumas reikalauja visos talpyklos apdorojimo poreikiams patenkinti
informacija.

Ekspertai pagalvojo apie GPU darbą lygiagrečiame skaičiavime ir
matematiką ir sugalvojo teoriją, į kurią daugelis mokslinių skaičiavimų daugeliu atžvilgių yra panašūs
3D grafikos apdorojimas. Daugelis ekspertų mano, kad pagrindinis veiksnys
plėtra GPGPU (Bendrosios paskirties skaičiavimas naudojant GPU – universalus
skaičiavimai naudojant vaizdo plokštę) buvo Brook GPU projekto atsiradimas 2003 m.

Projekto kūrėjai iš Stanfordo universiteto turėjo išspręsti sudėtingą
problema: aparatinė ir programinė įranga priversti grafikos adapterį gaminti
įvairūs skaičiavimai. Ir jiems pavyko. Naudojant bendrąją C kalbą,
Amerikiečių mokslininkai privertė GPU veikti kaip procesorius, pritaikytas tam
lygiagretus apdorojimas. Po Brook pasirodė keletas projektų apie VGA skaičiavimus,
pvz., Accelerator biblioteka, Brahma biblioteka, sistema
GPU++ metaprogramavimas ir kt.

CUDA!

Vystymosi perspektyvų nuojauta privertė AMD Ir NVIDIA
prikibkite prie Brook GPU kaip pitbulis. Jei praleisime rinkodaros politiką, tada
Teisingai viską įgyvendinę galite įsitvirtinti ne tik grafikos sektoriuje
rinkoje, bet ir kompiuterijoje (žiūrėkite specialias skaičiavimo korteles ir
serveriai Tesla su šimtais kelių procesorių), išstumiant įprastus procesorius.

Natūralu, kad „FPS lordų“ keliai išsiskyrė suklupę kiekvienas savaip.
kelią, tačiau pagrindinis principas liko nepakitęs – atlikti skaičiavimus
naudojant GPU. O dabar atidžiau pažvelgsime į „žaliąją“ technologiją - CUDA
(Apskaičiuokite vieningą įrenginių architektūrą).

Mūsų „herojės“ darbas yra pateikti API, dvi iš karto.
Pirmasis yra aukšto lygio, CUDA Runtime, kuris atstovauja funkcijas, kurios
yra suskaidomi į paprastesnius lygius ir perduodami žemesnei API – CUDA tvarkyklei. Taigi
kad frazė „aukšto lygio“ yra tam tikras procesas. Visa druska yra
tiksliai tvarkyklėje, o maloniai sukurtos bibliotekos padės ją gauti
kūrėjai NVIDIA: CUBLAS (matematinių skaičiavimų įrankiai) ir
FFT (skaičiavimas naudojant Furjė algoritmą). Na, pereikime prie praktikos
medžiagos dalys.

CUDA terminija

NVIDIA veikia su labai unikaliais CUDA API apibrėžimais. Jie
skiriasi nuo apibrėžimų, naudojamų dirbant su centriniu procesoriumi.

Siūlas– duomenų rinkinys, kurį reikia apdoroti (ne
reikalauja didelių apdorojimo išteklių).

Metmenys– 32 gijų grupė. Duomenys tvarkomi tik
deformuojasi, todėl deformacija yra minimalus duomenų kiekis.

Blokuoti– srautų rinkinys (nuo 64 iki 512) arba rinkinys
metmenys (nuo 2 iki 16).

Tinklelis yra blokų rinkinys. Šis duomenų padalijimas
naudojami tik našumui pagerinti. Taigi, jei numeris
kelių procesorių yra didelis, tada blokai bus vykdomi lygiagrečiai. Jei su
nesiseka su kortele (kūrėjai rekomenduoja naudoti
adapteris ne mažesnis kaip GeForce 8800 GTS 320 MB), tada duomenų blokai bus apdorojami
nuosekliai.

NVIDIA taip pat pristato tokias sąvokas kaip branduolys, šeimininkas
Ir prietaisas.

Mes dirbame!

Norėdami visiškai dirbti su CUDA, jums reikia:

1. Žinoti GPU šešėlių branduolių struktūrą, nuo programavimo esmės
susideda iš tolygaus apkrovos paskirstymo tarp jų.
2. Mokėti programuoti C aplinkoje, atsižvelgiant į kai kuriuos aspektus.

Kūrėjai NVIDIA kelis kartus atskleidė vaizdo plokštės „vidų“.
kitaip nei esame įpratę matyti. Taigi, norom nenorom, teks viską studijuoti
architektūros subtilybes. Pažvelkime į legendinio G80 „akmens“ struktūrą GeForce 8800
GTX.

Shader branduolį sudaro aštuonios TPC (Texture Processor Cluster) grupės
tekstūros procesoriai (taigi GeForce GTX 280- 15 branduolių, 8800 GTS
jų yra šeši 8600 – keturi ir pan.). Jie, savo ruožtu, susideda iš dviejų
srautinio perdavimo kelių procesorių (toliau – SM). SM (visi jie
16) susideda iš priekinės dalies (sprendžia instrukcijų skaitymo ir dekodavimo problemas) ir
back end (galutinė instrukcijų išvestis) konvejeriai, taip pat aštuoni skaliariniai SP (shader
procesorius) ir du SFU (super funkciniai blokai). Už kiekvieną taktą (vnt
laikas) priekinė dalis pasirenka metmenis ir jį apdoroja. Kad visi metmenys tekėtų
(priminsiu, jų yra 32) apdoroti, konvejerio gale reikia 32/8 = 4 ciklai.

Kiekvienas daugiaprocesorius turi vadinamąją bendrąją atmintį.
Jo dydis yra 16 kilobaitų ir suteikia programuotojui visišką laisvę
veiksmai. Išdalinkite kaip norite :). Bendra atmintis užtikrina ryšį tarp gijų
vienas blokas ir nėra skirtas dirbti su pikselių šešėliais.

SM taip pat gali pasiekti GDDR. Tam jiems buvo suteikta po 8 kilobaitus.
talpyklos atmintis, kurioje saugomi visi svarbiausi darbui reikalingi dalykai (pavyzdžiui, kompiuterija
konstantos).

Daugiaprocesorius turi 8192 registrus. Aktyvių blokų skaičius negali būti
daugiau nei aštuoni, o metmenų skaičius ne didesnis kaip 768/32 = 24. Iš to aišku, kad G80
gali apdoroti daugiausiai 32*16*24 = 12288 gijas per laiko vienetą. Jūs negalite padėti
atsižvelkite į šiuos skaičius optimizuodami programą ateityje (vienoje skalėje
– bloko dydis, kita vertus – siūlų skaičius). Parametrų balansas gali atlikti tam tikrą vaidmenį
todėl svarbus vaidmuo ateityje NVIDIA rekomenduoja naudoti blokus
su 128 arba 256 siūlais. 512 gijų blokas yra neefektyvus, nes turi
padidėjęs vėlavimas. Atsižvelgiant į visas GPU vaizdo plokštės struktūros plius subtilybes
geri programavimo įgūdžiai, galite kurti labai produktyviai
lygiagretaus skaičiavimo įrankis. Beje, apie programavimą...

Programavimas

„Kūrybiškumui“ su CUDA jums reikia GeForce vaizdo plokštė ne žemesnė
aštunta serija. SU

oficialioje svetainėje reikia atsisiųsti tris programinės įrangos paketus: tvarkyklę iš
CUDA palaikymas (kiekviena OS turi savo), pats CUDA SDK paketas (antrasis
beta versija) ir papildomos bibliotekos (CUDA įrankių rinkinys). Technologijos palaiko
operacinės sistemos Windows (XP ir Vista), Linux ir Mac OS X. Studijuoti I
pasirinko Vista Ultimate Edition x64 (žiūrėdamas į priekį, pasakysiu, kad sistema elgėsi
Tiesiog puikus). Šių eilučių rašymo metu tai buvo aktualu darbui
ForceWare tvarkyklė 177.35. Naudojamas kaip įrankių rinkinys
Borland C++ 6 Builder programinės įrangos paketas (nors bet kokia aplinka, kuri veikia su
C kalba).

Kalbą mokančiam žmogui bus lengva priprasti prie naujos aplinkos. Viskas, ko reikia, yra
prisiminti pagrindinius parametrus. Raktinis žodis _global_ (įdėtas prieš funkciją)
rodo, kad funkcija priklauso branduoliui. Jai paskambins centrinis
procesoriaus, o visas darbas vyks su GPU. _pasauliniam_ skambučiui reikia daugiau
konkrečią informaciją, būtent akių dydį, bloko dydį ir koks bus branduolys
taikomos. Pavyzdžiui, eilutė _global_ void saxpy_parallel<<>>, kur X –
tinklelio dydis, o Y yra bloko dydis, nurodo šiuos parametrus.

Simbolis _device_ reiškia, kad funkciją iškvies grafinis branduolys, dar žinomas kaip
laikysis visų nurodymų. Ši funkcija yra kelių procesorių atmintyje,
todėl neįmanoma sužinoti jos adreso. Priešdėlis _host_ reiškia, kad skambutis
o apdorojimas vyks tik dalyvaujant CPU. Reikia atsižvelgti į tai, kad _pasaulio_ ir
_įrenginiai_ negali skambinti vienas kitam ir negali prisiskambinti.

Be to, CUDA kalba turi keletą funkcijų, skirtų darbui su vaizdo atmintimi: cudafree
(atlaisvinti atmintį tarp GDDR ir RAM), cudamemcpy ir cudamemcpy2D (kopijavimas
atmintis tarp GDDR ir RAM) ir cudamalloc (atminties paskirstymas).

Visus programų kodus sudaro CUDA API. Pirmiausia jis paimamas
kodas, skirtas išskirtinai centriniam procesoriui ir taikomas
standartinis rinkinys ir kitas grafiniam adapteriui skirtas kodas
perrašyta į tarpinę kalbą PTX (panašiai kaip assembler).
nustatyti galimas klaidas. Po visų šių „šokių“ finalas
komandų vertimas (vertimas) į GPU/CPU suprantamą kalbą.

Studijų rinkinys

Beveik visi programavimo aspektai aprašyti pridedamoje dokumentacijoje
kartu su tvarkykle ir dviem programomis, taip pat kūrėjų svetainėje. Dydis
Straipsnio neužtenka jiems apibūdinti (suinteresuotas skaitytojas turėtų pridėti
šiek tiek pasistengti ir patiems studijuoti medžiagą).

CUDA SDK naršyklė buvo sukurta specialiai pradedantiesiems. Kiekvienas gali
pajuskite lygiagrečiojo skaičiavimo galią iš pirmų rankų (geriausias testas
stabilumas – pavyzdžiai veikia be artefaktų ar gedimų). Paraiška turi
daug orientacinių mini programų (61 „testas“). Kiekvienai patirtis yra
Išsami programos kodo dokumentacija ir PDF failai. Iš karto matyti, kad žmonės
esantys su savo kūryba naršyklėje dirba rimtą darbą.
Taip pat galite palyginti procesoriaus ir vaizdo plokštės greitį apdorojimo metu
duomenis. Pavyzdžiui, daugiamačių masyvų nuskaitymas vaizdo plokšte GeForce 8800
GT Sukuria 512 MB su bloku su 256 gijomis per 0,17109 milisekundės.
Ši technologija neatpažįsta SLI tandemų, todėl jei turite duetą ar trio,
prieš pradėdami dirbti išjunkite „suporavimo“ funkciją, kitaip CUDA matys tik vieną
prietaisas Dviejų branduolių AMD Athlon 64 X2(pagrindinis dažnis 3000 MHz) ta pati patirtis
praeina per 2,761528 milisekundės. Pasirodo, G92 yra daugiau nei 16 kartų
greičiau nei uola AMD! Kaip matote, tai toli gražu nėra ekstremali sistema
kartu su masių nemėgsta operacine sistema rodo gerą
rezultatus.

Be naršyklės, yra nemažai visuomenei naudingų programų. Adobe
savo gaminius pritaikė naujoms technologijoms. Dabar „Photoshop CS4“ yra pilnas
mažiausiai naudoja grafikos adapterių išteklius (reikia atsisiųsti specialų
Prijunkite). Su tokiomis programomis kaip Badaboom media converter ir RapiHD galite
iššifruoti vaizdo įrašą į MPEG-2 formatą. Tinka garso apdorojimui
Tinka nemokama „Accelero“ programa. CUDA API pritaikytos programinės įrangos kiekis,
neabejotinai augs.

Ir šiuo metu...

Tuo tarpu jūs skaitote šią medžiagą, procesoriaus rūpesčių darbininkai
kuria savo technologijas, skirtas integruoti GPU į CPU. Iš išorės AMD Visi
aišku: jie turi didžiulę patirtį, įgytą kartu su ATI.

„Mikroįrenginių“, „Fusion“, kūrimas susideda iš kelių branduolių
kodinis pavadinimas Buldozeris ir vaizdo lustas RV710 (Kong). Jų santykiai bus
atlikta per patobulintą HyperTransport autobusą. Priklausomai nuo
branduolių skaičius ir jų dažninės charakteristikos AMD planuoja sukurti visą kainą
„akmenų“ hierarchija. Taip pat planuojama gaminti procesorius nešiojamiesiems kompiuteriams (Falcon),
ir daugialypės terpės programėlėms („Bobcat“). Be to, tai technologijų taikymas
nešiojamuose įrenginiuose bus pirmasis iššūkis kanadiečiams. Su vystymusi
lygiagretus skaičiavimas, tokių „akmenų“ naudojimas turėtų būti labai populiarus.

Intelšiek tiek atsiliko nuo laiko su savo Larrabee. Produktai AMD,
jei nieko neatsitiks, parduotuvių lentynose jie pasirodys 2009 metų pabaigoje – anksti
2010 m. O priešo sprendimas paaiškės tik beveik dviese
metų.

Larrabee turės daug (skaitykite: šimtus) branduolių. Pradžioje
Taip pat bus produktų, skirtų 8 – 64 branduoliams. Jie labai panašūs į Pentium, bet
gana stipriai perdirbtas. Kiekvienas branduolys turi 256 kilobaitus L2 talpyklos
(jo dydis laikui bėgant didės). Santykiai bus vykdomi per
1024 bitų dvikryptė žiedinė magistralė. „Intel“ teigia, kad jų „vaikas“ bus
puikiai veikia su DirectX ir Open GL API (skirta Apple kūrėjams), todėl ne
nereikia jokios programinės įrangos įsikišimo.

Kodėl aš tau visa tai pasakiau? Akivaizdu, kad Larrabee ir Fusion neišstums
įprasti, stacionarūs procesoriai iš rinkos, kaip ir jie nebus išstumti iš rinkos
vaizdo plokštės. Žaidėjai ir ekstremalaus sporto entuziastai vis tiek išliks didžiausia svajonė
kelių branduolių procesorius ir kelių aukščiausios klasės VGA tandemas. Bet net kas
procesorių įmonės pereina prie lygiagretaus skaičiavimo, pagrįsto principais
panašus į GPGPU, daug ką pasako. Visų pirma, apie ką tokio
tokia technologija kaip CUDA turi teisę egzistuoti ir, greičiausiai, turės
labai populiarus.

Trumpa santrauka

Lygiagretusis skaičiavimas naudojant vaizdo plokštę yra tik geras įrankis
darbščio programuotojo rankose. Vargu ar perdirbėjams, vadovaujamiems Moore'o įstatymo
ateis galas. Įmonės NVIDIA laukia dar ilgas kelias
reklamuoti savo API masėms (tas pats pasakytina ir apie smegenis ATI/AMD).
Kaip bus, parodys ateitis. Taigi CUDA grįš :).

P.S. Rekomenduoju apsilankyti pradedantiesiems programuotojams ir besidomintiems
šios „virtualios įstaigos“:

NVIDIA oficiali svetainė ir svetainė
GPGPU.com. Visi
pateikta informacija yra anglų kalba, bet bent jau ačiū, kad jos nėra
kinų Taigi pirmyn! Tikiuosi, kad autorius jums bent šiek tiek padėjo
įdomios kelionės į CUDA tyrinėjimą!




Pasidalinkite šiuo straipsniu:
Panašūs straipsniai

2015 m. balandžio 17 d
Valdikliai prietaisų skydelyje Kaip uždaryti prietaisų skydelį

2015 m. balandžio 17 d
Tranzistorius KT815: KT ​​815 parametrai, išvestis ir analogai, kokiu režimu jis veikia

2015 m. balandžio 17 d
Programos apsaugos raktas nerastas 7

2015 m. balandžio 17 d
RLS – lankstus ir tikslus duomenų prieigos apribojimų derinimas