Onion Information

Page Clicks: 0

First Seen: 03/15/2024

Last Indexed: 09/18/2024

Domain Index Total: 397

Python modul PyCUDA: mjerenje performansi, profiliranje i optimizacija - Mjerenje performansi zrna - Da bi mjerili brzinu izvođenja zrna i operacija kopiranja memorije s domaćina na uređaj i obrnuto, koristimo profiler (engl. profiler ). PyCUDA podržava profiliranje , i ono se uključuje postavljanjem varijable okoline CUDA_PROFILE na vrijednost 1 , što možemo učiniti na način: Uzmemo li da je program.py sada kod koji zbraja dva vektora, datoteka cuda_profile_0.log koju dobivamo kao izlaz je sadržaja: Za usporedbu, ista datoteka na drugom domaćinu i drugom uređaju je sadržaja: Vidimo da je zbog jačeg grafičkog procesora vrijeme izvođenja pojedinih funkcije manje, tako da bi najjednostavnija metoda optimizacije bila "pokrenite kod na jačem grafičkom procesoru". Zadatak - Promijenite kod da zbraja dva vektora veličine 1600 elemenata i usporedite vrijeme izvođenja kad: Napišite kod koji množi matricu veličine (2048, 2084) s vektorom veličine 2048: Koja varijanta daje najbolje performanse? Vizualni profiler - Naredbom nvvp pokrećemo NVIDIA Visual Profiler. On je namijenjen za profiliranje CUDA C/C++ aplikacija, ali može profilirati i PyCUDA aplikacije na idući način: pod File/New Session podaci su - Environment: > Add - Zatim, Next > i na idućoj stranici Finish . Note - Ukoliko radite na udaljenom računalu, nvvp vam neće moći na ispravan način prikazati pomoć; iskoristite kopiju iste dokumentacije dostupnu putem weba . Optimizacija - GPU nudi najbolje performanse kada je potpuno zasićen, odnosno kada sve njegove jezgre računaju. Optimizacija se uglavnom vrši kombinacijom tri načina: Maksimizacija paralelizacije - Maksimizacija paralelizacije uključuje tri faktora: Korištenje svih jezgara zahtijeva veliki broj blokova (barem onoliko koliko ima multiprocessora na GPU uređaju). Na uređajima koji imaju računsku sposobnost 2.0+ (mi ih koristimo), moguće je istovremeno pokrenuti dva zrna. Ako imate potreba za serializacijom, odnosno procesne niti ovise međusobno o rezultatima izračuna, pokušajte zadržati međuovisne procesne niti na istom bloku: za to možete koristiti __syncthreads() umjesto odvojenih pokretanja zrna. Skrivanje kašnjenja eksploatira paralelnost između komponenata GPU-a: ako postoje računske osnove koje mogu biti otpremljene dok druge računske osnove čekaju, kašnjenje je time pokriveno. Skrivanje kašnjenja zahtijeva poznavanje otpreme instrukcija, što ovisi od računskoj sposobnosti uređaja. Uređaji s računskom sposobnošću 2.0 skrivaju kašnjenje na slijedeći način: Tipična instrukcija zahtijeva 22 instrukcijska ciklusa da se izvrši. To implicira da je potrebno barem 22 računske osnove (za sposobnost 2.x) da bi sakrili standardno sekvencijalno izvođenje (instrukcija ovisi o rezultatu izvođenja prethodne instrukcije). Pristup globalnoj i lokalnoj memoriji traje od 400 do 800 instrukcijskih ciklusa. Računske osnove potrebne da bi sakrile to kašnjenje ovise o gustoći između instrukcije koje ne ovise o pristupu memoriji i instrukcija koje ovise o tome. Kada se koristi sinkronizacija, sve računske osnove u bloku zastaju i njihovo kašnjenje je jednako kašnjenju procesne niti koja se najduže izvodi. To se pokriva pokretanjem više od jednog bloka po multiprocesoru. Zauzeće, koje se dobiva kao omjer rezidentnih računskih osnova i maksimalnog broja rezidentnih računskih osnova, može se izračunati korištenjem proračunske tablice CUDA Occupancy Calculator . Veće zauzeće pomaže u skrivanju kašnjenja zbog sinkronizacije. Minimizacija kašnjenja memorije - Pristup memoriji vrši se u transkacijama. Jedna transkacija pristupa 32, 64 ili 128 bajta, poravnato njihovoj veličini (primjerice, 128-bajtne transakcije mogu biti poravnate na 128-bajtnu granicu). Instrukcije pristupa memoriji automatski se slažu u transakcije. Broj transakcija potreban za pristup memoriji ovisi o: Jedna memorijska instrukcija se šalje za pristup broju bajtova koji je potencija broja 2 ( \(2^1, 2^2, \dots, 2^{16}\) ), koji su prirodno poravnati. Svi ostali pristupi dijele se u više od jedne instrukcije. Funkcija cudaMalloc() vraća memoriju koja je poravnata na 256 bajta. Poravnanje redaka za linearizirana 2D polja može se poboljšati korištenjem funkcije cudaMallocPitch() : Varijabla pitch sadržava veličinu danu u bajtovima. Elementu (redak, stupac) se tada može pristupiti s ((data_t*)((char*)devData + row*pitch))[col] . Kopiranje tako alocirane memorije vrši se pomoću funkcije cudaMemcpy2D() : gdje su destpitch , srcpitch i row_width dane u bajtovima i redom su: Pravila udruživanja - Pravila udruživanja za Fermi (računska sposobnost 2.x) su: Dijeljena memorija i konflikti banke - Dijenjena memorija sastoji se od memorijskih banaka. 32-bitne (4-bajtne) riječi su slijedno dodijeljene bankama. Širina pojasa je 32 bita kroz dva ciklusa sata. Dijeljena memorija je tipično jednako brza kao registri. Konflikti banke (kada cuda_runtime_api.h dvije ili više niti istovremeno pristupa istoj banci) uvode kašnjenje. Svaka polovica računske osnove vrši zahtijeve za memorijom neovisno. Procesne niti u prvom dijelu računske osnove ne mogu imati konflikte banke s procesnim nitima u drugom dijelu. Ako sve procesne niti pristupaju istoj riječi, podaci se predaju svima na broadcast način i nema konflikata banke. Zajednički pristup vrši se u koracima veličine 32 bita ( float , float2 , float3 , float4 itd.). Konflikti banke događaju se ako je korak paran (primjerice, float2 ili float4 ), a u slučaju da je neparan nema konflikata banke (primjerice, float , float3 ). Kod 8-bitnog and 16-bitnog pristupa memoriji, svi pristupi uzrokuju konflikte banke osim u slučaju kada su podaci isprepleteni (engl. interleaved ). Kod 64 ili više-bitnog pristupa konflikte banke nije moguće izbjeći. Računska sposobnost 2.x uvodi nekoliko promjena. Maksimizacija propusnosti instrukcija - Maksimizacija propusnosti instrukcija ima tri faktora: Uređaji koji imaju računsku sposobnost 1.x imaju mogućnost brzog množenja 24-bitnih cijelih brojeva. __mul24() može se koristiti umjesto standardnog produkta kada množenje ne izaziva aritmetički preljev. Uređaji koji imaju računsku sposobnost 2.x imaju mogućnost brzog množenja 32-bitnih cijelih brojeva, i __mul24() je sporiji od od standardnog produkta. Određene zastavice program prevoditelja upravljaju korištenjem funkcija kao što su logaritamska, eksponencijalna i trigonometrijske funkcije. ( --use_fast_math i bool parametri --ftz= , --prec-div= , --prec-sqrt= mijenjaju način baratanja denormalnim brojevima te preciznost dijeljenja i kvadratnog korjenovanja). Preporuke za optimizaciju koda - Sumarno, preporuke za optimizaciju koda su: Author: Vedran Miletić