Varianta cuda este diferită. Ce este NVIDIA CUDA? Conflicte de bancă de memorie partajată

Să ne întoarcem la 2003, când Intel și AMD erau într-o cursă comună pentru a găsi cel mai puternic procesor. În doar câțiva ani, ca urmare a acestei curse, vitezele de ceas au crescut semnificativ, mai ales după lansarea Intel Pentium 4.

Dar cursa se apropia rapid de limită. După un val de creșteri uriașe ale vitezei de ceas (între 2001 și 2003, viteza de ceas al Pentium 4 s-a dublat de la 1,5 la 3 GHz), utilizatorii au trebuit să se mulțumească cu zecimile de gigaherți pe care producătorii le-au putut stoarce (din 2003 până în 2005). , vitezele de ceas au crescut de la doar 3 la 3,8 GHz).

Chiar și arhitecturile optimizate pentru frecvențe înalte de ceas, precum Prescott, au început să întâmpine dificultăți, și de data aceasta nu doar cele de producție. Producătorii de cipuri pur și simplu au dat peste legile fizicii. Unii analiști au prezis chiar că Legea lui Moore va înceta să se aplice. Dar acest lucru nu s-a întâmplat. Sensul inițial al legii este adesea distorsionat, dar se referă la numărul de tranzistori de pe suprafața miezului de siliciu. Pentru o lungă perioadă de timp, o creștere a numărului de tranzistori într-un CPU a fost însoțită de o creștere corespunzătoare a performanței - ceea ce a dus la o denaturare a sensului. Dar atunci situația s-a complicat. Dezvoltatorii arhitecturii CPU s-au apropiat de legea reducerii creșterii: numărul de tranzistori care trebuiau adăugate pentru creșterea necesară a performanței a devenit din ce în ce mai mare, ducând la o fundătură.

În timp ce producătorii de procesoare își smulgeau părul încercând să găsească o soluție la problemele lor, producătorii de procesoare grafice au continuat să beneficieze remarcabil de beneficiile Legii lui Moore.

De ce nu au ajuns în aceeași fundătură ca și dezvoltatorii arhitecturii CPU? Motivul este foarte simplu: procesoarele centrale sunt concepute pentru a obține performanțe maxime pe un flux de instrucțiuni care procesează diverse date (atât numere întregi, cât și numere în virgulă mobilă), efectuează acces aleatoriu la memorie etc. Până acum, dezvoltatorii încearcă să ofere un paralelism mai mare al instrucțiunilor - adică să execute cât mai multe instrucțiuni în paralel. De exemplu, cu Pentium, a apărut execuția superscalar, când în anumite condiții era posibilă executarea a două instrucțiuni pe ciclu de ceas. Pentium Pro a primit execuția necorespunzătoare a instrucțiunilor, ceea ce a făcut posibilă optimizarea funcționării unităților de calcul. Problema este că există limitări evidente la executarea unui flux secvenţial de instrucţiuni în paralel, astfel încât creşterea orboasă a numărului de unităţi de calcul nu oferă niciun beneficiu, deoarece acestea vor fi încă inactive de cele mai multe ori.

În schimb, operarea GPU este relativ simplă. Constă în luarea unui grup de poligoane pe o parte și generarea unui grup de pixeli pe cealaltă parte. Poligoanele și pixelii sunt independenți unul de celălalt, astfel încât pot fi procesați în paralel. Astfel, într-un GPU este posibil să se aloce o mare parte a cristalului în unități de calcul, care, spre deosebire de CPU, vor fi efectiv utilizate.

Click pe poza pentru marire.

GPU diferă de CPU nu numai în acest fel. Accesul la memorie în GPU este foarte cuplat - dacă se citește un texel, atunci după câteva cicluri de ceas va fi citit texelul vecin; Când un pixel este înregistrat, cel vecin va fi înregistrat după câteva cicluri de ceas. Prin organizarea inteligentă a memoriei, puteți obține performanțe apropiate de debitul teoretic. Aceasta înseamnă că GPU-ul, spre deosebire de CPU, nu necesită un cache uriaș, deoarece rolul său este de a accelera operațiunile de texturare. Tot ceea ce este necesar sunt câțiva kiloocteți care conțin câțiva texeli utilizați în filtrele biliniare și triliniare.

Click pe poza pentru marire.

Trăiască GeForce FX!

Cele două lumi au rămas mult timp separate. Am folosit procesoare (sau chiar mai multe procesoare) pentru sarcini de birou și aplicații de internet, iar GPU-urile erau bune doar pentru a accelera randarea. Dar o caracteristică a schimbat totul: și anume, apariția GPU-urilor programabile. La început, CPU-urile nu aveau de ce să se teamă. Primele așa-numite GPU-uri programabile (NV20 și R200) ​​nu reprezentau cu greu o amenințare. Numărul de instrucțiuni din program a rămas limitat la aproximativ 10 și au lucrat pe tipuri de date foarte exotice, cum ar fi numere cu virgulă fixă ​​de 9 sau 12 biți.

Click pe poza pentru marire.

Dar legea lui Moore și-a arătat din nou partea cea mai bună. Creșterea numărului de tranzistori nu numai că a făcut posibilă creșterea numărului de unități de calcul, dar a îmbunătățit și flexibilitatea acestora. Apariția lui NV30 poate fi considerată un pas semnificativ înainte din mai multe motive. Desigur, jucătorilor nu prea le-au plăcut plăcile NV30, dar noile GPU-uri au început să se bazeze pe două caracteristici care au fost concepute pentru a schimba percepția GPU-urilor ca mai mult decât acceleratoare grafice.

• Suport pentru calcule cu virgulă mobilă cu precizie unică (chiar dacă nu a respectat standardul IEEE754);
• suport pentru mai mult de o mie de instrucțiuni.

Deci avem toate condițiile care pot atrage cercetători de pionierat care doresc întotdeauna să obțină putere de calcul suplimentară.

Ideea de a folosi acceleratoare grafice pentru calcule matematice nu este nouă. Primele încercări au fost făcute încă din anii 90 ai secolului trecut. Desigur, erau foarte primitive - limitate, în cea mai mare parte, la utilizarea unor funcții hardware, cum ar fi rasterizarea și Z-buffer-uri, pentru a accelera sarcini precum găsirea rutei sau inferența. Diagramele Voronoi .

Click pe poza pentru marire.

În 2003, odată cu apariția shader-urilor evoluate, s-a ajuns la o nouă bară - de data aceasta efectuând calcule matriceale. Acesta a fost anul în care o întreagă secțiune a SIGGRAPH („Calcule pe GPU”) a fost dedicată unei noi zone a IT. Această inițiativă timpurie a fost numită GPGPU (General-Purpose calcul on GPU). Iar un punct de cotitură timpuriu a fost apariția lui .

Pentru a înțelege rolul BrookGPU, trebuie să înțelegeți cum s-a întâmplat totul înainte de apariția sa. Singura modalitate de a obține resurse GPU în 2003 a fost să utilizați unul dintre cele două API-uri grafice - Direct3D sau OpenGL. În consecință, dezvoltatorii care doreau capabilități GPU pentru calcularea lor au trebuit să se bazeze pe cele două API-uri menționate. Problema este că nu au fost întotdeauna experți în programarea plăcilor video, iar acest lucru a complicat serios accesul la tehnologie. Dacă programatorii 3D operează cu shadere, texturi și fragmente, atunci specialiștii din domeniul programării paralele se bazează pe thread-uri, nuclee, scatters etc. Prin urmare, mai întâi a fost necesar să se facă analogii între cele două lumi.

• Flux este un flux de elemente de același tip în GPU poate fi reprezentat printr-o textură. În principiu, în programarea clasică există un astfel de analog ca o matrice.
• Nucleu- o funcție care va fi aplicată independent fiecărui element al fluxului; este echivalentul unui pixel shader. În programarea clasică, putem da o analogie a unei bucle - se aplică unui număr mare de elemente.
• Pentru a citi rezultatele aplicării nucleului pe un fir, trebuie creată o textură. Nu există un echivalent pe CPU, deoarece are acces complet la memorie.
• Controlul locației din memorie în care se va face scrierea (în operațiuni de împrăștiere) se face prin vertex shader, deoarece pixel shader-ul nu poate modifica coordonatele pixelului procesat.

După cum puteți vedea, chiar și luând în considerare analogiile de mai sus, sarcina nu pare simplă. Și Brook a venit în ajutor. Acest nume se referă la extensii ale limbajului C („C cu fluxuri”, „C cu fluxuri”), așa cum le-au numit dezvoltatorii de la Stanford. În esență, sarcina lui Brook a fost să ascundă de programator toate componentele API-ului 3D, ceea ce a făcut posibilă prezentarea GPU-ului ca un coprocesor pentru calculul paralel. Pentru a face acest lucru, compilatorul Brook a procesat un fișier .br cu cod C++ și extensii, apoi a generat cod C++ care a fost legat la o bibliotecă cu suport pentru diferite ieșiri (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Click pe poza pentru marire.

Brook are mai multe credite la credit, primul dintre care scoate GPGPU din umbră, astfel încât tehnologia să poată fi îmbrățișată de către mase. Deși după anunțarea proiectului, o serie de site-uri IT au raportat prea optimist că lansarea lui Brook pune la îndoială existența CPU-urilor, care vor fi înlocuite în curând cu GPU-uri mai puternice. Dar, după cum vedem, nici după cinci ani acest lucru nu s-a întâmplat. Sincer, nu credem că asta se va întâmpla vreodată. Pe de altă parte, privind evoluția de succes a procesoarelor, care sunt din ce în ce mai orientate spre paralelism (mai multe nuclee, tehnologie SMT multi-threading, extinderea blocurilor SIMD), precum și a GPU-urilor, care, dimpotrivă, devin din ce în ce mai multe și mai universal (suport pentru calcule cu virgulă mobilă) precizie unică, calcule întregi, suport pentru calcule cu precizie dublă), se pare că GPU și CPU se vor îmbina în curând pur și simplu. Ce se va întâmpla atunci? Vor fi absorbite GPU-urile de procesoare, așa cum sa întâmplat cu coprocesoarele matematice? Foarte posibil. Intel și AMD lucrează astăzi la proiecte similare. Dar încă se pot schimba multe.

Dar să revenim la subiectul nostru. Avantajul lui Brook a fost acela de a populariza conceptul de GPGPU, a simplificat semnificativ accesul la resursele GPU, ceea ce a permis din ce în ce mai multor utilizatori să stăpânească noul model de programare. Pe de altă parte, în ciuda tuturor calităților lui Brook, mai era un drum lung de parcurs înainte ca resursele GPU să poată fi folosite pentru calcul.

Una dintre probleme este legată de diferitele niveluri de abstractizare și, de asemenea, în special, de încărcarea suplimentară excesivă creată de API-ul 3D, care poate fi destul de vizibilă. Dar o problemă mai serioasă poate fi considerată o problemă de compatibilitate, cu care dezvoltatorii Brook nu au putut face nimic. Există o concurență acerbă între producătorii de GPU, așa că adesea își optimizează driverele. În timp ce astfel de optimizări sunt în mare parte bune pentru jucători, ele ar putea pune capăt compatibilității Brook peste noapte. Prin urmare, este greu de imaginat utilizarea acestui API în codul de producție care va funcționa undeva. Și pentru o lungă perioadă de timp, Brook a rămas în rezerva cercetătorilor și programatorilor amatori.

Succesul lui Brook a fost însă suficient pentru a atrage atenția ATI și nVidia, care au devenit interesați de o astfel de inițiativă, deoarece ar putea extinde piața, deschizând un nou sector important pentru companii.

Cercetătorii implicați inițial în proiectul Brook s-au alăturat rapid echipelor de dezvoltare din Santa Clara pentru a prezenta o strategie globală de dezvoltare a noii piețe. Ideea a fost de a crea o combinație de hardware și software potrivită pentru sarcinile GPGPU. Deoarece dezvoltatorii nVidia cunosc toate secretele GPU-urilor lor, ei nu se puteau baza pe API-ul grafic, ci comunica cu GPU-ul printr-un driver. Deși, desigur, acest lucru vine cu propriile sale probleme. Deci, echipa de dezvoltare CUDA (Compute Unified Device Architecture) a creat un set de straturi software pentru lucrul cu GPU-uri.

Click pe poza pentru marire.

După cum puteți vedea în diagramă, CUDA oferă două API-uri.

• API de nivel înalt: CUDA Runtime API;
• API de nivel scăzut: CUDA Driver API.

Deoarece API-ul de nivel înalt este implementat peste API-ul de nivel scăzut, fiecare apel la o funcție la nivel de execuție este defalcat în instrucțiuni mai simple pe care le procesează API-ul Driver. Vă rugăm să rețineți că cele două API-uri se exclud reciproc: un programator poate folosi una sau alta API, dar nu va fi posibilă combinarea apelurilor de funcții de la cele două API-uri. În general, termenul „API de nivel înalt” este relativ. Chiar și API-ul Runtime este de așa natură încât mulți l-ar considera de nivel scăzut; cu toate acestea, încă oferă funcții care sunt foarte convenabile pentru inițializarea sau gestionarea contextului. Dar nu vă așteptați la un nivel deosebit de ridicat de abstractizare - trebuie totuși să aveți un set bun de cunoștințe despre GPU-urile nVidia și cum funcționează acestea.

API-ul Driver este și mai dificil de utilizat; va fi nevoie de mai mult efort pentru a rula procesarea GPU. Pe de altă parte, un API de nivel scăzut este mai flexibil, oferind programatorului control suplimentar dacă este necesar. Două API-uri sunt capabile să lucreze cu resurse OpenGL sau Direct3D (numai a noua versiune de astăzi). Beneficiul acestei caracteristici este evident - CUDA poate fi folosit pentru a crea resurse (geometrie, texturi procedurale etc.) care pot fi transmise API-ului grafic sau, dimpotrivă, puteți avea API-ul 3D să trimită rezultate de randare către programul CUDA, care, la rândul său, va efectua post-procesare. Există multe exemple de astfel de interacțiuni, iar avantajul este că resursele continuă să fie stocate în memoria GPU-ului, nu trebuie să fie transferate prin magistrala PCI Express, care rămâne totuși un blocaj.

Cu toate acestea, trebuie remarcat faptul că partajarea resurselor în memoria video nu este întotdeauna ideală și poate duce la unele bătăi de cap. De exemplu, atunci când se schimbă rezoluția sau adâncimea culorii, datele grafice au prioritate. Prin urmare, dacă trebuie să creșteți resursele din memoria tampon, driverul va face acest lucru cu ușurință în detrimentul resurselor aplicațiilor CUDA, care pur și simplu se vor bloca cu o eroare. Desigur, nu foarte elegant, dar o astfel de situație nu ar trebui să se întâmple foarte des. Și din moment ce am început să vorbim despre dezavantaje: dacă doriți să utilizați mai multe GPU-uri pentru aplicațiile CUDA, atunci trebuie să dezactivați mai întâi modul SLI, altfel aplicațiile CUDA vor putea „vedea” doar un singur GPU.

