U prethodnom blogu započela sam priču o NVIDIA Parallel Nsight, softveru koji omogućava bolju integraciju CUDA arhitekture u Microsoft Visual Studio i koji znatno olakšava razvoj i analizu efikasnosti CUDA aplikacija. Parallel Nsight omogućava i debug-ovanje CUDA kernel funkcija, što bez ovog softvera uopšte nije bilo moguće u Windows okruženju. U ovom tekstu ću na jednom primeru da pokažem kako se pomoću Parallel Nsight-a debug-uje CUDA kernel funkcija.

Debug-ovanje CUDA programa

Ne znam da li ste nekada pokušali da debug-ujete neki CUDA program. Ako jeste, onda ste svakako primetili da ne možete da vidite vrednosti promenljivih koje su definisane u device memoriji, kao i da ne možete da pristupite kernel funkcijama sa Step Into opcijom – za klasičan debugger kernel funkcije predstavljaju crne kutije. Uz pomoć Parallel Nsight-a postaje moguće debug-ovati kernel funkcije, pratiti tok njihovog izvršavanja, pratiti promene vrednosti promenljivih. Za prikaz procesa debug-ovanja CUDA programa, koristiću već objašnjeni primer iz Bloga br. 3.

Kada se kompajlira source code i kreira exe verzija, treba pokrenuti Nsight Monitor, da bi uopšte bilo moguće pokrenuti Nsight Debugger i postaviti breakpoint unutar tela kernel funkcije. Breakpoint se može postaviti ili klikom miša pored željene linije ili pritiskom na taster F9 – oba načina se koriste i u standardnim C/C++ programima. Nakon toga sa linije menija Visual Studio okruženja odabrati opciju Nsight -> Start CUDA Debugging.

Parallel Nsight se preko Nsight Monitor-a konektuje na target mašinu (takva je procedura bez obzira da li su host i target mašina odvojeni računari ili su oba GPU uređaja u jednom računaru) i počinje sa izvršavanjem programa. Kada izvršavanje dođe do linije u kojoj je postavljen breakpoint, zaustavlja se rad programa i dobro poznatom žutom strelicom obeležava se linija do koje je program stigao sa izvršavanjem. Izborom opcije Step Over ili pritiskom tastera F10 prelazi se na sledeću liniju koda. Sve u svemu, sve prečice koje važe za debug-ovanje klasičnih programa u VS okruženju, važe i ovde.

Moguće je pratiti promene vrednosti promenljivih deklarisanih i alociranih u device memoriji. Na sledećoj slici je prikazano kako se u Watch prozoru mogu pratiti vrednosti članova niza d_a, koji se menjaju unutar kernel funkcije. Pošto je ovaj primer već objašnjavan u Blogu br. 3, neću se ponovo baviti objašnjavanjem, samo ću podsetiti da kernel funkcija u ovom primeru treba svim članovima niza da dodeli vrednost indeksa tog elementa. Dakle, vrednosti članova niza su 0,1,2,3,4…, što se vidi i u Watch prozoru, koji je prikazan na sledećoj slici:

Takođe, u Memory prozoru se mogu pratiti promene vrednosti koje su zapamćene na određenim lokacijama u memoriji. Ako se prilikom debug-ovanja u kodu označi promenljiva d_a, zatim se prevuče u prozor Memory 1, pojaviće se vrednosti koje su zapamćene na toj adresi u memoriji. Desnim klikom unutar ovog prozora i selektovanjem tipa promenljive d_a (u ovom slučaju, to je 4-byte Integer), prikazuju se vrednosti kao na slici:

Napomena: ove slike ilustruju vrednosti koje članovi niza imaju u određenom trenutku izvršavanja (u ovom slučaju to je sam kraj kernel funkcije). Izvršavanje kernel funkcije je prekinuto kada neka od niti (nije striktno definisano ni tačno koja nit, ni kom bloku pripada) dostigne breakpoint. Postoji način da se odabere i tačan blok niti u okviru koga treba pratiti izvršavanje i zaustaviti se na breakpoint-u. Takođe, postoji mogućnost postavljanja uslovnog breakpoint-a. Ovo se radi uobičajenim postupkom – desni klik na breakpoint, selektovanje opcije Condition i upisivanje željenog uslova. Treba voditi računa da se u uslovu koriste samo promenljive definisane u device memoriji, pošto su samo one “vidljive” u okviru kernel funkcija. Kada je definisan uslovni breakpoint, pored te linije stoji crveni krug, sa znakom plus:

U ovom primeru definisano je da treba prekinuti izvršavanje programa kada nit sa indeksom 8 dostigne obeleženu 12. liniju koda. U prozoru Breakpoints na listi se nalazi ovaj uslovni breakpoint koji je definisan za kernel funkciju:

Parallel Nsight ima implementiranu još jednu jako korisnu funkcionalnost – proveru memorijskog opsega (out-of-bounds memory checking). Ukoliko bi se malo izmenila kernel funkcija iz primera 3, tako da se linija sa proverom opsega stavi pod komentar, kernel bi izgledao ovako:

__global__ void myFirstKernel(int N, int *d_a)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
//    if (idx<N)
        d_a[idx] = idx;  
}

Ako se ponovo kompajlira i pokrene ovako izmenjeni program, moguće je da bi CUDA prilikom izvršavanja prijavila grešku. Možda se sećate da sam upravo u tom blogu br. 3, ali i u blogu o otkrivanju grešaka objašnjavala da bi bez ovakve provere opsega neke niti pristupale nedefinisanom memorijskom prostoru. U tom slučaju, došlo bi do neke nedefinisane reakcije – ili bi program jednostavno nastavio dalje sa radom, upisujući podatke u neki nedefinisani prostor u memoriji, ili bi prijavio grešku. I pritom reakcija ne bi bila ista prilikom svakog pokretanja… Ako je program jednostavniji, onda i nije veliki problem da se uz malo iskustva otkrije uzrok ovakvog ponašanja, ali ako je program kompleksniji, a pritom su kernel funkcije kao crna kutija, jako je teško otkriti šta je problem. Uz pomoć Parallel Nsight-a vrlo se lako prilikom debug-ovanja otkriva da kernel pokušava da pristupi nedefinisanoj memorijskoj lokaciji. U ovom konkretnom slučaju, ako se pokrene debug sa izmenjenom verzijom programa, izvršavanje se zaustavlja na definisanom breakpoint-u:

A kada se pokuša prelazak na sledeću liniju, pojavljuje se pop-up prozor u dnu ekrana, sa informacijom da je došlo do greške u pristupu memoriji:

Ovakve informacije su od velike koristi kada je potrebno tražiti greške, pogotovo ako se radi o kompleksnim programima sa desetinama kernel funkcija integrisanim u postojeće C++ projekte.

…

Predstavljanje Parallel Nsight alata za Visual Studio je nastavljeno u ovom tekstu. Pokazala sam kako  se debug-uju CUDA programi i koliko je jednostavnije otkrivati greške na ovaj način. NVIDIA je razvojem ovog softvera značajno pomogla programerima koji koriste Microsoft-ovo razvojno okruženje, a žele da se bave CUDA programiranjem. Sigurno je da će Nsight nastaviti da se razvija, tako da ću se možda u budućnosti ponovo vratiti na ovu temu, da predstavim neke novitete… U svakom slučaju, u sledećem blogu pisaću o još jednom alata za analizu efikasnosti CUDA programa, koji takođe stiže iz “kuhinje” NVIDIA korporacije.

Do skorog čitanja, T.Đ.