În cele din urmă, al treilea nivel de software este dedicat bibliotecilor - doi, mai exact.

• CUBLAS, care conține blocurile necesare pentru calculele de algebră liniară pe GPU;
• CUFFT, care suportă calculul transformărilor Fourier - un algoritm utilizat pe scară largă în domeniul procesării semnalului.

Înainte de a ne aprofunda în CUDA, să definim o serie de termeni împrăștiați în documentația nVidia. Compania a ales o terminologie foarte specifică cu care este greu de obișnuit. În primul rând, observăm că firîn CUDA nu are același înțeles ca un fir CPU și, de asemenea, nu este echivalentul unui fir din articolele noastre GPU. Firul GPU în acest caz este setul de bază de date care trebuie procesat. Spre deosebire de firele de execuție CPU, firele de execuție CUDA sunt foarte „ușoare”, adică schimbarea contextului între două fire de execuție nu este o operațiune care necesită mult resurse.

Al doilea termen des întâlnit în documentația CUDA este urzeală. Nu există confuzie aici, deoarece nu există un analog în rusă (cu excepția cazului în care sunteți fan al Start Trek sau al jocului Warhammer). De fapt, termenul provine din industria textilă, unde firul de bătătură este tras printr-un fir de urzeală care este întins pe un războaie. Un warp în CUDA este un grup de 32 de fire și reprezintă cantitatea minimă de date procesate într-un mod SIMD în multiprocesoarele CUDA.

Dar o astfel de „granulare” nu este întotdeauna convenabilă pentru programator. Prin urmare, în CUDA, în loc să lucrezi direct cu warps, poți lucra cu blocuri, care conține de la 64 la 512 fire.

În cele din urmă, aceste blocuri se reunesc în grilă. Avantajul acestei grupări este că numărul de blocuri procesate simultan de GPU este strâns legat de resursele hardware, așa cum vom vedea mai jos. Gruparea blocurilor în grile vă permite să abstrageți complet această limitare și să aplicați nucleul la mai multe fire într-un singur apel, fără a vă face griji cu privire la resursele fixe. Bibliotecile CUDA sunt responsabile de toate acestea. În plus, un astfel de model se scalează bine. Dacă GPU-ul are puține resurse, va executa blocuri secvenţial. Dacă numărul de procesoare de calcul este mare, atunci blocurile pot fi executate în paralel. Adică, același cod poate rula pe GPU-uri entry-level, precum și pe modele de top și chiar viitoare.

Mai există câțiva termeni în API-ul CUDA care indică CPU ( gazdă/gazdă) și GPU ( dispozitiv). Dacă această mică introducere nu te-a speriat, atunci este timpul să arunci o privire mai atentă la CUDA.

Dacă citiți în mod regulat Ghidul hardware al lui Tom, atunci arhitectura celor mai recente GPU-uri de la nVidia vă este familiară. Dacă nu, vă recomandăm să citiți articolul „. nVidia GeForce GTX 260 și 280: nouă generație de plăci video„Când vine vorba de CUDA, Nvidia prezintă arhitectura puțin diferit, dezvăluind câteva detalii care erau ascunse anterior.

După cum puteți vedea din ilustrația de mai sus, nucleul shader nVidia este format din mai multe clustere de procesoare de textură (Cluster de procesoare de texturi, TPC). Placa video 8800 GTX, de exemplu, folosea opt clustere, 8800 GTS - șase etc. Fiecare cluster constă în esență dintr-un bloc de textură și două multiprocesor de streaming. Acestea din urmă includ începutul conductei (front end), care citește și decodifică instrucțiunile, precum și trimiterea acestora pentru execuție, și sfârșitul conductei (back end), constând din opt dispozitive de calcul și două dispozitive superfuncționale. SFU (Super Function Unit), unde instrucțiunile sunt executate folosind principiul SIMD, adică o singură instrucțiune este aplicată tuturor firelor din warp. nVidia numește această metodă de execuție SIMT(o singură instrucțiune, mai multe fire, o singură instrucțiune, mai multe fire). Este important de menționat că capătul transportorului funcționează cu o frecvență de două ori mai mare decât începutul său. În practică, aceasta înseamnă că piesa apare de două ori mai „largă” decât este de fapt (adică ca un bloc SIMD cu 16 canale în loc de unul cu opt canale). Multiprocesoarele de streaming funcționează după cum urmează: la fiecare ciclu de ceas, începutul conductei selectează un warp gata de execuție și începe executarea instrucțiunii. Pentru ca instrucțiunea să se aplice tuturor celor 32 de fire din warp, sfârșitul conductei va necesita patru cicluri de ceas, dar, deoarece rulează la o frecvență de două ori mai mare decât cea de la început, va necesita doar două cicluri de ceas (în ceea ce privește începutul conducta). Prin urmare, pentru ca începutul conductei să nu oprească un ciclu de ceas, iar hardware-ul să fie încărcat maxim, într-un caz ideal, puteți alterna instrucțiuni în fiecare ciclu de ceas - o instrucțiune clasică într-un ciclu de ceas și o instrucțiune pentru SFU în altul. .

Fiecare multiprocesor are un set specific de resurse care merită înțeles. Există o zonă mică de memorie numită „Memorie partajată”, 16 kbytes per multiprocesor. Aceasta nu este deloc o memorie cache: programatorul o poate folosi la propria discreție. Adică avem ceva aproape de Magazinul Local de SPU-uri pe procesoarele Cell. Acest detaliu este foarte interesant, deoarece subliniază că CUDA este o combinație de tehnologii software și hardware. Această zonă de memorie nu este folosită pentru pixel shaders, ceea ce Nvidia subliniază inteligent „nu ne place pixelii care vorbesc între ei”.

Această zonă de memorie deschide posibilitatea schimbului de informații între fire într-un singur bloc. Este important să subliniem această limitare: toate firele dintr-un bloc sunt garantate a fi executate de un singur multiprocesor. În schimb, alocarea blocurilor către diferite multiprocesoare nu este specificată deloc, iar două fire de execuție din blocuri diferite nu pot face schimb de informații între ele în timpul execuției. Adică, folosirea memoriei partajate nu este atât de ușoară. Cu toate acestea, memoria partajată este în continuare justificată, cu excepția cazurilor în care mai multe fire de execuție încearcă să acceseze aceeași bancă de memorie, provocând un conflict. În alte situații, accesul la memoria partajată este la fel de rapid ca și accesul la registre.

Memoria partajată nu este singura memorie pe care o pot accesa multiprocesoarele. Pot folosi memoria video, dar cu o lățime de bandă mai mică și o latență mai mare. Prin urmare, pentru a reduce frecvența de acces la această memorie, nVidia a echipat multiprocesoarele cu un cache (aproximativ 8 KB per multiprocesor) care stochează constante și texturi.

Multiprocesorul are 8.192 de registre care sunt comune tuturor firelor de execuție ale tuturor blocurilor active pe multiprocesor. Numărul de blocuri active per multiprocesor nu poate depăși opt, iar numărul warps-urilor active este limitat la 24 (768 fire). Prin urmare, 8800 GTX poate gestiona până la 12.288 de fire simultan. Toate aceste constrângeri merită menționate deoarece permit optimizarea algoritmului pe baza resurselor disponibile.

Optimizarea unui program CUDA constă astfel în obținerea unui echilibru optim între numărul de blocuri și dimensiunea acestora. Mai multe fire de execuție per bloc vor fi utile pentru reducerea latenței memoriei, dar numărul de registre disponibile pe fir este, de asemenea, redus. Mai mult, un bloc de 512 fire de execuție va fi ineficient deoarece doar un bloc poate fi activ pe un multiprocesor, rezultând o pierdere de 256 fire de execuție. Prin urmare, nVidia recomandă utilizarea blocurilor de 128 sau 256 de fire, ceea ce oferă compromisul optim între latența redusă și numărul de registre pentru majoritatea nucleelor/kernel-urilor.

Din punct de vedere software, CUDA constă dintr-un set de extensii la limbajul C, care amintește de BrookGPU, precum și mai multe apeluri API specifice. Printre extensii se numără specificatorii de tip legate de funcții și variabile. Este important să rețineți cuvântul cheie __global__, care, atunci când este dat înaintea funcției, arată că aceasta din urmă aparține nucleului - această funcție va fi apelată de CPU și va fi executată pe GPU. Prefix __dispozitiv__ indică faptul că funcția va fi executată pe GPU (care, apropo, este ceea ce CUDA numește „dispozitiv”), dar poate fi apelată doar din GPU (cu alte cuvinte, dintr-o altă funcție __device__ sau dintr-o funcție __global__) . În sfârșit, prefixul __gazdă__ opțional, denotă o funcție care este apelată de CPU și executată de CPU - cu alte cuvinte, o funcție obișnuită.

Există o serie de restricții asociate cu funcțiile __device__ și __global__: ele nu pot fi recursive (adică, nu se pot numi) și nu pot avea un număr variabil de argumente. În cele din urmă, deoarece funcțiile __device__ sunt localizate în spațiul de memorie GPU, este logic ca adresa lor să nu poată fi obținută. Variabilele au, de asemenea, o serie de calificative care indică zona de memorie în care vor fi stocate. Variabilă cu prefix __distribuit__înseamnă că va fi stocat în memoria partajată a multiprocesorului de streaming. Apelul funcției __global__ este ușor diferit. Chestia este că, atunci când apelați, trebuie să specificați configurația de execuție - mai precis, dimensiunea grilei la care va fi aplicat nucleul, precum și dimensiunea fiecărui bloc. Să luăm, de exemplu, un nucleu cu următoarea semnătură.

__global__ void Func(parametru float*);

Se va numi ca

Func<<< Dg, Db >>>(parametru);

unde Dg este dimensiunea grilei și Db este dimensiunea blocului. Aceste două variabile se referă la un nou tip de vector introdus cu CUDA.

API-ul CUDA conține funcții pentru lucrul cu memoria în VRAM: cudaMalloc pentru alocarea memoriei, cudaFree pentru eliberarea acesteia și cudaMemcpy pentru copierea memoriei între RAM și VRAM și invers.

Vom încheia această recenzie cu un mod foarte interesant în care este compilat un program CUDA: compilarea se realizează în mai multe etape. În primul rând, codul specific CPU este extras și transmis compilatorului standard. Codul destinat GPU-ului este mai întâi convertit în limbajul intermediar PTX. Este similar cu un asamblator și vă permite să examinați codul căutând potențiale ineficiențe. În cele din urmă, ultima fază este de a traduce limbajul intermediar în instrucțiuni specifice GPU și de a crea un fișier binar.

După ce mă uit la documentația nVidia, sunt tentat să încerc săptămâna aceasta CUDA. Într-adevăr, ce ar putea fi mai bun decât evaluarea unui API prin crearea propriului program? Acesta este momentul în care majoritatea problemelor ar trebui să iasă la suprafață, chiar dacă totul arată perfect pe hârtie. De asemenea, practica va arăta cel mai bine cât de bine ați înțeles toate principiile subliniate în documentația CUDA.

Este destul de ușor să te implici într-un astfel de proiect. Astăzi există un număr mare de instrumente gratuite, dar de înaltă calitate disponibile pentru descărcare. Pentru testul nostru am folosit Visual C++ Express 2005, care are tot ce ne trebuie. Cea mai grea parte a fost să găsești un program care să nu dureze săptămâni pentru a se porta pe GPU și, totuși, să fie suficient de interesant încât eforturile noastre să nu fie în zadar. În final, am ales o bucată de cod care ia o hartă de înălțime și calculează o hartă normală corespunzătoare. Nu vom intra în detaliu despre această funcție, deoarece nu este deloc interesantă în acest articol. Pe scurt, programul se ocupă de curbura zonelor: pentru fiecare pixel al imaginii inițiale, suprapunem o matrice care determină culoarea pixelului rezultat din imaginea generată din pixelii adiacenți, folosind o formulă mai mult sau mai puțin complexă. Avantajul acestei funcții este că este foarte ușor de paralelizat, astfel încât acest test arată perfect capacitățile CUDA.

Un alt avantaj este că avem deja o implementare CPU, așa că putem compara rezultatul cu versiunea CUDA - fără a reinventa roata.

Să repetăm ​​încă o dată că scopul testului a fost de a face cunoștință cu utilitățile CUDA SDK și nu de a compara versiunile pentru CPU și GPU. Deoarece aceasta a fost prima noastră încercare de a crea un program CUDA, nu ne așteptam cu adevărat la performanță ridicată. Deoarece această parte a codului nu este critică, versiunea pentru CPU nu a fost optimizată, așa că o comparație directă a rezultatelor nu este deloc interesantă.

Performanţă

Cu toate acestea, am măsurat timpul de execuție pentru a vedea dacă există un avantaj în utilizarea CUDA chiar și cu cea mai brută implementare sau dacă am avea nevoie de practică lungă și obositoare pentru a obține vreun beneficiu atunci când folosim GPU. Mașina de testare a fost luată din laboratorul nostru de dezvoltare - un laptop cu un procesor Core 2 Duo T5450 și o placă grafică GeForce 8600M GT care rulează Vista. Acesta este departe de a fi un supercomputer, dar rezultatele sunt foarte interesante, deoarece testul nu este „personalizat” pentru GPU. Este întotdeauna plăcut să vezi că Nvidia arată câștiguri uriașe la sistemele cu GPU-uri monstru și lățime de bandă multă, dar, în practică, multe dintre cele 70 de milioane de GPU-uri capabile CUDA de pe piața PC-urilor de astăzi nu sunt la fel de puternice, motiv pentru care testul nostru încă rămâne valabil. .

Pentru o imagine de 2048 x 2048 pixeli, am obținut următoarele rezultate.

• Fire CPU 1: 1.419 ms;
• CPU 2 fire: 749 ms;
• CPU 4 fire: 593 ms
• GPU (8600M GT) blocuri de 256 fire: 109 ms;
• GPU (8600M GT) blocuri de 128 fire: 94 ms;
• Blocuri GPU (8800 GTX) de 128 fire/256 fire: 31 ms.

Din rezultate se pot trage mai multe concluzii. Să începem cu faptul că, în ciuda discuțiilor despre lenea evidentă a programatorilor, am modificat versiunea inițială a procesorului pentru mai multe fire. După cum am menționat deja, codul este ideal pentru această situație - tot ceea ce este necesar este să împărțiți imaginea inițială în atâtea zone câte fire există. Vă rugăm să rețineți că trecerea de la un fir la două pe procesorul nostru dual-core a dus la o accelerare aproape liniară, ceea ce indică și natura paralelă a programului de testare. În mod destul de neașteptat, versiunea cu patru fire s-a dovedit a fi, de asemenea, mai rapidă, deși pe procesorul nostru acest lucru este foarte ciudat - dimpotrivă, ne-am putea aștepta la o scădere a eficienței din cauza suprasolicitarii gestionării firelor suplimentare. Cum poate fi explicat acest rezultat? Este greu de spus, dar programatorul de fire Windows poate fi de vină; în orice caz, repetăm ​​rezultatul. Cu texturi mai mici (512x512), câștigul din threading nu a fost la fel de pronunțat (aproximativ 35% față de 100%), iar comportamentul versiunii cu patru fire a fost mai logic, fără nicio creștere față de versiunea cu două fire. GPU-ul era încă mai rapid, dar nu la fel de mult mai rapid (8600M GT a fost de trei ori mai rapid decât versiunea dual-thread).

Click pe poza pentru marire.

A doua observație semnificativă este că chiar și cea mai lentă implementare GPU a fost de aproape șase ori mai rapidă decât versiunea CPU cu cea mai mare performanță. Pentru primul program și versiunea neoptimizată a algoritmului, rezultatul este foarte încurajator. Vă rugăm să rețineți că am obținut rezultate semnificativ mai bune pe blocuri mici, deși intuiția poate sugera altfel. Explicația este simplă - programul nostru folosește 14 registre pe fir, iar cu blocuri cu 256 de fire, sunt necesare 3.584 de registre pe bloc și sunt necesare 768 de fire pentru a încărca complet procesorul, așa cum am arătat. În cazul nostru, aceasta înseamnă trei blocuri sau 10.572 de registre. Dar multiprocesorul are doar 8.192 de registre, deci nu poate menține decât două blocuri active. În schimb, cu blocuri de 128 fire, avem nevoie de 1.792 de registre pe bloc; Dacă 8.192 este împărțit la 1.792 și rotunjit la cel mai apropiat număr întreg, obținem patru blocuri. În practică, numărul de fire va fi același (512 per multiprocesor, deși teoretic sunt necesare 768 pentru o încărcare completă), dar creșterea numărului de blocuri oferă GPU-ului avantajul flexibilității în accesarea memoriei - atunci când există o operație cu întârzieri mari, puteți începe executarea instrucțiunilor dintr-un alt bloc în timp ce așteptați primirea rezultatelor. Patru blocuri reduc în mod clar latența, mai ales că programul nostru utilizează mai multe accesări la memorie.

Analiză

În cele din urmă, în ciuda celor spuse mai sus, nu am rezistat tentației și am rulat programul pe 8800 GTX, care a fost de trei ori mai rapid decât 8600, indiferent de dimensiunea blocului. Ai putea crede că în practică pe arhitecturile corespunzătoare rezultatul ar fi de patru sau mai multe ori mai mare: 128 de procesoare ALU/shader față de 32 și o viteză de ceas mai mare (1,35 GHz față de 950 MHz), dar nu a funcționat așa. Cel mai probabil, accesul la memorie a fost factorul limitativ. Pentru a fi mai precis, imaginea inițială este accesată ca o matrice CUDA multidimensională - un termen destul de complex pentru ceea ce nu este altceva decât o textură. Dar există mai multe avantaje.

• accesele beneficiază de cache-ul texturii;
• folosim modul de împachetare, în care nu este nevoie să procesăm limitele imaginii, spre deosebire de versiunea CPU.

În plus, putem beneficia de filtrare „gratuită” cu adresare normalizată între în schimb și , dar în cazul nostru este puțin probabil ca acest lucru să fie util. După cum știți, 8600 are 16 unități de textură, față de 32 pentru 8800 GTX. Prin urmare, raportul dintre cele două arhitecturi este de doar doi la unu. Adăugați la aceasta diferența de frecvențe și obținem un raport de (32 x 0,575) / (16 x 0,475) = 2,4 - aproape de raportul „trei la unu” pe care l-am obținut de fapt. Această teorie explică, de asemenea, de ce dimensiunea blocului nu se schimbă prea mult la G80, deoarece ALU se bazează încă pe blocurile de textură.

Click pe poza pentru marire.

Pe lângă rezultatele promițătoare, prima noastră cunoștință cu CUDA a mers foarte bine, având în vedere condițiile nu cele mai favorabile alese. Dezvoltarea pe un laptop Vista înseamnă că va trebui să utilizați CUDA SDK 2.0, încă în beta, cu driverul 174.55, care este și beta. În ciuda acestui fapt, nu putem raporta surprize neplăcute - doar erori inițiale în timpul primei depanări, când programul nostru încă foarte defectuos a încercat să abordeze memoria în afara spațiului alocat.

Monitorul a început să pâlpâie sălbatic, apoi ecranul s-a înnegrit... până când Vista a rulat serviciul de reparații șoferi și totul a fost bine. Dar este totuși oarecum surprinzător să observați acest lucru dacă sunteți obișnuit să vedeți o eroare tipică de eroare de segmentare în programele standard precum al nostru. În sfârșit, o mică critică la adresa nVidia: în toată documentația disponibilă pentru CUDA, nu există un ghid mic care să vă spună pas cu pas cum să configurați un mediu de dezvoltare pentru Visual Studio. De fapt, aceasta nu este o problemă mare, deoarece SDK-ul are un set complet de exemple pe care le puteți studia pentru a înțelege cadrul pentru aplicațiile CUDA, dar un ghid pentru începători ar fi util.

Click pe poza pentru marire.

Nvidia a introdus CUDA odată cu lansarea GeForce 8800. Și la momentul respectiv promisiunile păreau foarte tentante, dar ne-am ținut entuziasmul până când l-am testat efectiv. Într-adevăr, la acea vreme părea mai degrabă a marca teritoriul pentru a rămâne pe valul GPGPU. Fără un SDK disponibil, este greu de spus că aceasta nu este doar o altă greșeală de marketing care nu va funcționa. Nu este prima dată când o inițiativă bună a fost anunțată prea devreme și nu a reușit să vadă lumina zilei la momentul respectiv din cauza lipsei de sprijin – mai ales într-un sector atât de competitiv. Acum, la un an și jumătate de la anunț, putem spune cu încredere că nVidia s-a ținut de cuvânt.

SDK-ul a apărut rapid în beta la începutul anului 2007, iar de atunci a fost rapid actualizat, ceea ce demonstrează importanța acestui proiect pentru nVidia. Astăzi, CUDA se dezvoltă foarte frumos: SDK-ul este deja disponibil în versiunea beta 2.0 pentru sistemele de operare majore (Windows XP și Vista, Linux, precum și 1.1 pentru Mac OS X), iar nVidia a dedicat o întreagă secțiune a site-ului pentru dezvoltatori.

La un nivel mai profesionist, impresia de la primii pasi cu CUDA s-a dovedit a fi foarte pozitiva. Chiar dacă sunteți familiarizat cu arhitectura GPU, vă puteți da seama cu ușurință. Când un API pare clar la prima vedere, începi imediat să crezi că vei obține rezultate convingătoare. Dar nu se va pierde timpul de calcul din cauza numeroaselor transferuri de la CPU la GPU? Și cum să folosiți aceste mii de fire fără primitivă de sincronizare? Am început experimentele noastre având în vedere toate aceste preocupări. Dar s-au disipat rapid când prima versiune a algoritmului nostru, deși foarte banală, s-a dovedit a fi semnificativ mai rapidă decât pe procesor.

Deci CUDA nu este o baghetă magică pentru cercetătorii care doresc să convingă conducerea universității să le cumpere o GeForce. CUDA este deja o tehnologie complet accesibilă care poate fi folosită de orice programator cu cunoștințe de C, atâta timp cât sunt dispuși să depună timp și efort pentru a se obișnui cu noua paradigmă de programare. Acest efort nu va fi irosit dacă algoritmii dvs. sunt bine paralelizați. De asemenea, am dori să mulțumim nVidia pentru că a furnizat documentație completă și de înaltă calitate, unde programatorii începători CUDA vor găsi răspunsuri.

De ce are nevoie CUDA pentru a deveni un API recunoscut? Într-un cuvânt: portabilitate. Știm că viitorul IT constă în calculul paralel - astăzi toată lumea se pregătește deja pentru astfel de schimbări și toate inițiativele, atât software, cât și hardware, sunt îndreptate în această direcție. Totuși, în momentul de față, dacă ne uităm la dezvoltarea paradigmelor, suntem încă la stadiul inițial: creăm fire manual și încercăm să planificăm accesul la resursele partajate; Toate acestea pot fi rezolvate cumva dacă numărul de nuclee poate fi numărat pe degetele unei mâini. Dar peste câțiva ani, când numărul procesoarelor se va număra la sute, această posibilitate nu va mai exista. Odată cu lansarea CUDA, nVidia a făcut primul pas în rezolvarea acestei probleme - dar, desigur, această soluție este potrivită numai pentru GPU-urile acestei companii și chiar și atunci nu pentru toată lumea. Doar GF8 și 9 (și derivatele lor Quadro/Tesla) pot rula programe CUDA astăzi. Și noua linie 260/280, desigur.

Click pe poza pentru marire.

Nvidia se poate lăuda că a vândut 70 de milioane de GPU-uri compatibile cu CUDA în întreaga lume, dar asta încă nu este suficient pentru a deveni standardul de facto. Ținând cont de faptul că concurenții nu stau cu mâinile în brațe. AMD oferă propriul SDK (Stream Computing), iar Intel a anunțat o soluție (Ct), deși nu este încă disponibilă. Urmează un război al standardelor și în mod clar nu va mai fi loc pe piață pentru trei concurenți până când un alt jucător precum Microsoft va veni cu un API comun, care cu siguranță va ușura viața dezvoltatorilor.

Prin urmare, nVidia are multe dificultăți în a obține aprobarea CUDA. Deși din punct de vedere tehnologic ne aflăm, fără îndoială, cu o soluție de succes, rămâne totuși să convingem dezvoltatorii de perspectivele acesteia - și acest lucru nu va fi ușor. Cu toate acestea, judecând după multe anunțuri și știri recente referitoare la API, viitorul nu pare sumbru.

Și este conceput pentru a traduce codul gazdă (principal, codul de control) și codul dispozitivului (codul hardware) (fișiere cu extensia .cu) în fișiere obiect potrivite pentru procesul de asamblare a programului sau bibliotecii finale în orice mediu de programare, de exemplu în NetBeans.

Arhitectura CUDA utilizează un model de memorie grilă, modelarea firelor de cluster și instrucțiuni SIMD. Aplicabil nu numai pentru calculul grafic de înaltă performanță, ci și pentru diferite calcule științifice folosind plăcile video nVidia. Oamenii de știință și cercetătorii folosesc pe scară largă CUDA într-o varietate de domenii, inclusiv astrofizică, biologie și chimie computațională, modelare dinamică a fluidelor, interacțiuni electromagnetice, tomografie computerizată, analiză seismică și multe altele. CUDA are capacitatea de a se conecta la aplicații folosind OpenGL și Direct3D. CUDA este un software multiplatform pentru sisteme de operare precum Linux, Mac OS X și Windows.

Pe 22 martie 2010, nVidia a lansat CUDA Toolkit 3.0, care conținea suport pentru OpenCL.

Echipamente

Platforma CUDA a apărut pentru prima dată pe piață odată cu lansarea cipului NVIDIA G80 de generația a opta și a devenit prezentă în toate seriile ulterioare de cipuri grafice care sunt utilizate în familiile de acceleratoare GeForce, Quadro și NVidia Tesla.

Prima serie de hardware care suportă SDK-ul CUDA, G8x, avea un procesor vectorial cu precizie simplă pe 32 de biți care folosea SDK-ul CUDA ca API (CUDA acceptă tipul dublu C, dar precizia sa a fost acum redusă la 32 de biți virgulă mobilă). Procesoarele GT200 ulterioare au suport pentru precizie de 64 de biți (numai SFU), dar performanța este semnificativ mai slabă decât pentru precizia de 32 de biți (datorită faptului că există doar două SFU-uri pe multiprocesor de flux, în timp ce există opt procesoare scalare). GPU-ul organizează multithreading hardware, ceea ce vă permite să utilizați toate resursele GPU-ului. Astfel, se deschide perspectiva de a transfera funcțiile acceleratorului fizic către acceleratorul grafic (un exemplu de implementare este nVidia PhysX). De asemenea, deschide posibilități largi de utilizare a hardware-ului de grafică pe computer pentru a efectua calcule complexe non-grafice: de exemplu, în biologia computațională și în alte ramuri ale științei.

Avantaje

În comparație cu abordarea tradițională de organizare a calculului de uz general prin API-uri grafice, arhitectura CUDA are următoarele avantaje în acest domeniu:

Restricții

• Toate funcțiile executabile pe dispozitiv nu acceptă recursiunea (CUDA Toolkit 3.1 acceptă pointeri și recursivitate) și au alte limitări

GPU-uri și acceleratoare grafice acceptate

Lista dispozitivelor de la producătorul de echipamente Nvidia cu suport deplin declarat pentru tehnologia CUDA este furnizată pe site-ul oficial Nvidia: Produse GPU CUDA-Enabled (engleză).

De fapt, următoarele periferice suportă în prezent tehnologia CUDA pe piața hardware pentru PC:

Versiune specificație	GPU	Plăci video
1.0	G80, G92, G92b, G94, G94b	GeForce 8800GTX/Ultra, 9400GT, 9600GT, 9800GT, Tesla C/D/S870, FX4/5600, 360M, GT 420
1.1	G86, G84, G98, G96, G96b, G94, G94b, G92, G92b	GeForce 8400GS/GT, 8600GT/GTS, 8800GT/GTS, 9600 GSO, 9800GTX/GX2, GTS 250, GT 120/30/40, FX 4/570, 3/580, 17/18/4700, 17/18/4700x /370M, 3/5/770M, 16/17/27/28/36/37/3800M, NVS420/50
1.2	GT218, GT216, GT215	GeForce 210, GT 220/40, FX380 LP, 1800M, 370/380M, NVS 2/3100M
1.3	GT200, GT200b	GeForce GTX 260, GTX 275, GTX 280, GTX 285, GTX 295, Tesla C/M1060, S1070, Quadro CX, FX 3/4/5800
2.0	GF100, GF110	GEFORCE (GF100) GTX 465, GTX 470, GTX 480, Tesla C2050, C2070, S/M2050/70, Quadro Plex 7000, Quadro 4000, 6000, GeForce (GF110) GTX 5.80, TX GTX 5.80, TX TX59, GTX59, GTX59 0
2.1	GF104, GF114, GF116, GF108, GF106	GeForce 610M, GT 430, GT 440, GTS 450, GTX 460, GTX 550 Ti, GTX 560, GTX 560 Ti, 500M, Quadro 600, 2000
3.0	GK104, GK106, GK107	GeForce GTX 690, GTX 680, GTX 670, GTX 660 Ti, GTX 660, GTX 650 Ti, GTX 650, GT 640, GeForce GTX 680MX, GeForce GTX 680M, GeForce GTX 675MX, GeForce GTX 675MX, GeForce GTX6060, GeForce GTX 680MX , GeForce GT 645M, GeForce GT 640M
3.5	GK110

Nvidia GeForce pentru computere desktop GeForce GTX 590 GeForce GTX 580 GeForce GTX 570 GeForce GTX 560 Ti GeForce GTX 560 GeForce GTX 550 Ti GeForce GTX 520 GeForce GTX 480 GeForce GTX 470 GeForce GTX 465 GeForce GTX 460 GeForce GTS 450 GeForce GTX 295 GeForce GTX 285 GeForce GTX 280 GeForce GTX 275 GeForce GTX 260 GeForce GTS 250 GeForce GT 240 GeForce GT 220 GeForce 210 GeForce GTS 150 GeForce GT 130 GeForce GT 120 GeForce G100 GeForce 9800 GX2 GeForce 9800 GTX+ GeForce 9800 GTX GeForce 9800 GT GeForce 9600 GSO GeForce 9600 GT GeForce 9500 GT GeForce 9400 GT GeForce 9400 mGPU GeForce 9300 mGPU GeForce 8800 GTS 512 GeForce 8800 GT GeForce 8600 GTS GeForce 8600 GT GeForce 8500 GT GeForce 8400GS

Nvidia GeForce pentru computere mobile GeForce GTX 580M GeForce GTX 570M GeForce GTX 560M GeForce GT 555M GeForce GT 540M GeForce GT 525M GeForce GT 520M GeForce GTX 485M GeForce GTX 480M GeForce GTX 470M GeForce GTX 460M GeForce GT 445M GeForce GT 435M GeForce GT 425M GeForce GT 420M GeForce GT 415M GeForce GTX 285M GeForce GTX 280M GeForce GTX 260M GeForce GTS 360M GeForce GTS 350M GeForce GTS 160M GeForce GTS 150M GeForce GT 335M GeForce GT 330M GeForce GT 325M GeForce GT 240M GeForce GT 130M GeForce G210M GeForce G110M GeForce G105M GeForce 310M GeForce 305M GeForce 9800M GTX GeForce 9800M GT GeForce 9800M GTS GeForce 9700M GTS GeForce 9700M GT GeForce 9650MGS GeForce 9600M GT GeForce 9600MGS GeForce 9500MGS GeForce 9500M G GeForce 9300MGS GeForce 9300M G GeForce 9200MGS GeForce 9100M G GeForce 8800M GTS GeForce 8700M GT GeForce 8600M GT GeForce 8600MGS GeForce 8400M GT GeForce 8400MGS

Nvidia Tesla * Tesla C2050/C2070 Tesla M2050/M2070/M2090 Tesla S2050 Tesla S1070 Tesla M1060 Tesla C1060 Tesla C870 Tesla D870 Tesla S870

Nvidia Quadro pentru computere desktop Quadro 6000 Quadro 5000 Quadro 4000 Quadro 2000 Quadro 600 Quadro FX 5800 Quadro FX 5600 Quadro FX 4800 Quadro FX 4700 X2 Quadro FX 4600 Quadro FX 3700 Quadro FX 1700 Quadro FX 570 Quadro FX 470 Quadro FX 380 Low Profile Quadro FX 370 Quadro FX 370 Low Profile Quadro CX Quadro NVS 450 Quadro NVS 420 Quadro NVS 290 Quadro Plex 2100 D4 Quadro Plex 2200 D2 Quadro Plex 2100 S4 Quadro Plex 1000 Model IV

Nvidia Quadro pentru computere mobile Quadro 5010M Quadro 5000M Quadro 4000M Quadro 3000M Quadro 2000M Quadro 1000M Quadro FX 3800M Quadro FX 3700M Quadro FX 3600M Quadro FX 2800M Quadro FX 2700M Quadro FX 1800M Quadro FX 1700M Quadro FX 1600M Quadro FX 880M Quadro FX 770M Quadro FX 570M Quadro FX 380M Quadro FX 370M Quadro FX 360M Quadro NVS 5100M Quadro NVS 4200M Quadro NVS 3100M Quadro NVS 2100M Quadro NVS 320M Quadro NVS 160M Quadro NVS 150M Quadro NVS 140M Quadro NVS 135M Quadro NVS 130M

• Modelele Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 permit calcule GPU cu dublă precizie.

Caracteristici și specificații ale diferitelor versiuni

Suport pentru funcții (funcțiile nelistate sunt acceptat pentru toate capabilitățile de calcul)	Capacitate de calcul (versiune)
	1.0	1.1	1.2	1.3	2.x
Cuvinte de 32 de biți în memoria globală	Nu	Da
valori în virgulă mobilă în memoria globală
Funcții atomice întregi care funcționează Cuvinte pe 32 de biți în memoria partajată	Nu		Da
atomicExch() funcționează pe 32 de biți valori în virgulă mobilă în memoria partajată
Funcții atomice întregi care funcționează Cuvinte pe 64 de biți în memoria globală
Funcții de vot Warp
Operații în virgulă mobilă cu dublă precizie	Nu			Da
Funcții atomice care funcționează pe 64 de biți valori întregi în memoria partajată	Nu				Da
Adunarea atomică în virgulă mobilă funcționează Cuvinte pe 32 de biți în memorie globală și partajată
_vot()
_sistem_threadfence()
_syncthreads_count(), _syncthreads_and(), _syncthreads_or()
Funcții de suprafață
Grilă 3D de blocuri de filet

Specificatii tehnice	Capacitate de calcul (versiune)
	1.0	1.1	1.2	1.3	2.x
Dimensiunea maximă a rețelei blocurilor de filet	2				3
Dimensiunea maximă x, y sau z a unei rețele de blocuri de filet	65535
Dimensiunea maximă a blocului de filet	3
Dimensiunea maximă x sau y a unui bloc	512				1024
Dimensiunea z maximă a unui bloc	64
Numărul maxim de fire pe bloc	512				1024
Dimensiunea urzelii	32
Numărul maxim de blocuri rezidente per multiprocesor	8
Numărul maxim de warps rezidenți per multiprocesor	24		32		48
Numărul maxim de fire rezidente per multiprocesor	768		1024		1536
Numărul de registre pe 32 de biți per multiprocesor	8K		16 K		32 K
Cantitatea maximă de memorie partajată per multiprocesor	16 KB				48 KB
Numărul de bănci de memorie partajată	16				32
Cantitatea de memorie locală per fir	16 KB				512 KB
Dimensiunea memoriei constantă	64 KB
Set de lucru cache per multiprocesor pentru memorie constantă	8 KB
Set de lucru cache per multiprocesor pentru memoria texturii	În funcție de dispozitiv, între 6 KB și 8 KB
Lățimea maximă pentru textura 1D	8192				32768
Lățimea maximă pentru textura 1D referință legată de memoria liniară	2 27
Lățimea maximă și numărul de straturi pentru o referință de textură stratificată 1D	8192 x 512				16384 x 2048
Lățimea și înălțimea maximă pentru 2D referință de textură legată de memorie liniară sau o matrice CUDA	65536 x 32768				65536 x 65535
Lățimea, înălțimea și numărul maxime de straturi pentru o referință de textură stratificată 2D	8192 x 8192 x 512				16384 x 16384 x 2048
Latime, inaltime si adancime maxime pentru o referință de textură 3D legată la liniar memorie sau o matrice CUDA	2048 x 2048 x 2048
Număr maxim de texturi care poate fi legat de un nucleu	128
Lățimea maximă pentru o suprafață 1D referință legată la o matrice CUDA	Nu sprijinit				8192
Lățimea și înălțimea maximă pentru un 2D referință de suprafață legată de o matrice CUDA					8192 x 8192
Numar maxim de suprafete care poate fi legat de un nucleu					8
Numărul maxim de instrucțiuni per nucleu	2 milioane

Exemplu

CudaArray* cu_array;< float , 2 >textură< float>tex; // Alocați matrice cudaMalloc( & cu_array, cudaCreateChannelDesc(), latime, inaltime);<<< gridDim, blockDim, 0 >// Copiați datele imaginii în matrice cudaMemcpy( cu_array, imagine, width* height, cudaMemcpyHostToDevice) ;

Importă pycuda.driver ca drv import numpy drv.init() dev = drv.Device(0) ctx = dev.make_context() mod = drv.SourceModule( """ __global__ void multiply_them(float *dest, float *a, float *b) ( const int i = threadIdx.x; dest[i] = a[i] * b[i]; ) """) multiply_them = mod.get_function ("multiply_them" ) a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 ) dest = numpy (a) multiply_them( drv.Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

CUDA ca subiect în universități

În decembrie 2009, modelul software CUDA este predat în 269 de universități din întreaga lume. În Rusia, cursurile de formare despre CUDA sunt susținute la Universitatea Politehnică din Sankt Petersburg, Universitatea de Stat din Yaroslavl. P. G. Demidov, Moscova, Nijni Novgorod, Sankt Petersburg, Tver, Kazan, Novosibirsk, Universitatea Tehnică de Stat Novosibirsk, Universitățile de Stat Omsk și Perm, Universitatea Internațională a Naturii Societății și Man „Dubna”, Universitatea de Stat de Energie Ivanovo, Universitatea de Stat Belgorod , MSTU-i. Bauman, Universitatea Tehnică Chimică Rusă numită după. Mendeleev, Centrul Interregional de Supercomputer RAS, . În plus, în decembrie 2009, a fost anunțat că a început să funcționeze primul centru științific și educațional rus „Parallel Computing”, situat în orașul Dubna, ale cărui sarcini includ instruire și consultări privind rezolvarea problemelor complexe de calcul pe GPU-uri.

În Ucraina, cursurile despre CUDA sunt predate la Institutul de Analiză de Sistem din Kiev.

Legături

Resurse oficiale

• Zona CUDA (rusă) - site-ul oficial CUDA
• CUDA GPU Computing (engleză) - forumuri web oficiale dedicate calculului CUDA

Resurse neoficiale

Hardware-ul lui Tom

• Dmitri Cekanov. nVidia CUDA: calcul pe o placă video sau moartea procesorului? . Tom's Hardware (22 iunie 2008).
• Dmitri Cekanov. nVidia CUDA: Benchmarking aplicații GPU pentru piața de masă. Tom's Hardware (19 mai 2009 Arhivat din original la 4 martie 2012. Consultat la 19 mai 2009).

iXBT.com

• Alexey Berillo. NVIDIA CUDA - calcul non-grafic pe GPU-uri. Partea 1. iXBT.com (23 septembrie 2008). Arhivat din original pe 4 martie 2012. Consultat la 20 ianuarie 2009.
• Alexey Berillo. NVIDIA CUDA - calcul non-grafic pe GPU-uri. Partea 2. iXBT.com (22 octombrie 2008). - Exemple de implementare a NVIDIA CUDA. Arhivat din original pe 4 martie 2012. Consultat la 20 ianuarie 2009.

Alte resurse

• Boreskov Alexey Viktorovici. Bazele CUDA (20 ianuarie 2009). Arhivat din original pe 4 martie 2012. Consultat la 20 ianuarie 2009.
• Vladimir Frolov. Introducere în tehnologia CUDA. Revista online „Computer Graphics and Multimedia” (19 decembrie 2008). Arhivat din original pe 4 martie 2012. Consultat la 28 octombrie 2009.
• Igor Oskolkov. NVIDIA CUDA este un bilet accesibil în lumea computerelor mari. Computerra (30 aprilie 2009). Preluat la 3 mai 2009.
• Vladimir Frolov. Introducere în Tehnologia CUDA (1 august 2009). Arhivat din original pe 4 martie 2012. Consultat la 3 aprilie 2010.
• GPGPU.ru. Utilizarea plăcilor video pentru calcul
• . Centrul de calcul paralel

Note

Vezi de asemenea

Nvidia
Grafic procesoare	Devreme NV1 NV2 familia RIVA TNT TNT2 Familia GeForce GeForce 256 GeForce 2 GeForce 3 GeForce 4 GeForce FX GeForce 6 GeForce 7 GeForce 8 GeForce 9 GeForce 100 GeForce 200 GeForce 400 GeForce 500 GeForce 600

De zeci de ani este în vigoare Legea lui Moore, care spune că la fiecare doi ani numărul de tranzistori de pe un cip se va dubla. Cu toate acestea, acest lucru a fost în 1965, iar în ultimii 5 ani ideea de multi-core fizice în procesoarele de clasă consumatorilor a început să se dezvolte rapid: în 2005, Intel a introdus Pentium D, iar AMD a introdus Athlon X2. Pe atunci, aplicațiile care foloseau 2 nuclee puteau fi numărate pe degetele unei mâini. Cu toate acestea, următoarea generație de procesoare Intel, care a făcut o revoluție, avea exact 2 nuclee fizice. Mai mult, seria Quad a apărut în ianuarie 2007, moment în care Moore însuși a recunoscut că legea sa va înceta în curând să se aplice.

Ce acum? Procesoarele dual-core chiar și în sistemele de birou bugetare, și 4 nuclee fizice au devenit norma și asta în doar 2-3 ani. Frecvența procesoarelor nu este crescută, dar arhitectura este îmbunătățită, numărul de nuclee fizice și virtuale este crescut. Cu toate acestea, ideea de a folosi adaptoare video echipate cu zeci sau chiar sute de „unități” de calcul există de mult timp.

Și deși perspectivele pentru GPU computing sunt enorme, cea mai populară soluție este Nvidia CUDA, care este gratuită, are multă documentație și este în general foarte ușor de implementat, nu există multe aplicații care folosesc această tehnologie. Practic, acestea sunt tot felul de calcule specializate, de care utilizatorul obișnuit în majoritatea cazurilor nu îi pasă. Dar există și programe concepute pentru utilizatorul de masă și despre ele vom vorbi în acest articol.

În primul rând, puțin despre tehnologia în sine și cu ce este folosită. Deoarece Când scriu un articol, mă concentrez pe o gamă largă de cititori, așa că voi încerca să-l explic într-un limbaj accesibil fără termeni complexi și oarecum pe scurt.

CUDA(English Compute Unified Device Architecture) este o arhitectură software și hardware care vă permite să efectuați calcule folosind procesoare grafice NVIDIA care acceptă tehnologia GPGPU (calcul arbitrar pe plăci video). Arhitectura CUDA a apărut pentru prima dată pe piață odată cu lansarea cipului NVIDIA de generația a opta - G80 și este prezentă în toate seriile ulterioare de cipuri grafice utilizate în familiile de acceleratoare GeForce, Quadro și Tesla. (c) Wikipedia.org

Fluxurile primite sunt procesate independent unul de celălalt, de exemplu. paralel.

Există o împărțire în 3 niveluri:

Grilă– miez. Conține o matrice una/două/tridimensională de blocuri.

Bloc– conține multe fire. Firele diferitelor blocuri nu pot interacționa între ele. De ce a fost necesar să se introducă blocuri? Fiecare bloc este în esență responsabil pentru propria sa sarcină secundară. De exemplu, o imagine mare (care este o matrice) poate fi împărțită în mai multe părți mai mici (matrici) și poate lucra cu fiecare parte a imaginii în paralel.

Fir– curgere. Firele dintr-un singur bloc pot interacționa fie prin intermediul memoriei partajate, care, apropo, este mult mai rapidă decât memoria globală, fie prin instrumente de sincronizare a firelor.

Urzeală este o uniune de fire care interacționează între ele, pentru toate GPU-urile moderne dimensiunea Warp este de 32. Urmează jumătate urzeală, care este jumătate din urzeală, pentru că Accesul la memorie are loc de obicei separat pentru prima și a doua jumătate a warp-ului.

După cum puteți vedea, această arhitectură este excelentă pentru paralelizarea sarcinilor. Și deși programarea se realizează în limbajul C cu unele restricții, în realitate nu totul este atât de simplu, pentru că... nu totul poate fi paralelizat. De asemenea, nu există funcții standard pentru generarea de numere aleatoare (sau inițializare, toate acestea trebuie implementate separat). Și deși există o mulțime de opțiuni gata făcute, nimic din toate acestea nu aduce bucurie. Capacitatea de a folosi recursiunea a apărut relativ recent.

Pentru claritate, a fost scris un mic program de consolă (pentru a minimiza codul) care efectuează operațiuni cu două matrice de tip float, adică. cu valori non-întregi. Din motivele expuse mai sus, inițializarea (umplerea matricei cu diferite valori arbitrare) a fost efectuată de CPU. În continuare, au fost efectuate 25 de operații diferite pe elementele corespunzătoare din fiecare matrice, iar rezultatele intermediare au fost înregistrate în a treia matrice. Dimensiunea matricei s-a schimbat, rezultatele sunt următoarele:

Au fost efectuate în total 4 teste:

1024 de elemente în fiecare matrice:

Se vede clar că, cu un număr atât de mic de elemente, calculul paralel este de puțin folos, deoarece Calculele în sine sunt mult mai rapide decât pregătirea lor.

4096 elemente în fiecare matrice:

Și acum puteți vedea că placa video efectuează operațiuni pe matrice de 3 ori mai rapid decât procesorul. Mai mult, timpul de execuție a acestui test pe placa video nu a crescut (o scădere ușoară a timpului poate fi atribuită unei erori).

Acum există 12288 de elemente în fiecare matrice:

Separarea plăcii video a crescut de 2 ori. Din nou, este de remarcat faptul că timpul de execuție pe placa video a crescut
nesemnificativ, dar pe procesor de mai mult de 3 ori, i.e. proporțional cu complexitatea sarcinii.

Și ultimul test este de 36864 de elemente în fiecare matrice:

În acest caz, accelerația atinge valori impresionante - de aproape 22 de ori mai rapid pe o placă video. Și din nou, timpul de execuție pe placa video a crescut ușor, dar pe procesor - necesarul de 3 ori, ceea ce este din nou proporțional cu complexitatea sarcinii.

Dacă continui să complici calculele, placa video câștigă din ce în ce mai mult. Deși exemplul este oarecum exagerat, situația generală arată clar. Dar, așa cum am menționat mai sus, nu totul poate fi paralelizat. De exemplu, calcularea Pi. Există doar exemple scrise folosind metoda Monte Carlo, dar precizia calculului este de 7 zecimale, adică. plutitor obișnuit. Pentru a crește acuratețea calculelor, este necesară o aritmetică lungă și aici apar probleme, deoarece Este foarte, foarte dificil să implementezi acest lucru în mod eficient. Nu am putut găsi exemple pe Internet care să folosească CUDA și să calculeze Pi la 1 milion de zecimale. S-au încercat să scrie o astfel de aplicație, dar cea mai simplă și eficientă metodă de calcul al Pi este algoritmul Brent-Salamin sau formula Gauss. Cunoscutul SuperPI cel mai probabil (judecând după viteza de funcționare și numărul de iterații) folosește formula Gauss. Și, judecând după
Datorită faptului că SuperPI este cu un singur thread, a lipsei de exemple sub CUDA și a eșecului încercărilor mele, este imposibil să paralelizez efectiv numărarea Pi.

Apropo, puteți observa cum crește sarcina pe GPU în timpul calculelor, precum și alocarea memoriei.

Acum să trecem la beneficiile mai practice ale CUDA, și anume la programele existente în prezent care folosesc această tehnologie. În cea mai mare parte, acestea sunt tot felul de convertoare și editoare audio/video.

În testare au fost utilizate 3 fișiere video diferite:

*Istoria realizării filmului Avatar - 1920x1080, MPEG4, h.264.

*Seria „Minți-mă” - 1280x720, MPEG4, h.264.

*Seria „It’s Always Sunny in Philadelphia” - 624x464, xvid.

Containerul și dimensiunea primelor două fișiere a fost .mkv și 1,55 GB, iar ultimul a fost .avi și 272 MB.

Să începem cu un produs foarte senzațional și popular - Badaboom. Versiunea folosita - 1.2.1.74 . Costul programului este $29.90 .

Interfața programului este simplă și intuitivă - în stânga selectăm fișierul sau discul sursă, iar în dreapta - dispozitivul necesar pentru care îl vom codifica. Există și un mod de utilizator în care parametrii sunt setați manual, ceea ce am folosit.

În primul rând, să ne uităm la cât de rapid și eficient este codificat videoclipul „în sine”, adică. aceeași rezoluție și aproximativ aceeași dimensiune. Vom măsura viteza în fps și nu în timpul scurs - în acest fel este mai convenabil să comparăm și să calculăm cât de mult va fi comprimat un videoclip de lungime arbitrară. Deoarece Astăzi luăm în considerare tehnologia „verde”, atunci graficele vor fi corespunzătoare -)

Viteza de codificare depinde direct de calitate, acest lucru este evident. Este de remarcat faptul că rezoluția luminii (să-i spunem în mod tradițional SD) nu este o problemă pentru Badaboom - viteza de codificare este de 5,5 ori mai mare decât rata de cadre video originală (24 fps). Și chiar și videoclipurile grele de 1080p sunt convertite de program în timp real. Este de remarcat faptul că calitatea videoclipului final este foarte apropiată de materialul video original, adică Badaboom codifică foarte, foarte eficient.

Dar de obicei transferă video la o rezoluție mai mică, să vedem cum stau lucrurile în acest mod. Pe măsură ce rezoluția a scăzut, a scăzut și rata de biți video. A fost 9500 kbps pentru fișierul de ieșire 1080p, 4100 kbps pentru 720p și 2400 kbps pentru 720x404. Alegerea a fost făcută pe baza unui raport rezonabil dimensiune/calitate.

Nu sunt necesare comentarii. Dacă extrageți un 720p la o calitate SD obișnuită, atunci transcodarea unui film de 2 ore va dura aproximativ 30 de minute. Și, în același timp, încărcarea procesorului va fi nesemnificativă, poți să-ți faci treaba fără a simți disconfort.

Ce se întâmplă dacă convertiți videoclipul într-un format pentru un dispozitiv mobil? Pentru a face acest lucru, selectați profilul iPhone (bitrate 1 Mbit/s, 480x320) și uitați-vă la viteza de codare:

Trebuie să spun ceva? Un film de două ore la calitate iPhone normală este transcodat în mai puțin de 15 minute. Cu calitate HD este mai dificil, dar totuși foarte rapid. Principalul lucru este că calitatea materialului video de ieșire rămâne la un nivel destul de ridicat atunci când este vizualizat pe un afișaj al telefonului.

În general, impresiile de la Badaboom sunt pozitive, viteza de operare este plăcută, iar interfața este simplă și clară. Au fost remediate tot felul de erori din versiunile anterioare (am folosit beta în 2008). Cu excepția unui singur lucru - calea către fișierul sursă, precum și către folderul în care este salvat videoclipul finalizat, nu ar trebui să conțină litere rusești. Dar în comparație cu avantajele programului, acest dezavantaj este nesemnificativ.

Următorul pe rând vom avea Super LoiLoScope. Pentru versiunea obișnuită ei întreabă 3.280 de ruble, iar pentru versiunea touch, care acceptă controlul tactil în Windows 7, ei cer la fel de mult 4.440 de ruble. Să încercăm să ne dăm seama de ce dezvoltatorul dorește astfel de bani și de ce editorul video are nevoie de suport multitouch. Ultima versiune folosită - 1.8.3.3 .

Este destul de dificil să descrii interfața programului în cuvinte, așa că am decis să fac un scurt videoclip. Voi spune imediat că, la fel ca toate convertoarele video pentru CUDA, accelerarea GPU este acceptată numai pentru ieșirea video în MPEG4 cu codecul h.264.

În timpul codificării, sarcina procesorului este de 100%, dar acest lucru nu provoacă disconfort. Browserul și alte aplicații ușoare nu încetinesc.

Acum să trecem la performanță. Pentru început, totul este la fel ca și cu Badaboom - transcodând videoclipul într-unul similar în calitate.

Rezultatele sunt mult mai bune decât Badaboom. Calitatea este si ea excelenta, diferenta cu originalul se observa doar prin compararea ramelor in perechi sub lupa.

Uau, aici LoiloScope îl depășește pe Badaboom de 2,5 ori. În același timp, puteți tăia și codifica cu ușurință un alt videoclip în paralel, să citiți știri și chiar să vizionați filme și chiar și FullHD se redă fără probleme, deși încărcarea procesorului este maximă.

Acum să încercăm să facem un videoclip pentru un dispozitiv mobil, să numim profilul la fel cum a fost numit în Badaboom - iPhone (480x320, 1 Mbit/s):

Nu există nicio eroare. Totul a fost reverificat de mai multe ori, de fiecare dată rezultatul a fost același. Cel mai probabil, acest lucru se întâmplă din simplul motiv că fișierul SD a fost înregistrat cu un codec diferit și într-un container diferit. La transcodare, videoclipul este mai întâi decodat, împărțit în matrice de o anumită dimensiune și comprimat. Decodorul ASP folosit în cazul xvid este mai lent decât AVC (pentru h.264) când se decodează în paralel. Cu toate acestea, 192 fps este de 8 ori mai rapidă decât viteza videoclipului original, o serie de 23 de minute este comprimată în mai puțin de 4 minute. Situația s-a repetat cu alte fișiere comprimate în xvid/DivX.

LoiloScope Am lăsat doar impresii plăcute - interfața, în ciuda caracterului său neobișnuit, este convenabilă și funcțională, iar viteza de funcționare este dincolo de laudă. Funcționalitatea relativ slabă este oarecum frustrantă, dar adesea, cu o instalare simplă, trebuie doar să ajustați ușor culorile, să faceți tranziții netede și să adăugați text, iar LoiloScope face o treabă excelentă în acest sens. Prețul este, de asemenea, oarecum înspăimântător - mai mult de 100 de dolari pentru versiunea obișnuită este normal pentru țările străine, dar astfel de cifre ni se par încă puțin sălbatice. Deși, recunosc că, dacă eu, de exemplu, am filmat și editat des videoclipuri acasă, s-ar putea să mă fi gândit să-l cumpăr. În același timp, apropo, am verificat posibilitatea de a edita conținut HD (sau mai bine zis AVCHD) direct de pe o cameră video fără a converti mai întâi în alt format LoiloScope nu a scos la iveală nicio problemă cu fișiere precum .mts;

Tehnologia CUDA

Vladimir Frolov,[email protected]

Adnotare

Articolul vorbește despre tehnologia CUDA, care permite unui programator să folosească plăcile video ca unități de calcul puternice.

Instrumentele oferite de Nvidia fac posibilă scrierea de programe de unități de procesare grafică (GPU) într-un subset al limbajului C++. Acest lucru scutește programatorul de necesitatea de a utiliza shadere și de a înțelege funcționarea conductei grafice. Articolul oferă exemple de programare folosind CUDA și diverse tehnici de optimizare.

1. Introducere

Dezvoltarea tehnologiei de calcul a progresat rapid în ultimele decenii. Atât de repede încât dezvoltatorii de procesoare au ajuns deja aproape în așa-numitul „punctul mort al siliconului”. O creștere nestăpânită a frecvenței ceasului a devenit imposibilă din cauza unui număr de motive tehnologice serioase.

Acesta este parțial motivul pentru care toți producătorii de sisteme de calcul moderne se îndreaptă spre creșterea numărului de procesoare și nuclee, mai degrabă decât spre creșterea frecvenței unui procesor. Numărul de nuclee de unități centrale de procesare (CPU) în sistemele avansate este acum deja de 8.

Un alt motiv este viteza relativ scăzută a memoriei RAM. Indiferent cât de repede funcționează procesorul, blocajele, după cum arată practica, nu sunt operațiuni aritmetice, ci mai degrabă accesări nereușite la memorie - pierderile de cache.

Cu toate acestea, dacă priviți către procesoarele grafice GPU (Graphics Processing Unit), acestea au luat calea paralelismului mult mai devreme. În plăcile video de astăzi, de exemplu în GF8800GTX, numărul de procesoare poate ajunge la 128. Performanța unor astfel de sisteme, cu o programare pricepută, poate fi destul de semnificativă (Fig. 1).

Orez. 1. Numărul de operații în virgulă mobilă pentru CPU și GPU

Când primele plăci video au apărut pentru prima dată pe piață, acestea erau dispozitive destul de simple (în comparație cu procesorul central) foarte specializate, concepute pentru a scuti procesorul de sarcina vizualizării datelor bidimensionale. Odată cu dezvoltarea industriei jocurilor de noroc și apariția unor astfel de jocuri tridimensionale precum Doom (Fig. 2) și Wolfenstein 3D (Fig. 3), a apărut nevoia de vizualizare 3D.

Figurile 2,3. Jocuri Doom și Wolfenstein 3D

Programatorii nu au avut de ales în algoritmul de randare, iar pentru a crește flexibilitatea au apărut shadere - programe mici executate de placa video pentru fiecare vârf sau pentru fiecare pixel. Sarcinile lor au inclus transformări peste vârfuri și umbrire - calculul luminii într-un punct, de exemplu, folosind modelul Phong.

Deși shaders-urile au parcurs un drum lung în aceste zile, ar trebui să se înțeleagă că au fost concepute pentru sarcini foarte specializate, cum ar fi transformarea 3D și rasterizarea. În timp ce GPU-urile evoluează către sisteme multiprocesoare de uz general, limbajele shader rămân foarte specializate.

Ele pot fi comparate cu FORTRAN în sensul că, la fel ca FORTRAN, au fost primele, dar concepute pentru a rezolva un singur tip de problemă. Shaders sunt de puțin folos pentru rezolvarea oricăror alte probleme, altele decât transformările 3D și rasterizarea, la fel cum FORTRAN nu este convenabil pentru rezolvarea problemelor care nu sunt legate de calcule numerice.

Astăzi există tendința de a folosi plăcile video într-un mod neconvențional pentru a rezolva probleme din domeniile mecanicii cuantice, inteligenței artificiale, calculelor fizice, criptografiei, vizualizării corecte fizic, reconstrucției din fotografii, recunoașterii etc. Aceste sarcini sunt incomod de rezolvat în cadrul API-urilor grafice (DirectX, OpenGL), deoarece aceste API-uri au fost create pentru aplicații complet diferite.

Dezvoltarea programării cu scop general pe GPU (General Programming on GPU, GPGPU) a condus în mod logic la apariția tehnologiilor care vizează o gamă mai largă de sarcini decât rasterizarea. Drept urmare, Nvidia a creat Compute Unified Device Architecture (sau CUDA pe scurt), iar compania concurentă ATI a creat tehnologia STREAM.

Trebuie remarcat faptul că, la momentul scrierii acestui articol, tehnologia STREAM era cu mult în urmă CUDA în dezvoltare și, prin urmare, nu va fi luată în considerare aici. Ne vom concentra pe CUDA, o tehnologie GPGPU care vă permite să scrieți programe într-un subset al limbajului C++.

2. Diferența fundamentală între CPU și GPU

Să aruncăm o privire rapidă la unele dintre diferențele semnificative dintre zonele și caracteristicile aplicațiilor CPU și plăci video.

2.1. Posibilitati

CPU este proiectat inițial pentru rezolvarea problemelor generale și funcționează cu memorie adresabilă aleatoriu. Programele de pe CPU pot accesa direct orice celule de memorie liniare și omogene.

Acesta nu este cazul GPU-urilor. După cum veți afla citind acest articol, CUDA are până la 6 tipuri de memorie. Puteți citi din orice celulă care este accesibilă fizic, dar nu puteți scrie în toate celulele. Motivul este că GPU-ul este în orice caz un dispozitiv specific conceput pentru scopuri specifice. Această limitare a fost introdusă pentru a crește viteza anumitor algoritmi și a reduce costul echipamentelor.

2.2. Performanța memoriei

O problemă perenă cu majoritatea sistemelor de calcul este că memoria este mai lentă decât procesorul. Producătorii de procesoare rezolvă această problemă prin introducerea cache-urilor. Zonele de memorie cele mai frecvent utilizate sunt plasate în memorie sau în memoria cache, funcționând la frecvența procesorului. Acest lucru vă permite să economisiți timp atunci când accesați datele cele mai frecvent utilizate și să încărcați procesorul cu calculele reale.

Rețineți că cache-urile sunt în esență transparente pentru programator. Atât la citire, cât și la scriere, datele nu merg direct în RAM, ci trec prin cache. Acest lucru permite, în special, să citiți rapid o valoare imediat după ce o scrieți.

GPU-urile (aici ne referim la plăcile video din seria a opta GF) au și ele cache și sunt și ele importante, dar acest mecanism nu este la fel de puternic ca pe un procesor. În primul rând, nu toate tipurile de memorie sunt stocate în cache, iar în al doilea rând, cache-urile sunt doar pentru citire.

Pe GPU, accesele lente la memorie sunt ascunse folosind calculul paralel. În timp ce unele sarcini așteaptă date, altele lucrează, gata de calcule. Acesta este unul dintre principiile principale ale CUDA, care poate îmbunătăți considerabil performanța sistemului în ansamblu.

3. CUDA Core

3.1. Model de streaming

Arhitectura de calcul CUDA se bazează pe concepto singură comandă pentru multe date(Single Instruction Multiple Data, SIMD) și concept multiprocesor.

Conceptul SIMD înseamnă că o instrucțiune poate procesa mai multe date simultan. De exemplu, comanda addps pe Pentium 3 și pe procesoarele Pentium mai noi vă permite să adăugați simultan 4 numere în virgulă mobilă cu precizie unică.

Un multiprocesor este un procesor SIMD multi-core care permite executarea unei singure instrucțiuni pe toate nucleele la un moment dat. Fiecare nucleu multiprocesor este scalar, adică nu suportă operații pur vector.

Înainte de a continua, să introducem câteva definiții. Rețineți că în acest articol ne vom referi la dispozitiv și la gazdă într-un mod complet diferit de ceea ce sunt obișnuiți majoritatea programatorilor. Vom folosi acești termeni pentru a evita discrepanțe cu documentația CUDA.

În articolul nostru, prin dispozitiv vom înțelege un adaptor video care acceptă driverul CUDA, sau un alt dispozitiv specializat conceput pentru executarea de programe care utilizează CUDA (cum ar fi NVIDIA Tesla). În articolul nostru, vom considera GPU-ul doar ca pe un dispozitiv logic, evitând detalii specifice de implementare.

Vom numi o gazdă un program din memoria RAM obișnuită a computerului care utilizează procesorul și efectuează funcții de control pentru lucrul cu dispozitivul.

De fapt, partea din programul dumneavoastră care rulează pe CPU este gazdă, și placa ta video - dispozitiv. În mod logic, dispozitivul poate fi reprezentat ca un set de multiprocesoare (Fig. 4) plus un driver CUDA.

Orez. 4. Dispozitiv

Să presupunem că vrem să rulăm o anumită procedură pe dispozitivul nostru în N fire de execuție (adică vrem să-i paralelizăm activitatea). Conform documentației CUDA, să numim această procedură kernel.

O caracteristică a arhitecturii CUDA este organizarea bloc-grid, care este neobișnuită pentru aplicațiile multi-threaded (Fig. 5). În acest caz, driverul CUDA distribuie independent resursele dispozitivului între fire.

Orez. 5. Organizarea firelor

În fig. 5. Nucleul este desemnat ca Kernel. Toate firele care execută acest nucleu sunt combinate în blocuri (Block), iar blocurile, la rândul lor, sunt combinate într-o grilă (Grid).

După cum se poate observa în Figura 5, indicii bidimensionali sunt utilizați pentru a identifica firele. Dezvoltatorii CUDA au oferit posibilitatea de a lucra cu indici tridimensionali, bidimensionali sau simpli (unidimensionali), în funcție de ceea ce este mai convenabil pentru programator.

În general, indicii sunt vectori tridimensionali. Pentru fiecare fir se vor cunoaste urmatoarele: indicele firului din interiorul blocului threadIdx si indexul blocului din interiorul grilei blockIdx. La pornire, toate firele de execuție vor diferi doar în acești indecși. De fapt, prin intermediul acestor indici programatorul exercită controlul, determinând ce parte din datele sale este procesată în fiecare fir.

Răspunsul la întrebarea de ce dezvoltatorii au ales această organizație nu este banal. Unul dintre motive este că un bloc este garantat pentru a fi executat pe una dispozitiv multiprocesor, dar un multiprocesor poate executa mai multe blocuri diferite. Motivele rămase vor deveni clare mai târziu în articol.

Un bloc de sarcini (threads) este executat pe un multiprocesor în părți, sau pool-uri, numite warps. Dimensiunea warp actuală a plăcilor video cu suport CUDA este de 32 de fire. Sarcinile din interiorul pool-ului warp sunt executate în stil SIMD, de exemplu. Toate firele din warp pot executa doar o instrucțiune la un moment dat.

Un avertisment trebuie făcut aici. În arhitecturile moderne la momentul scrierii acestui articol, numărul de procesoare din interiorul unui multiprocesor este de 8, nu de 32. Rezultă că nu întregul warp este executat simultan, acesta este împărțit în 4 părți, care se execută secvenţial (întrucât procesoarele). sunt scalare).

Dar, în primul rând, dezvoltatorii CUDA nu reglementează strict dimensiunea warp. În lucrările lor, ei menționează parametrul mărimea warp, și nu numărul 32. În al doilea rând, din punct de vedere logic, warp este uniunea minimă de fire despre care putem spune că toate firele din interiorul său sunt executate simultan - și la în același timp, nu există ipoteze cu privire la restul, sistemul nu se va face.

3.1.1. Ramificare

Apare imediat întrebarea: dacă în același moment, toate firele din interiorul unui warp execută aceeași instrucțiune, atunci cum rămâne cu ramurile? La urma urmei, dacă există o ramură în codul programului, atunci instrucțiunile vor fi diferite. O soluție standard pentru programarea SIMD este utilizată aici (Figura 6).

Orez. 6. Organizarea ramificarii in SIMD

Să presupunem că avem următorul cod:

dacă (cond)B;

În cazul SISD (Single Instruction Single Data), executăm instrucțiunea A, verificăm condiția, apoi executăm instrucțiunile B și D (dacă condiția este adevărată).

Să avem acum 10 fire care rulează în stil SIMD. În toate cele 10 fire executăm instrucțiunea A, apoi verificăm condiția cond și reiese că în 9 din 10 fire este adevărat, iar într-un fir este fals.

Este clar că nu putem lansa 9 fire de execuție pentru a executa operatorul B, iar cel rămas pentru a executa operatorul C, deoarece în toate firele de execuție se poate executa o singură instrucțiune în același timp. În acest caz, trebuie să faceți acest lucru: mai întâi, „omorâm” firul separat, astfel încât să nu strice datele nimănui și executăm cele 9 fire rămase. Apoi „omorâm” 9 fire care au executat operatorul B și trecem printr-un fir cu operatorul C. După aceasta, firele se îmbină din nou și execută operatorul D în același timp.

Rezultatul trist este că nu numai că sunt irosite resursele procesorului pentru șlefuirea biților goale în firele împărțite, dar ceea ce este mult mai rău este că vom ajunge să fim nevoiți să executăm AMBELE ramuri.

Cu toate acestea, nu totul este atât de rău pe cât ar părea la prima vedere. Un avantaj foarte mare al tehnologiei este că aceste trucuri sunt efectuate dinamic de driverul CUDA și sunt complet transparente pentru programator. În același timp, atunci când se ocupă de comenzile SSE ale procesoarelor moderne (mai ales în cazul încercării de a executa simultan 4 copii ale algoritmului), programatorul însuși trebuie să aibă grijă de detalii: combinați datele în quad-uri, nu uitați de aliniere , și, în general, scrie la un nivel scăzut, de fapt, cum în assembler.

Din toate cele de mai sus, rezultă o concluzie foarte importantă. Ramurile nu provoacă degradarea performanței în sine. Singurele ramuri dăunătoare sunt cele în care firele diverg în cadrul aceluiași bazin de fir de urzeală. Mai mult, dacă firele diverg într-un singur bloc, dar în bazine de urzeală diferite sau în blocuri diferite, acest lucru nu are absolut niciun efect.

3.1.2. Comunicarea între fire

La momentul scrierii acestui articol, orice interacțiune între fire (sincronizare și schimb de date) era posibilă doar în cadrul unui bloc. Adică, este imposibil să se organizeze interacțiunea între fire de execuție ale diferitelor blocuri folosind doar capabilități documentate.

În ceea ce privește caracteristicile nedocumentate, utilizarea lor este foarte descurajată. Motivul pentru aceasta este că se bazează pe caracteristicile hardware specifice ale unui anumit sistem.

Sincronizarea tuturor sarcinilor dintr-un bloc se realizează prin apelarea funcției __synchtreads. Schimbul de date este posibil prin intermediul memoriei partajate, deoarece este comun tuturor sarcinilor din bloc.

3.2. Memorie

CUDA distinge șase tipuri de memorie (Fig. 7). Acestea sunt registre, locale, globale, partajate, constante și memorie de textură.

Această abundență se datorează specificului plăcii video și scopului său principal, precum și dorinței dezvoltatorilor de a face sistemul cât mai ieftin, sacrificând în diverse cazuri fie versatilitatea, fie viteza.

Orez. 7. Tipuri de memorie în CUDA

3.2.0. Registrele

Ori de câte ori este posibil, compilatorul încearcă să plaseze toate variabilele funcției locale în registre. Astfel de variabile sunt accesate cu viteza maxima. În arhitectura actuală, există 8192 de registre pe 32 de biți disponibile per multiprocesor. Pentru a determina câte registre sunt disponibile pentru un fir, trebuie să împărțiți acest număr (8192) la dimensiunea blocului (numărul de fire din acesta).

Cu împărțirea obișnuită de 64 de fire pe bloc, rezultatul este doar 128 de registre (există câteva criterii obiective, dar 64 este potrivit în medie pentru multe sarcini). În realitate, nvcc nu va aloca niciodată 128 de registre. De obicei, nu dă mai mult de 40, iar variabilele rămase vor intra în memoria locală. Acest lucru se întâmplă deoarece mai multe blocuri pot fi executate pe un multiprocesor. Compilatorul încearcă să maximizeze numărul de blocuri de lucru simultan. Pentru o eficiență mai mare, ar trebui să încercați să ocupați mai puțin de 32 de registre. Apoi, teoretic, 4 blocuri (8 warps, dacă 64 fire într-un bloc) pot fi lansate pe un multiprocesor. Totuși, aici ar trebui să țineți cont și de cantitatea de memorie partajată ocupată de fire, deoarece dacă un bloc ocupă întreaga memorie partajată, două astfel de blocuri nu pot fi executate simultan pe un multiprocesor.

3.2.1. Memoria locală

În cazurile în care datele procedurii locale sunt prea mari sau compilatorul nu poate calcula un pas constant de acces pentru ele, le poate pune în memoria locală. Acest lucru poate fi facilitat, de exemplu, prin turnarea de indicatori pentru tipuri de dimensiuni diferite.

Din punct de vedere fizic, memoria locală este analogă memoriei globale și funcționează cu aceeași viteză. La momentul scrierii, nu există mecanisme care să împiedice în mod explicit compilatorul să folosească memoria locală pentru anumite variabile. Deoarece este destul de dificil să controlați memoria locală, este mai bine să nu o utilizați deloc (vezi secțiunea 4 „Recomandări de optimizare”).

3.2.2. Memoria globală

În documentația CUDA ca una dintre principalele realizăritehnologia oferă posibilitatea unei adrese arbitrare a memoriei globale. Adică, puteți citi din orice celulă de memorie și puteți scrie și într-o celulă arbitrară (de obicei nu este cazul pe un GPU).

Cu toate acestea, în acest caz trebuie să plătiți pentru versatilitate cu viteza. Memoria globală nu este stocată în cache. Funcționează foarte lent, numărul de accesări la memoria globală ar trebui minimizat în orice caz.

Memoria globală este necesară în primul rând pentru a stoca rezultatele unui program înainte de a le trimite către gazdă (în DRAM obișnuit). Motivul pentru aceasta este că memoria globală este singurul tip de memorie în care poți scrie orice.

Variabilele declarate cu calificativul __global__ sunt alocate memoriei globale. Memoria globală poate fi, de asemenea, alocată dinamic prin apelarea cudaMalloc(void* mem, int size) pe gazdă. Această funcție nu poate fi apelată de pe dispozitiv. Rezultă că alocarea memoriei ar trebui să fie gestionată de programul gazdă care rulează pe CPU. Datele de la gazdă pot fi trimise către dispozitiv apelând funcția cudaMemcpy:

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

Puteți efectua procedura inversă exact în același mod:

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

Acest apel se face și de la gazdă.

Când lucrați cu memoria globală, este important să urmați regula coalescerii. Ideea principală este că firele de execuție trebuie să acceseze celule de memorie consecutive, de 4, 8 sau 16 octeți. În acest caz, primul fir trebuie să acceseze o adresă aliniată la o limită de 4, 8 sau, respectiv, 16 octeți. Adresele returnate de cudaMalloc sunt aliniate pe o limită de cel puțin 256 de octeți.

3.2.3. Memoria partajată

Memoria partajată nu poate fi stocată în cache, dar este rapidă. Este recomandat să îl utilizați ca cache gestionat. Doar 16 KB de memorie partajată sunt disponibile pentru fiecare multiprocesor. Împărțind acest număr la numărul de sarcini din bloc, obținem cantitatea maximă de memorie partajată disponibilă pe fir (dacă intenționați să o utilizați independent în toate firele).

O caracteristică distinctivă a memoriei partajate este că este abordată în mod egal pentru toate sarcinile dintr-un bloc (Fig. 7). Rezultă că poate fi folosit pentru a face schimb de date între fire de execuție dintr-un singur bloc.

Este garantat că în timpul execuției blocului pe un multiprocesor, conținutul memoriei partajate va fi păstrat. Cu toate acestea, după ce un bloc a fost schimbat pe un multiprocesor, nu este garantat că conținutul vechiului bloc va fi păstrat. Prin urmare, nu ar trebui să încercați să sincronizați sarcinile între blocuri, lăsând orice date în memoria partajată și sperând în siguranță.

Variabilele declarate cu calificativul __shared__ sunt alocate în memoria partajată.

Shared__ float mem_shared;

Trebuie subliniat încă o dată că există o singură memorie partajată pentru un bloc. Prin urmare, dacă trebuie să-l utilizați pur și simplu ca cache gestionat, ar trebui să accesați diferite elemente ale matricei, de exemplu, astfel:

float x = mem_shared;

Unde threadIdx.x este indexul x al firului din interiorul blocului.

3.2.4. Memoria constantă

Memoria constantă este stocată în cache, așa cum se poate vedea în Fig. 4. Cache-ul există într-o singură copie pentru un multiprocesor, ceea ce înseamnă că este comun pentru toate sarcinile din bloc. Pe gazdă, puteți scrie ceva în memoria constantă apelând funcția cudaMemcpyToSymbol. De pe dispozitiv, memoria constantă este doar pentru citire.

Memoria constantă este foarte convenabilă de utilizat. Puteți plasa date de orice tip în el și le puteți citi folosind o sarcină simplă.

#definiți N 100

Constant__ int gpu_buffer[N];

void gazdă_funcție()

int cpu_buffer[N];

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

// __global__ înseamnă device_kernel este un nucleu care poate fi rulat pe GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; EROARE! memoria constantă este doar pentru citire

Deoarece memoria constantă utilizează un cache, accesul la aceasta este în general destul de rapid. Singurul dezavantaj, dar foarte mare, al memoriei constante este că dimensiunea sa este de doar 64 KB (pentru întregul dispozitiv). Acest lucru implică faptul că are sens să stocați doar o cantitate mică de date utilizate frecvent în memoria contextuală.

3.2.5. Memoria texturii

Memoria texturii este stocată în cache (Fig. 4). Există un singur cache pentru fiecare multiprocesor, ceea ce înseamnă că acest cache este comun tuturor sarcinilor din bloc.

Numele memoriei texturii (și, din păcate, funcționalității) este moștenit de la conceptele de „textură” și „texturare”. Texturarea este procesul de aplicare a unei texturi (doar o imagine) unui poligon în timpul procesului de rasterizare. Memoria texturii este optimizată pentru eșantionarea datelor 2D și are următoarele capacități:

preluarea rapidă a valorilor de dimensiune fixă ​​(octet, cuvânt, cuvânt dublu sau patru) dintr-o matrice unidimensională sau bidimensională;

adresare normalizată cu numere float în intervalul . Apoi le puteți selecta folosind adresarea normalizată. Valoarea rezultată va fi un cuvânt de tip float4 mapat la interval;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // alocă memorie în GPU

// setarea parametrilor de textură textura

Texture.addressMode = cudaAddressModeWrap; // modul Înfășurați

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //cea mai apropiată valoare

Texture.normalized = fals; // nu folosiți adresarea normalizată

CudaBindTexture(0, texture, gpu_memory, N ) // de acum înainte această memorie va fi considerată memorie de textură


CudaMemcpy(gpu_memory, cpu_buffer, N*sizeof(uint 4), cudaMemcpyHostToDevice ); // copiază datele înGPU

// __global__ înseamnă că device_kernel este nucleul care trebuie paralelizat

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // puteți selecta datele doar în acest fel!

uint4 b = tex1Dfetch(texture,1);

int c = a.x * b.y;

...

3.3. Exemplu simplu

Ca exemplu simplu, luați în considerare programul cppIntegration din SDK-ul CUDA. Demonstrează tehnici de lucru cu CUDA, precum și utilizarea nvcc (un subset special al compilatorului C++ de la Nvidia) în combinație cu MS Visual Studio, ceea ce simplifică foarte mult dezvoltarea programelor pe CUDA.

4.1. Împărțiți-vă sarcina corect

Nu toate sarcinile sunt potrivite pentru arhitecturile SIMD. Dacă sarcina dvs. nu este potrivită pentru aceasta, este posibil să nu merite să utilizați un GPU. Dar dacă sunteți hotărât să utilizați un GPU, ar trebui să încercați să descompuneți algoritmul în bucăți care pot fi executate eficient în stil SIMD. Dacă este necesar, schimbați algoritmul pentru a vă rezolva problema, veniți cu unul nou - unul care s-ar potrivi bine cu SIMD. Un exemplu de zonă potrivită pentru utilizarea GPU-ului este implementarea adăugării piramidale a elementelor matricei.

4.2. Selectarea tipului de memorie

Plasați datele în textură sau în memorie constantă dacă toate sarcinile din același bloc accesează aceeași locație de memorie sau zone apropiate. Datele 2D pot fi procesate eficient folosind funcțiile text2Dfetch și text2D. Memoria texturii este optimizată special pentru eșantionarea 2D.

Utilizați memoria globală în combinație cu memoria partajată dacă toate sarcinile accesează aleatoriu locații de memorie diferite, larg separate (cu adrese sau coordonate foarte diferite dacă acestea sunt date 2D/3D).

memorie globală => memorie partajată

Syncthreads();

Procesați datele în memoria partajată

Syncthreads();

memorie globală<= разделяемая память

4.3. Activați contoarele de memorie

Indicatorul compilatorului --ptxas-options=-v vă permite să spuneți exact cât și ce fel de memorie (registre, partajată, locală, constantă) utilizați. Dacă compilatorul folosește memoria locală, cu siguranță știi despre asta. Analiza datelor privind cantitatea și tipurile de memorie utilizate vă poate ajuta foarte mult în optimizarea programului.

4.4. Încercați să minimizați utilizarea registrelor și a memoriei partajate

Cu cât nucleul folosește mai multe registre sau memorie partajată, cu atât mai puține fire (sau mai degrabă warps) pot fi executate simultan pe un multiprocesor, deoarece Resursele multiprocesor sunt limitate. Prin urmare, o ușoară creștere a ocupării registrelor sau a memoriei partajate poate duce în unele cazuri la o scădere a performanței la jumătate - tocmai pentru că acum exact jumătate din câte warp-uri sunt executate simultan pe un multiprocesor.

4.5. Memoria partajată în loc de locală

Dacă compilatorul Nvidia din anumite motive a alocat date memoriei locale (de obicei acest lucru se observă printr-o scădere foarte mare a performanței în locurile în care nu este nimic consumator de resurse), aflați exact ce date au ajuns în memoria locală și plasați-le în partajare. memorie).

Adesea, compilatorul plasează o variabilă în memoria locală dacă nu este folosită frecvent. De exemplu, acesta este un fel de acumulator în care acumulezi o valoare în timp ce calculezi ceva într-o buclă. Dacă bucla este mare în ceea ce privește volumul codului (dar nu și timpul de execuție!), atunci compilatorul poate plasa acumulatorul în memoria locală, deoarece este folosit relativ rar, iar registrele sunt puține. Pierderea de performanță în acest caz poate fi vizibilă.

Dacă utilizați foarte rar o variabilă, este mai bine să o plasați în mod explicit în memoria globală.

Deși poate părea convenabil pentru compilator să plaseze automat astfel de variabile în memoria locală, de fapt nu este așa. Va fi dificil să găsiți blocajul în timpul modificărilor ulterioare ale programului dacă variabila începe să fie folosită mai des. Compilatorul poate sau nu transfera o astfel de variabilă în memoria de înregistrare. Dacă modificatorul __global__ este specificat în mod explicit, este mai probabil ca programatorul să-i acorde atenție.

4.6. Bucle de derulare

Derularea buclei este o tehnică de performanță standard în multe sisteme. Esența sa este de a efectua mai multe acțiuni la fiecare iterație, reducând astfel numărul total de iterații și, prin urmare, numărul de ramuri condiționate pe care procesorul va trebui să le efectueze.

Iată cum să derulați bucla pentru găsirea sumei unei matrice (de exemplu, un număr întreg):

int a[N]; int suma;

pentru (int i=0;i

Desigur, buclele pot fi derulate și manual (așa cum se arată mai sus), dar aceasta este o muncă neproductivă. Este mult mai bine să folosiți șabloane C++ în combinație cu funcții inline.

șablon

clasa ArraySumm

Dispozitiv__ static T exec(const T* arr) ( return arr + ArraySumm (arr+1); )

șablon

clasa ArraySumm<0,T>

Dispozitiv__ static T exec(const T* arr) ( returnează 0; )

pentru (int i=0;i

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

Trebuie remarcată o caracteristică interesantă a compilatorului nvcc. Compilatorul va folosi întotdeauna funcții în linie de tip __device__ în mod implicit (există o directivă specială __noinline__ pentru a înlocui acest lucru).

Prin urmare, puteți fi sigur că un exemplu ca cel de mai sus se va desfășura într-o simplă succesiune de declarații și nu va fi în niciun fel inferior ca eficiență codului scris de mână. Cu toate acestea, în cazul general (nu nvcc) nu puteți fi sigur de acest lucru, deoarece inline este doar o indicație pentru compilator pe care o poate ignora. Prin urmare, nu este garantat că funcțiile dvs. vor fi integrate.

4.7. Alinierea datelor și eșantionarea pe 16 octeți

Aliniați structurile de date pe granițele de 16 octeți. În acest caz, compilatorul va putea folosi instrucțiuni speciale pentru ei care încarcă date pe 16 octeți la un moment dat.

Dacă structura este de 8 octeți sau mai puțin, o puteți alinia la 8 octeți. Dar, în acest caz, puteți selecta două variabile simultan combinând două variabile de 8 octeți într-o structură utilizând o unire sau un indicator. Castingul trebuie folosit cu prudență deoarece compilatorul poate plasa datele în memoria locală mai degrabă decât în ​​registre.

4.8. Conflicte de bancă de memorie partajată

Memoria partajată este organizată sub forma a 16 (în total!) bănci de memorie cu un pas de 4 octeți. În timpul execuției unui pool de fire warp pe un multiprocesor, acesta este împărțit în două jumătăți (dacă warp-size = 32) din 16 fire, care accesează la rândul lor memoria partajată.

Sarcinile din diferite jumătăți ale warp nu intră în conflict cu memoria partajată. Datorită faptului că sarcinile unei jumătăți din pool-ul warp vor accesa aceleași bănci de memorie, vor apărea coliziuni și, ca urmare, o scădere a performanței. Sarcinile aflate la o jumătate din warp pot accesa diferite părți ale memoriei partajate într-un anumit pas.

Pașii optimi sunt 4, 12, 28, ..., 2^n-4 octeți (Fig. 8).



Orez. 8. Pași optimi.

Pașii neoptimi sunt 1, 8, 16, 32, ..., 2^n octeți (Fig. 9).



Orez. 9. Pași suboptimali

4.9. Minimizarea mișcărilor datelor gazdei<=>Dispozitiv

Încercați să transferați rezultatele intermediare către gazdă pentru procesare folosind CPU cât mai puțin posibil. Implementați, dacă nu întregul algoritm, atunci cel puțin partea sa principală pe GPU, lăsând doar sarcinile de control la CPU.

5. Biblioteca de matematică portabilă CPU/GPU

Autorul acestui articol a scris o bibliotecă portabilă MGML_MATH pentru lucrul cu obiecte spațiale simple, al cărei cod este operativ atât pe dispozitiv, cât și pe gazdă.

Biblioteca MGML_MATH poate fi folosită ca un cadru pentru scrierea sistemelor portabile CPU/GPU (sau hibride) pentru calcularea problemelor fizice, grafice sau alte probleme spațiale. Principalul său avantaj este că același cod poate fi folosit atât pe CPU, cât și pe GPU și, în același timp, viteza este principala cerință pentru bibliotecă.

6 . Literatură

Chris Kaspersky. Tehnici de optimizare a programelor. Utilizarea eficientă a memoriei. - Sankt Petersburg: BHV-Petersburg, 2003. - 464 p.: ill.

Ghid de programare CUDA 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

Ghid de programare CUDA 1.1. pagina 14-15

Ghid de programare CUDA 1.1. pagina 48

Conform teoriei evoluției lui Darwin, prima maimuță (dacă
mai exact - homo antecessor, predecesor uman) transformat ulterior în
în noi. Centre de calculatoare de mai multe tone cu o mie sau mai multe tuburi radio,
care ocupă încăperi întregi au fost înlocuite cu laptopuri de jumătate de kilogram, care, apropo,
nu va fi inferior ca performanță față de primul. Mașinile de scris antediluviene au devenit
în imprimarea orice și pe orice (chiar și pe corpul uman)
dispozitive multifuncționale. Giganții procesoarelor s-au hotărât brusc să se ridice
miez grafic în „piatră”. Și plăcile video au început nu numai să arate o imagine cu
FPS acceptabil și calitate grafică, dar efectuează și tot felul de calcule. Da
totusi cum sa produci! Tehnologia de calcul multi-threaded folosind GPU-uri va fi discutată.

De ce GPU?

Mă întreb de ce au decis să transfere toată puterea de calcul către grafică
adaptor? După cum puteți vedea, procesoarele sunt încă la modă și este puțin probabil să renunțe la căldură
loc. Dar GPU-ul are câțiva ași în mânecă, împreună cu un joker și câteva mâneci
suficient. Un procesor central modern este conceput pentru a atinge maximum
performanță la procesarea datelor întregi și în virgulă mobilă
virgulă, fără să vă faceți griji în special cu privire la procesarea paralelă a informațiilor. În același timp
timp, arhitectura plăcii video vă permite să „paralelizați” rapid și fără probleme
prelucrarea datelor. Pe de o parte, poligoanele sunt calculate (datorită transportorului 3D),
pe de altă parte, procesarea texturii pixelilor. Este clar că există un „armonios
defalcarea sarcinii din miezul cardului. În plus, performanța memoriei și a procesorului video
mai optim decât combinația „RAM-cache-procesor”. Momentul o unitate de date
în placa video începe să fie procesată de un procesor de flux GPU, altul
unitate este încărcată în paralel în alta și, în principiu, este ușor de realizat
sarcină GPU comparabilă cu lățimea de bandă a magistralei,
oricum, pentru ca acest lucru sa se intample, transportoarele trebuie sa fie incarcate uniform, fara
orice tranziții și ramuri condiționate. Procesorul central, în virtutea acestuia
versatilitatea necesită un cache complet pentru nevoile sale de procesare
informaţii.

Expertii s-au gândit la munca GPU-urilor în calculul paralel și
matematică și a venit cu teoria că multe calcule științifice sunt în multe privințe similare cu
Procesare grafică 3D. Mulți experți consideră că factorul fundamental în
dezvoltare GPGPU (Calcul de uz general pe GPU – universal
calcule folosind o placă video) a fost apariția proiectului Brook GPU în 2003.

Creatorii proiectului de la Universitatea Stanford au trebuit să rezolve o problemă dificilă
problemă: hardware și software pentru a forța adaptorul grafic să producă
calcule diverse. Și au reușit. Folosind limbajul generic C,
Oamenii de știință americani au făcut ca GPU-ul să funcționeze ca un procesor, ajustat pentru
prelucrare paralelă. După Brook, au apărut o serie de proiecte privind calculele VGA,
cum ar fi biblioteca Accelerator, biblioteca Brahma, sistem
Metaprogramarea GPU++ și altele.

CUDA!

Premoniţia perspectivelor de dezvoltare a forţat AMDŞi NVIDIA
se agață de GPU-ul Brook ca un pitbull. Dacă omitem politica de marketing, atunci
Implementând totul corect, puteți obține un punct de sprijin nu numai în sectorul grafic
piață, dar și în informatică (uitați-vă la carduri speciale de calcul și
servere Tesla cu sute de multiprocesoare), înlocuind procesoarele obișnuite.

Desigur, „lorzii FPS” s-au despărțit la piatră de poticnire, fiecare în felul său.
cale, dar principiul de bază a rămas neschimbat - a face calcule
folosind GPU. Și acum vom arunca o privire mai atentă la tehnologia „verde” - CUDA
(Arhitectura dispozitivului unificat de calcul).

Sarcina „eroinei” noastre este să furnizăm un API, două deodată.
Primul este la nivel înalt, CUDA Runtime, care reprezintă funcțiile care
sunt defalcate în niveluri mai simple și trecute la API-ul inferior - Driver CUDA. Aşa
că expresia „nivel înalt” este o întindere de aplicat procesului. Toată sarea este
exact în driver, iar bibliotecile create cu amabilitate vă vor ajuta să îl obțineți
dezvoltatori NVIDIA: CUBLAS (instrumente pentru calcule matematice) și
FFT (calcul folosind algoritmul Fourier). Ei bine, să trecem la practic
părți ale materialului.

Terminologie CUDA

NVIDIA operează cu definiții foarte unice pentru API-ul CUDA. Ei
diferă de definițiile folosite pentru lucrul cu un procesor central.

Fir– un set de date care trebuie prelucrate (nu
necesită resurse mari de procesare).

Urzeală– un grup de 32 de fire. Datele sunt prelucrate numai
warps, prin urmare, warp este cantitatea minimă de date.

Bloc– un set de fluxuri (de la 64 la 512) sau un set
urzeală (de la 2 la 16).

Grilă este o colecție de blocuri. Această diviziune de date
folosit exclusiv pentru a îmbunătăți performanța. Deci, dacă numărul
multiprocesoarele sunt mari, atunci blocurile vor fi executate în paralel. Dacă cu
fără noroc cu cardul (dezvoltatorii recomandă utilizarea
adaptorul nu mai mic decât GeForce 8800 GTS 320 MB), atunci blocurile de date vor fi procesate
secvenţial.

NVIDIA introduce, de asemenea, concepte precum nucleu, gazdă
Şi dispozitiv.

Lucrăm!

Pentru a lucra pe deplin cu CUDA aveți nevoie de:

1. Cunoașteți structura nucleelor ​​GPU shader, încă de la esența programării
constă în distribuirea uniformă a sarcinii între ele.
2. Sa poata programa in mediul C, tinand cont de unele aspecte.

Dezvoltatori NVIDIA a dezvăluit „interiorul” plăcii video de mai multe ori
altfel decât suntem obișnuiți să vedem. Deci, vrând-nevrând, va trebui să studiezi totul
subtilitățile arhitecturii. Să ne uităm la structura legendarei „piatre” G80 GeForce 8800
GTX.

Nucleul shader este format din opt clustere TPC (Texture Processor Cluster).
procesoare de textură (deci, GeForce GTX 280– 15 nuclee, 8800 GTS
sunt șase 8600 – patru etc.). Acestea, la rândul lor, sunt formate din două
multiprocesoare de streaming (denumite în continuare SM). SM (toți
16) constă din front end (rezolvă problemele de citire și decodare a instrucțiunilor) și
conducte back end (ieșirea finală a instrucțiunilor), precum și opt SP scalari (shader
procesor) și două SFU (unități super-funcționale). Pentru fiecare ritm (unitate
timp) front-end-ul selectează warp-ul și îl prelucrează. Pentru ca toată warp să curgă
(să vă reamintesc, sunt 32) procesate, 32/8 = sunt necesare 4 cicluri la capătul transportorului.

Fiecare multiprocesor are ceea ce se numește memorie partajată.
Dimensiunea sa este de 16 kiloocteți și oferă programatorului libertate deplină
actiuni. Distribuiți cum doriți :). Memoria partajată asigură comunicarea între fire
un bloc și nu este destinat să funcționeze cu pixel shaders.

SM-urile pot accesa și GDDR. Pentru a face acest lucru, li s-au dat 8 kilobytes fiecare.
memorie cache care stochează toate cele mai importante lucruri pentru muncă (de exemplu, calcul
constante).

Multiprocesorul are 8192 de registre. Numărul de blocuri active nu poate fi
mai mult de opt, iar numărul de warps nu este mai mare de 768/32 = 24. Din aceasta este clar că G80
poate procesa maximum 32*16*24 = 12288 fire pe unitatea de timp. Nu te poți abține decât
luați în considerare aceste cifre atunci când optimizați programul în viitor (la o scară
– dimensiunea blocului, pe de altă parte – numărul de fire). Echilibrul parametrilor poate juca un rol
rol important în viitor, așadar NVIDIA recomandă utilizarea blocurilor
cu 128 sau 256 fire. Un bloc de 512 fire este ineficient pentru că are
întârzieri crescute. Luând în considerare toate subtilitățile structurii plăcii video GPU plus
bune abilități de programare, puteți crea foarte productiv
instrument pentru calcul paralel. Apropo, despre programare...

Programare

Pentru „creativitate” cu CUDA ai nevoie Placa video GeForce nu mai jos
episodul opt. CU

site-ul oficial trebuie să descărcați trei pachete software: driver de la
Suport CUDA (fiecare sistem de operare are propriul său), pachetul SDK CUDA în sine (al doilea
versiunea beta) și biblioteci suplimentare (kit de instrumente CUDA). Tehnologia suportă
sisteme de operare Windows (XP și Vista), Linux și Mac OS X. A studia I
a ales Vista Ultimate Edition x64 (privind în viitor, voi spune că sistemul s-a comportat
pur și simplu excelent). La momentul scrierii acestor rânduri, era relevant pentru muncă
Driver ForceWare 177.35. Folosit ca un set de instrumente
Pachetul software Borland C++ 6 Builder (deși orice mediu care funcționează cu
limba C).

Va fi ușor pentru o persoană care cunoaște limba să se obișnuiască cu un mediu nou. Tot ceea ce este necesar este
amintiți-vă parametrii de bază. Cuvânt cheie _global_ (plasat înaintea funcției)
indică faptul că funcția aparține nucleului. Ea va fi chemată de centrală
procesor și toată munca se va întâmpla pe GPU. Apelul _global_ necesită mai mult
detalii specifice, și anume dimensiunea rețelei, dimensiunea blocului și ce va fi nucleul
aplicat. De exemplu, linia _global_ void saxpy_parallel<<>>, unde X –
dimensiunea grilei, iar Y este dimensiunea blocului, specifică acești parametri.

Simbolul _device_ înseamnă că funcția va fi apelată de nucleul grafic, cunoscut și ca
va urma toate instrucțiunile. Această funcție este localizată în memoria multiprocesorului,
prin urmare, este imposibil să obții adresa ei. Prefixul _gazdă_ înseamnă că apelul
iar procesarea va avea loc numai cu participarea CPU. Trebuie avut în vedere faptul că _global_ și
_dispozitivele_ nu se pot suna între ele și nu se pot suna singure.

De asemenea, limbajul pentru CUDA are o serie de funcții pentru lucrul cu memoria video: cudafree
(eliberarea memoriei între GDDR și RAM), cudamemcpy și cudamemcpy2D (copiere
memorie între GDDR și RAM) și cudamalloc (alocarea memoriei).

Toate codurile de program sunt compilate de API-ul CUDA. Mai întâi se ia
cod destinat exclusiv procesorului central și supus
compilare standard și alt cod destinat adaptorului grafic,
rescris în limbajul intermediar PTX (la fel ca assembler) pentru
identificarea eventualelor erori. După toate aceste „dansuri” finala
traducerea (traducerea) comenzilor într-un limbaj ușor de înțeles pentru GPU/CPU.

Trusa de studiu

Aproape toate aspectele programării sunt descrise în documentația care urmează
împreună cu driverul și două aplicații, precum și pe site-ul dezvoltatorilor. Dimensiune
articolul nu este suficient pentru a le descrie (cititorul interesat ar trebui să atașeze
puțin efort și studiați singur materialul).

CUDA SDK Browser a fost dezvoltat special pentru începători. Oricine poate
simțiți direct puterea calculului paralel (cel mai bun test pentru
stabilitate – exemplele funcționează fără artefacte sau accidente). Aplicația are
un număr mare de mini-programe orientative (61 de „teste”). Pentru fiecare experiență există
Documentație detaliată a codului programului plus fișiere PDF. Este imediat evident că oamenii
cei prezenți cu creațiile lor în browser fac treabă serioasă.
De asemenea, puteți compara viteza procesorului și a plăcii video în timpul procesării
date. De exemplu, scanarea matricelor multidimensionale cu o placă video GeForce 8800
GT Produce 512 MB cu un bloc cu 256 fire în 0,17109 milisecunde.
Tehnologia nu recunoaște tandemurile SLI, așa că dacă aveți un duo sau trio,
dezactivați funcția de „împerechere” înainte de a lucra, altfel CUDA va vedea doar una
dispozitiv Dual core AMD Athlon 64 X2(frecvența de bază 3000 MHz) aceeași experiență
trece în 2,761528 milisecunde. Se pare că G92 este de peste 16 ori
mai repede decât o stâncă AMD! După cum puteți vedea, acesta este departe de a fi un sistem extrem
în tandem cu un sistem de operare neiubit de mase arată bine
rezultate.

Pe lângă browser, există o serie de programe utile societății. Chirpici
și-a adaptat produsele la noile tehnologii. Acum Photoshop CS4 este complet
folosește cel puțin resursele adaptoarelor grafice (trebuie să descărcați o versiune specială
plugin). Cu programe precum Badaboom media converter și RapiHD puteți
decodați video în format MPEG-2. Bun pentru procesarea audio
Utilitarul gratuit Accelero este potrivit. Cantitatea de software adaptată pentru API-ul CUDA,
va crește fără îndoială.

Și în acest moment...

Între timp, citiți acest material, muncitori din problemele de procesor
își dezvoltă propriile tehnologii pentru a integra GPU-uri în procesoare. Din exterior AMD Toate
este clar: au o experiență extraordinară dobândită împreună cu ATI.

Crearea de „microdispozitive”, Fusion, va consta din mai multe nuclee sub
Nume de cod Buldozer și cip video RV710 (Kong). Relația lor va fi
realizat prin intermediul autobuzului HyperTransport îmbunătățit. În funcție de
numărul de nuclee și caracteristicile de frecvență ale acestora AMD intenționează să creeze un preț întreg
ierarhia „pietrelor”. De asemenea, este planificată să producă procesoare pentru laptopuri (Falcon),
și pentru gadgeturi multimedia (Bobcat). Mai mult, este aplicarea tehnologiei
în dispozitive portabile va fi provocarea inițială pentru canadieni. Cu dezvoltarea
calculul paralel, utilizarea unor astfel de „pietre” ar trebui să fie foarte populară.

Intel puțin în urmă cu Larrabee-ul său. Produse AMD,
dacă nu se întâmplă nimic, acestea vor apărea pe rafturile magazinelor la sfârșitul anului 2009 - devreme
2010. Iar decizia inamicului va ieși la iveală doar în aproape două
an.

Larrabee va avea un număr mare (a se citi: sute) de nuclee. La început
Vor exista și produse concepute pentru 8 – 64 de nuclee. Sunt foarte asemănătoare cu Pentium, dar
destul de puternic refăcut. Fiecare nucleu are 256 kilobytes de cache L2
(dimensiunea sa va crește în timp). Relația se va realiza prin
Bus inel bidirecțional de 1024 de biți. Intel spune că „copilul” lor va fi
funcționează perfect cu DirectX și Open GL API (pentru dezvoltatorii Apple), deci nu
nu este necesară nicio intervenție software.

De ce ți-am spus toate astea? Este evident că Larrabee și Fusion nu se vor deplasa
procesoare obișnuite, staționare de pe piață, așa cum nu vor fi forțați să iasă de pe piață
plăci video. Pentru jucătorii și pasionații de sporturi extreme, visul suprem va rămâne în continuare
CPU multi-core și un tandem de mai multe VGA-uri de top. Dar ce chiar
companiile de procesoare trec la calculul paralel pe baza principiilor
similar cu GPGPU, spune multe. În special, despre ce așa
tehnologie precum CUDA are dreptul de a exista și, cel mai probabil, va exista
foarte popular.

Un scurt rezumat

Calculul paralel folosind o placă video este doar un instrument bun
în mâinile unui programator harnic. Cu greu pentru procesoarele conduse de legea lui Moore
va veni sfarsitul. Companii NVIDIA mai este mult de parcurs
promovarea API-ului său în masă (același lucru se poate spune despre creație ATI/AMD).
Cum va fi, viitorul va arăta. Deci CUDA se va întoarce :).

P.S. Recomand să viziteze programatorii începători și persoanele interesate
următoarele „unități virtuale”:

Site-ul web și site-ul oficial NVIDIA
GPGPU.com. Toate
informațiile furnizate sunt în engleză, dar măcar vă mulțumesc că nu sunt în
chinez Așa că du-te! Sper că autorul te-a ajutat măcar puțin
călătorii interesante în explorarea CUDA!