nvidia cuda tehnologija. NVidia CUDA: GPU se puni i CPU umire? Host uređaj minimizacije migracije podataka

Taj ínshi. Međutim, traženje kombinacije "CUDA skeniranje", nakon što sam vidio samo 2 članka, ni na koji način ne odgovara algoritmu skeniranja na GPU-u - ali to je jedan od samih osnovnih algoritama. Dakle, duboko udahnuvši, osvrnimo se na tečaj Udacity - Uvod u paralelno programiranje i usuđujem se napisati novu seriju članaka o CUDA-i. Reći ću vam još jednom da će serija biti utemeljena na svom tijeku, a čak i ako imate sat vremena, proći ćete jogu bogatije. Trenutno su u planu sljedeći članci:
1. dio: Odredba.
Dio 2: GPU hardver i obrasci paralelne komunikacije.
Dio 3: Temeljni GPU algoritmi: smanjenje, skeniranje i histogram.
Dio 4: Temeljni GPU algoritmi: kompaktno, segmentirano skeniranje, sortiranje. Praktična primjena nekih algoritama.
Dio 5: Optimizacija GPU softvera.
Dio 6: Primjer paralelizacije sljedećih algoritama.
Dio 7: Dodatne teme paralelnog programiranja, dinamički paralelizam.

Zatrimka vs građevinska dozvola

Prije svega, možete li staviti skin ispred zamrzavanja GPU-a da dovršite svoje zadatke - a za takve svrhe, dobar GPU, ako ga trebate zamrznuti? Za vidpovídí síd označavaju 2 pojma:
Zatrimka(latencija) - sat tijekom kojeg je jedna instrukcija/operacija završena.
Kapacitet protoka- broj instrukcija/operacija koje se broje po satu.
Jednostavan primjer: možda osobni automobil brzine 90 km/godišnje i kapaciteta 4 osobe, te autobus kapaciteta 60 km/godišnje i kapaciteta 20 osoba. Ako operacija prihvaća kretanje 1 osobe po 1 kilometru, tada je kašnjenje osobnog automobila 3600/90 = 40 s - za nekoliko sekundi 1 osoba je na pola puta do 1 kilometra, propusnost automobila je 4/ 40 = 0,1 operacija / sekundi; kašnjenje sabirnice - 3600/60=60s, propusnost sabirnice - 20/60=0,3(3) operacije/sekundi.
Dakle, od CPU - tse car, GPU - bus: može doći do velike blokade, ali također i do velike propusne moći. Što se tiče vremena određene operacije skina, ono nije toliko važno koliko broj tih operacija u sekundi - samo pogledajte GPU.

Osnovno razumijevanje i pojmovi CUDA

• Privitak (uređaj)- GPU. Uloga "ispadnika" je pobjednička - manje je vjerojatno da će raditi oni koji uopće mogu koristiti CPU.
• domaćin (domaćin)- CPU. Vikonu igra ulogu - pokretanje zadatka na aneksu, gledanje memorije na aneksu, premještanje memorije / iz dodatka. I tako, CUDA pjevanje to kaže kao prilog, tako da se domaćin može sjetiti svojeg sjećanja.
• Zrno- Zadatak koji pokreće domaćin na gospodarskoj zgradi.

Glavne faze CUDA programa

1. Domaćin vidi potrebnu količinu memorije na aneksu.
2. Host će kopirati podatke iz moje memorije u moju memoriju.
3. Domaćin započinje pjevanje kernela na gospodarskoj zgradi.
4. Pričvrstite jezgre.
5. Priložit ću kopiju rezultata iz memorije u svoju memoriju.

Jezgre

1. Srednje jezgre mogu prepoznati "identifikator" ili, jednostavnije, naizgled, položaj toka, koji je zarazno pobjednički - pobjedničkim položajem postižemo da je jedna jezgra izvediva s različitim traperom ugara u toku, u koji radi. Prije govora, takva organizacija paralelnih izračuna naziva se SIMD (Single Instruction Multiple Data) - ako nekoliko procesora dovršava istu operaciju u isto vrijeme, iako na različitim podacima.
2. U nekim slučajevima kodovi jezgre trebaju promijeniti različite načine sinkronizacije.

• Na poleđini su postavljene dimenzije tzv. grida (mreže), u 3D koordinatama: mreža_x, mreža_y, mreža_z. Kao rezultat toga, mreža se formira od mreža_x*mreža_y*mreža_z blokovi.
• Zatim se bloku mijenja veličina u 3D koordinatama: blok_x, blok_y, blok_z. Kao rezultat toga, blok će biti presavijen blok_x*blok_y*blok_z teče. Oče, molim te grid_x*grid_y*grid_z*block_x*block_y*block_z teče. Važno poštovanje - maksimalan broj protoka u jednom bloku je obrubljen i pohranjen po GPU modelu - tipične vrijednosti su 512 (stariji modeli) i 1024 (noviji modeli).
• Srednja jezgra dostupna promjena threadIdxі blockIdx s poljima x, y, z- smrad za osvetu 3D koordinate toka u bloku i bloku na mjestu je jasno. Promjene su također dostupne blokDimі gridDim s istim poljima - promijenite veličinu bloka i mreža je jasna.

Pišemo program na CUDA-i

glavni.cpp

#uključi #uključi #uključi #uključi #uključi #uključi #uključi #uključi #include "openMP.hpp" #include "CUDA_wrappers.hpp" #include "common/image_helpers.hpp" using namespace cv; korištenje imenskog prostora std; int main(int argc, char** argv) ( koristeći prostor imena std::chrono; if(argc != 2) ( cout<<" Usage: convert_to_grayscale imagefile" << endl; return -1; } Mat image, imageGray; uchar4 *imageArray; unsigned char *imageGrayArray; prepareImagePointers(argv, image, &imageArray, imageGray, &imageGrayArray, CV_8UC1); int numRows = image.rows, numCols = image.cols; auto start = system_clock::now(); RGBtoGrayscaleOpenMP(imageArray, imageGrayArray, numRows, numCols); auto duration = duration_cast(system_clock::now() - početak); cout<<"OpenMP time (ms):" << duration.count() << endl; memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols); RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols); return 0; }

pripremitiImagePointers

šablona void readyImagePointers(const char * const inputImageFileName, cv::Mat& inputImage, T1** inputImageArray, cv::Mat& outputImage, T2** outputImageArray, const int outputImageType) ( using namespace std; using namespace cv; inputImage = imread IMREAD_COLOR), if (inputImage.empty()) (cerr<< "Couldn"t open input file." << endl; exit(1); } //allocate memory for the output outputImage.create(inputImage.rows, inputImage.cols, outputImageType); cvtColor(inputImage, inputImage, cv::COLOR_BGR2BGRA); *inputImageArray = (T1*)inputImage.ptr(0); *izlazni niz slika = (T2*) izlazna slika.ptr (0); }

openMP.hpp

#uključi #uključi #uključi void RGBtoGrayscaleOpenMP(uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) ( #pragma omp parallel for collapse(2) for (int i = 0; i< numRows; ++i) { for (int j = 0; j < numCols; ++j) { const uchar4 pixel = imageArray; imageGrayArray = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } } }

Tekst priloga

void RGBtoGrayscaleCUDA(const uchar4 * const h_imageRGBA, unsigned char* const h_imageGray, size_t numRows, size_t numCols) ( uchar4 *d_imageRGBA; unsigned char *d_imageGray; const size_t num // izbrisana memorija na dodatku za provjeru pokretanjaCudaErrors(cudaMalgeRG( &d_imageRG) , sizeof(uchar4) * numPixels)); provjera GPU-aCudaErrors(cudaMemcpy(d_imageRGBA, h_imageRGBA, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));

Tekst priloga

#uključi #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__) predložak void check (T err, const char * const func, const char * const file, const int line) ( if (err! = cudaSuccess) ( std::cerr<< "CUDA error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString(err) << " " << func << std::endl; exit(1); } }

Tekst priloga

dim3 veličina bloka; dim3 veličina mreže; int NitNum; cudaEvent_t početak, zaustavljanje; cudaEventCreate(&start); cudaEventCreate(&stop); NitNum = 1024; Veličina bloka = dim3(Brojniti, 1, 1); veličina rešetke = dim3(brojCols/threadNum+1, numRows, 1); cudaEventRecord(start); rgba_u_sivim tonovima_jednostavno<<>>(d_imageRGBA, d_imageGray, numRows, numCols); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); provjeriCudaErrors(cudaGetLastError()); float milisekunde = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::out<< "CUDA time simple (ms): " << milliseconds << std::endl;

rgba_u_sivim tonovima_jednostavno

Global__ void rgba_to_grayscale_simple(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols) ( int y = blockDim.y*blockIdx.y + threadIdx.x; .x, if (x>=numCols || y >=numRows) return, const int offset = y*numCols+x; const uchar4 pixel = d_imageRGBA; 0.114f*pixel.z;

Tekst priloga

checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost)); cudaFree(d_imageGray); cudaFree(d_imageRGBA);

Tekst priloga

Procesor: Intel Core(TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4GB, 1600 MHz.
OS: OS X 10.9.5.
Prevodilac: g++ (GCC) 4.9.2 20141029.
CUDA kompajler: Cuda alati za kompilaciju, izdanje 6.0, V6.0.1.
Podržana verzija OpenMP-a: OpenMP 4.0.

rgba_to_grayscale_optimized

#define WARP_SIZE 32 __global__ void rgba_to_grayscale_optimized(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols, int elemsPerThread) ( int yx; x*blockIdx.x + x threadIdx.x; const int start = (petlja/_ WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x;for (int i=loop_start, j=0; j

Tekst priloga

NitNum=128; const int elemsPerThread = 16; Veličina bloka = dim3(Brojniti, 1, 1); gridSize = dim3(numCols/(threadNum*elemsPerThread) + 1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_optimized<<>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); provjeriCudaErrors(cudaGetLastError()); milisekunde = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::out<< "CUDA time optimized (ms): " << milliseconds << std::endl;

Oznake: Dodajte oznake

- Skup sučelja s niskim programiranjem ( API) za stvaranje igora i drugih visoko produktivnih multimedijskih dodataka. Uključite visokoproduktivno podrezivanje 2D- І 3D-grafika, zvuk i uvod.

Direct3D (D3D) - sučelje za gledanje trivi svjetova primitivci(Geometrijska tijela). Unesi .

OpenGL(Vid engleski. Otvori grafičku biblioteku, Doslovno - otvorena grafička biblioteka) - specifikacija koja definira neovisnu vrstu višeplatformskog programabilnog softverskog sučelja za pisanje dodataka, koji koriste dvodimenzionalnu i trodimenzionalnu računalnu grafiku. Uključuje više od 250 funkcija za slikanje sklopivih scena iz tri svijeta iz jednostavnih primitiva. Pobjeda za stvaranje video planina, virtualna stvarnost, vizualizacija u znanstvenom istraživanju. Na peronu Windows natjecati se sa .

OpenCL(Vid engleski. Otvoreni računalni jezik, Doslovno - vodkrita mov izračun) - okvir(Okvir programskog sustava) za pisanje računalnih programa povezanih s paralelnim nabojima na raznim grafikama ( GPU) i ( ). U okvir OpenCL unesite dodatke za programiranje mov i programiranje sučelja ( API). OpenCL osigurati paralelizam na razini instrukcija i na razini podataka te na implementaciji tehnologije GPGPU.

GPGPU(kratki engleski) Grafičke jedinice opće namjene, doslovno - GPU globalno prepoznavanje) - tehnika korištenja grafičkog procesora video kartice za globalne izračune, kako to zvuči.

shader(engl. shader) - program za stvaranje sjena na sintetiziranim slikama, za isticanje trivijalne grafike preostalih parametara objekta slike. U pravilu uključuje određeni stupanj preklapanja, opis gline i ružičastog svjetla, prekrivanje teksture, presavijeno i lomljeno, sjenčanje, pomicanje površine i efekte naknadne obrade. Preklopne površine mogu se vizualizirati uz pomoć jednostavnih geometrijskih oblika.

prikazivanje(engl. prikazivanje) - vizualizacija, u računalnoj grafici, proces uzimanja slike za model uz pomoć softvera.

SDK(kratki engleski) Komplet za razvoj softvera) - Skup alata za razvoj softverske sigurnosti.

CPU(kratki engleski) Središnja procesorska jedinica, Doslovno - središnji / glavni / dodatak za brojanje glava) - središnji (mikro); prilog, koji vikonu strojne upute; Dio hardverske sigurnosti koji je zadužen za proračun operacija (postavke od strane operativnog sustava i aplikacijskog softvera) i koordinaciju rada svih priključaka.

GPU(kratki engleski) Jedinica za grafičku obradu, Doslovno - grafički računski dodatak) - grafički procesor; okremy privitke ili igraće konzole, koje vikonu grafički prikaz (vizualizacija). Moderni grafički procesori mogu učinkovito obraditi i realistično prikazati računalnu grafiku. Grafički procesor u današnjim video adapterima zastosovuetsya kao pričvršćivač trivivirnoy grafike, proteo yogo može se uvijati u nekim načinima i za izračun ( GPGPU).

problema CPU

Dugo vremena povećanje produktivnosti tradicionalnih bilo je najvažnije zbog naknadnog povećanja frekvencije takta (oko 80% produktivnosti samog početka frekvencije sata) s jednosatnim povećanjem broja tranzistora u jednom kristalu. Međutim, dalji pomak taktne frekvencije (na taktnoj frekvenciji većoj od 3,8 GHz, čipovi se jednostavno pregrijavaju!) nailazi na niz temeljnih fizičkih prepreka (krhotine tehnološkog procesa mogu biti vrlo blizu veličine atoma : , A veličina atoma silicija je približno 0,543 nm):

Prvo, da biste promijenili ekspanziju kristala i povećali frekvenciju takta, povećajte protok tranzistora. Tse vede do povećanja smanjene napetosti i povećanja gubitka topline;

Na drugi način, povećanje najviše taktne frekvencije često se pokreće kroz zastoje tijekom prijelaza na memoriju, budući da pristup memoriji ne odgovara rastućim taktnim frekvencijama;

Treće, za neke dodatke tradicionalnim uzastopnim arhitekturama, oni postaju neučinkoviti zbog povećanja frekvencije takta kroz takozvanu "von Neumannovu srednju školu" - smanjenje produktivnosti kao rezultat uzastopnog tijeka izračuna. Kod kojih postoji otporno-kapacitivno ometanje prijenosa signala, bilo s dodatnim uskim prostorom, bilo s povećanjem frekvencije takta.

Rosvitok GPU

U isto vrijeme, razvoj GPU:

…

Opadanje lišća 2008 - Intel uveo liniju od 4 jezgre Intel Core i7, koji se temelje na mikroarhitekturi nove generacije Nehalem. Procesori rade na radnom taktu od 26-32 GHz. Vikonani iza 45-nm procesne tehnologije.

Grudi 2008 – rozpochalis isporuke 4-core AMD Phenom II 940(kodno ime - Deneb). Proizvedeno na frekvenciji od 3 GHz, proizvedeno za 45-nm procesnu tehnologiju.

…

Traven 2009. godine - društvo AMD predstavio verziju grafičkog procesora ATI Radeon HD 4890 s frekvencije takta jezgre povećane s 850 MHz na 1 GHz. Tse prvi grafički procesor koji radi na 1 GHz. Broj naprezanja čipa i povećanje frekvencije porasli su s 1,36 na 1,6 teraflopa. Procesor za zamjenu 800 (!) brojećih jezgri, podržava video memoriju GDDR5, DirectX 10.1, ATI CrossFireX i sve druge tehnologije koje pokreću moderne modele video kartica. Pripreme čipova temeljene na 55-nm tehnologiji.

Osnovne ovlasti GPU

Vídminnimi riža GPU(ispravljeno iz ) є:

- arhitektura, maksimalno usmjerena na povećanje fleksibilnosti tekstura i sklopivih grafičkih objekata;

- vršni tlak tipičan GPU bogatiji, niži ;

– voditelji posebne arhitekture pokretne trake, GPU učinkovitiji u obradi grafičkih informacija, niži.

"Kriza žanra"

"Kriza žanra" za sazrijevali do 2005. godine, - pojavili su se i sami. Ali, bez obzira na razvoj tehnologije, rast produktivnosti izrazito smanjena. Produktivnost vodenog sata GPU nastaviti rasti. Da, do 2003. i revolucionarna ideja se iskristalizirala - vikoristovuvat za potrebe proračuna míts grafike. Grafički procesori počeli su se aktivno koristiti za "negrafička" izračunavanja (simulacija fizike, obrada signala, računalna matematika/geometrija, operacije s bazama podataka, računalna biologija, računalna ekonomija, računalna znanost).

Glavni problem je što nije postojalo standardno sučelje za programiranje GPU. Rozrobniki su koristili OpenGL ili Direct3D ali bilo je preteško. Korporacija NVIDIA(jedna od najvećih kolekcija grafičkih, medijskih i komunikacijskih procesora, kao i medijskih procesora bez strelice; osnovana 1993.) bavila se razvojem jednog ručnog standarda - i predstavila tehnologiju CUDA.

Kako je počelo

2006. godine - NVIDIA demonstrirajući CUDA™; početak revolucije u računanju GPU.

2007. godine - NVIDIA arhitektura izdanja CUDA(Pochatovska verzija CUDA SDK bula predstavljena 15. veljače 2007.); nominacija "Najbolji novitet" u časopisu Popularna znanost i "Vybír chitachív" víd vidannya HPCWire.

2008. godine – tehnologija NVIDIA CUDA osvojio je nominaciju "Tehnička perevaga" u PC magazin.

Što je CUDA

CUDA(kratki engleski) Compute Unified Device Architecture, Doslovno - jedinstveni izračun arhitekture gospodarskih zgrada) - arhitektura (slijed softverskih i hardverskih objekata), koja omogućuje vibracije GPU obračun tajne ispovijedi, u svom GPU Praktički igra ulogu iscrpljujućeg spivprocesora.

Tehnologija NVIDIA CUDA™- srž razvoja mog programiranja C, koji trgovcima omogućuje izradu programa za presavijanje računskih zadataka kraćih od sat vremena, za izračunavanje intenziteta grafičkih procesora. Svijet već prakticira milijune GPU uz podršku CUDA, te tisuće programera već koriste alate (bez troškova!) CUDA za brzo dodavanje i za najnaprednije zadatke uštede resursa – video kodiranje, audio snimanje, istraživanje nafte i plina, modeliranje proizvoda, medicinsko snimanje i znanstveno istraživanje.

CUDA trgovcu daje mogućnost organiziranja sudske istrage s pristupom skupu uputa za grafičko ubrzanje i njegovanje sjećanja, organiziranje novog paralelnog obračuna. Grafički priskoryuvach íz pídtrimkoy CUDA postaje tvrdo ožičena programirana arhitektura, slična današnjoj. Svejedno, naručitelju dajemo nizak, nizak i visok pristup posjedu, roblyach CUDA nužna osnova za ozbiljne alate visoke kvalitete, kao što su kompajleri, kalkulatori, matematičke knjižnice, softverske platforme.

Uralsky, vodeći stručnjak za tehnologije NVIDIA, porívnyuyuchi GPUі reci ovako: - Tse poshlyahovik. Vín í̈zdit zavzhdí taj skíz, ali ne prebrzo. ALI GPU- Ovaj sportski auto. Samo nećete ići nigdje na prljavo skupo vino, ali dajte garni pokrivač, i pokazat ćete svu svoju švedskost, kao poshlyahovik o kakvom nisam ni sanjao! ..».

Mogućnost tehnologije CUDA

CUDA tehnologija

Volodimir Frolov,[e-mail zaštićen]

Sažetak

U članku se govori o tehnologiji CUDA, koja programeru omogućuje hakiranje video kartica kao specifičniji broj. Nvidijini alati omogućuju pisanje programa za grafički procesor (GPU) u C++ podskupovima. To će pomoći programeru s potrebom korištenja shadera i razumijevanja procesa robotskog grafičkog cjevovoda. U članku je predstavljena primjena programiranja s CUDA alternativama i raznim metodama optimizacije.

1. Uvod

Razvoj tehnologija brojanja i dalje je desetke godina brzim tempom. Podovi su brzi, ali u isto vrijeme, prodavači procesora praktički su otišli do takozvane "silikonske gluhe kute". Nezamislivo povećanje frekvencije takta postalo je nemoguće zbog brojnih ozbiljnih tehnoloških razloga.

Zato svi algoritmi modernih sustava brojanja idu za povećanjem broja procesora i jezgri, a ne za povećanjem frekvencije jednog procesora. Broj jezgri središnje procesorske jedinice (CPU) u naprednim sustavima veći je od 8.

Drugi razlog je očito niska brzina robota i operativne memorije. Kao da procesor ne radi brzo, u uskim prostorima, kao što pokazuje praksa, to nisu aritmetičke operacije, već bliske promašajima predmemorije memorije.

Prote se čuditi bik grafičkim procesorima GPU-a (Graphics Processing Unit), tu su, putem paralelizma, išli bogato ranije. U trenutnim video karticama, na primjer GF8800GTX, broj procesora može doseći 128. Produktivnost takvih sustava, ako se pravilno programira, može biti značajna (slika 1).

Riža. 1. Broj operacija s pomičnim zarezom za CPU i GPU

Da su se prve grafičke kartice upravo pojavile u prodaji, smrad bi činili jednostavni (zajedno sa središnjim procesorom) visokospecijalizirani dodaci, koji su bili namijenjeni uklanjanju procesora iz vizualizacije podataka dva svijeta. S razvojem industrije igara i pojavom tako trivijalnih igara kao što su Doom (Slika 2) i Wolfenstein 3D (Slika 3) vinilna potražnja za 3D vizualizacijom.

Bebe 2.3. Igre Doom i Wolfenstein 3D

Prve Voodoo video kartice, (1996.) i sve do 2001., stvorila je tvrtka 3Dfx neposredno prije 2001. u GPU-u, samo popravljajući skup operacija na ulaznim podacima.

Programeri nisu imali izbora algoritama za vizualizaciju, a povećana je fleksibilnost, pojavili su se shaderi - mali programi koji vizualiziraju vrh kože i piksel kože s video karticom. Njihovi su zadaci uključivali transformacije preko vrhova i sjenčanje-osvjetljavanje u točkama, na primjer, za model Phong.

Želeći ponekad, shaderi su oduzeli čak i snažan razvoj, nakon što su shvatili da je smrad razvijen za visokoškolske ustanove trivi- mernih transformacija i rosterizacije. U isto vrijeme, kako se GPU-ovi razvijaju u višenamjenskim bogatim procesorskim sustavima, filmski shaderi su preplavljeni visokim specijalizacijama.

Moguće je usporediti s mojim FORTRAN-om na činjenicu da su smrdi, poput FORTRAN-a, bili prvi, ali su bili prepoznati za vyrishennya samo jednu vrstu zadatka. Shaderi nisu prikladni za usavršavanje bilo kojih drugih zadataka, osim trivijalnih transformacija i rasterizacija, poput FORTRAN-a, nisu prikladni za dovršetak zadatka koji nije povezan s numeričkim rasporedom.

Danas postoji trend netradicionalnih video kartica za vizualizaciju u ormarima kvantne mehanike, inteligencije komada, fizičkih rekonstrukcija, kriptografije, fizički ispravne vizualizacije, rekonstrukcije iz fotografija, identifikacije. Qi zavdannya neručno je vibrirao na granicama grafičkih API-ja (DirectX, OpenGL), krhotine qi API-ja stvorili su drugi zastosuvani.

Razvoj općeg programiranja na GPU-u (General Programming on GPU, GPGPU) logično je doveo do opravdanja tehnologija usmjerenih na šire područje proizvodnje, nižu rosterizaciju. Kao rezultat toga, Nvidia je stvorila tehnologiju Compute Unified Device Architecture (ili skraćeno CUDA), a ATI, koji se natječe, stvorio je tehnologiju STREAM.

Treba napomenuti da je u trenutku pisanja ovog članka STREAM tehnologija bila izrazito naklonjena razvoju CUDA-e, a to se ovdje ne može vidjeti. Usredotočeni smo na CUDA - GPGPU tehnologiju, koja vam omogućuje pisanje programa u više C++ filmova.

2. Glavna razlika između CPU-a i GPU-a

Pogledajmo ukratko pojedinosti o razlikama između regija i specifičnim značajkama središnjeg procesora i video kartice.

2.1. Mogućnosti

CPU hrpa privitaka za izvođenje glavnog plana i rad na memoriji, što je prikladno riješeno. Programi na CPU-u mogu se bez prekida preuzeti u sredinu linearne ili homogene memorije.

Za GPU, nije loše. Kao što znate, nakon što ste pročitali ovaj članak, CUDA može vidjeti 6 vrsta memorije. Možete čitati iz bilo koje sredine, fizički dostupne, ali zapisujte - ne u sredini. Razlog leži u činjenici da je GPU, u svakom slučaju, specifičan dodatak, prepoznajemo ga za određene namjene. Tse obezhennya zaprovadzhennja za zbílshennya svydkostí roboti singhnyh algogorívív í zvízhennya vartosti í obladannya.

2.2. Švedski kod pamćenja

Isti problem računalnih sustava je da je memorija učinkovitija od procesora. CPU hakeri na neki način krše predmemorije. Većinu vremena pomak memorije provodi se u super-operativnoj ili cache memoriji, koja radi na frekvenciji procesora. Tse vam omogućuje da odvojite sat vremena uz smrt do smrti, koja je najčešće pobjednička, i zavantazhit procesor na najbolji mogući način.

S poštovanjem, za programera, predmemorije su praktički čiste. Kao i kod čitanja, tako i kod pisanja podaci se ne unose jednokratno u operativnu memoriju, već prolaze kroz predmemorije. Dopustite mi, zokrema, da brzo pročitam značenje dana nakon unosa.

Na GPU-u (ovdje možete koristiti video kartice 8. serije GF) predmemorije su također važne, ali mehanizam nije tako težak kao na CPU-u. Na prvi način, unovčavanje pune opsjednutosti vrstama memorije, ali na drugačiji način, predmemorije se prakticiraju samo od čitanja.

Na GPU-u ima više nego dovoljno vremena da se ne zaboravite platiti za dodatne paralelne izračune. Za sada, jedan zavdannya provjerava podatke, pratsyut ínshí, spreman za izračun. Ovo je jedno od glavnih načela CUDA-e, koje vam omogućuje znatno povećanje produktivnosti sustava u cjelini.

3. CUDA jezgra

3.1. Model protoka

Numerička arhitektura CUDA-e temelji se na konceptujedan tim za anonimne podatke(Jedna instrukcija s više podataka, SIMD) višeprocesorski.

Može se koristiti koncept SIMD-a, da jedna instrukcija omogućuje prikupljanje anonimnih podataka odjednom. Na primjer, naredba addps u procesoru Pentium 3 i novijim modelima Pentiuma omogućuje dodavanje 4 broja jednostruke preciznosti s pomičnim zarezom odjednom.

Višeprocesor je višejezgreni SIMD procesor koji ima samo jednu instrukciju na svim jezgrama. Koža višeprocesorske jezgre nije skalarna, dakle. ne podržava vektorske operacije na čist način.

Prije Tima, kako dalje, uvest ćemo par termina. Značajno, pod privitkom i domaćinom ove statistike, to nisu oni prije kojih je većina programera pozvala. Koristit ćemo takve izraze kako bismo izbjegli razlike u odnosu na CUDA dokumentaciju.

Pod uređajem (uređajem), naš članak ima razuman video adapter koji podržava CUDA upravljački program ili druge specijalizacije privitaka, zadataka za programiranje programa koji koriste CUDA (na primjer, poput NVIDIA Tesla). U našem članku možemo vidjeti da je GPU manje nalik logičnom posjedu, jedinstvenom za specifične detalje implementacije.

Domaćin (host) je naziv programa u glavnoj operativnoj memoriji računala, nadjačavajući CPU i nadjačavajući ključne funkcije robota s privitkom.

Zapravo, onaj dio vašeg programa koji radi na CPU je domaćin, i tvoja video kartica - privitak. Logično, možete se prijaviti kao skup multiprocesora (mala 4) plus CUDA drajver.

Riža. 4. Prilog

Recimo da želimo pokrenuti proceduru u N niti na našem proširenju (zato želimo paralelizirati robota). Što se tiče CUDA dokumentacije, nazovimo proceduru kernel.

Osobitost CUDA arhitekture je organizacija blok-sita, koja nije ograničena bogatim aditivima za protok (slika 5). CUDA upravljački program neovisno distribuira resurse i gradi između niti.

Riža. 5. Organizacija tokova

Na sl. 5. jezgra je označena kao Kernel. Sve niti koje pogađaju jezgru kombiniraju se u blokove (Block), a blokovi se, po svojoj prirodi, kombiniraju u mrežu (Grid).

Kao što se može vidjeti na slici 5, indeksi dva svijeta koriste se za identifikaciju tokova. CUDA rozrobnici su dali mogućnost rada s trivum, two-world ili jednostavnim (single-world) indeksima, ovisno o tome koliko je programer spretan.

U divljem tipu, indeks je trivijalan po vektorima. Za dermalnu nit bit će dano sljedeće: indeks niti u sredini threadIdx bloka i indeks bloka u sredini blockIdx mreže. Prilikom pokretanja, sve će se niti osvježiti bez više indeksa. Zapravo, putem qi indeksa, programator kontrolira kontrolu, što znači da se dio toga obrađuje u znoju kože.

Dokazi za opskrbu, zašto su trgovci sebi oduzeli takvu organizaciju, nisu trivijalni. Čini se da je jedan od razloga to što će jedan blok zajamčeno pobijediti na jedan Dodat ću multiprocesor, ali jedan multiprocesor može osvojiti gomilu raznih blokova. Ostali razlozi za raščišćavanje dati su u satu članka.

Blok zadataka (tijekovi) se pobjeđuje na multiprocesoru dijelovima ili skupovima koji se nazivaju warp. Proširenje warpa za trenutni trenutak u video karticama s podrškom za CUDA je do 32 toka. Naredbe u sredini warp skupa postavljene su na SIMD stilove, tj. sve niti u sredini warpa mogu imati samo jednu instrukciju u isto vrijeme.

Evo sljedećeg upozorenja. U arhitekturama koje su trenutno u trenutku pisanja ovog članka, broj procesora u sredini jednog multiprocesora je 8, a ne 32. Zbog toga je jasno da se ne pobjeđuju svi warpi odjednom, oni su razbijeni na 4. dijelovi, tuku se sekvencijalno (jer procesi) .

Ali, prije svega, trgovci CUDA ne reguliraju veličinu osnove. U njihovim robotima, smrad postavlja parametar veličine osnove, a ne broj 32. Na drugačiji način, s logičke točke gledišta, sama osnova je minimalni bazen tokova, za koji se može reći da su svi tokovi u sredina ciklusa računa se odjednom - i sa svakim danom, dopustite riješiti sustav neće biti slomljen.

3.1.1. Odsoljavanje

Pa, vi krivite hranu: ako baš u tom trenutku svi tokovi u sredini warpa tipkaju upravo tu instrukciju, kako se onda možete raščistiti? Čak i ako se programski kod očisti, upute će biti drugačije. Ovdje na scenu stupa standardno SIMD programsko rješenje (slika 6).

Riža. 6. Organizacija otklanjanja pogrešaka u SIMD-u

Hajde sljedeći kod:

ako (uvjet)B;

U slučaju SISD (Single Instruction Single Data), operator A se vikonira, um ga ponovno provjerava, zatim se operatori B i D vikoniraju (tako da je um istinit).

Hajde sada imati 10 streamova koji su napisani u stilu SIMD-a. U svih 10 tokova pada nam operator A, zatim provjeravamo mentalni uvjet i čini se da je u 9 od 10 tokova točan, au jednom je loš.

Shvatio sam da ne možemo pokrenuti 9 niti za izvođenje operatora B i jednu za izvođenje operatora C, tako da sve niti mogu izvoditi samo jednu po jednu instrukciju. Ovoj vapadki valja još dodati ovako: lonac se “zabije” na zatiljak, da se razbistri, da vino ne daje danak, i uzme se 9 potoka, koji se ostave. van. "Ubacimo" 9 dretvi koje pobjeđuju operator B, i prođimo jednu dretu s operatorom C. Sljedeće dretve se ponovno kombiniraju i pobjeđuju operator D sve odjednom.

Postoji sažeti rezultat: ne samo da se resursi procesora iscrpljuju na prazno mljevenje bitaka u tokovima koji su prekinuti, toliko bogatiji, bit ćemo osramoćeni kao rezultat poraza OBIGI gílki.

Ipak, nije sve tako loše, kao što se vidi na prvi pogled. Kao velika prednost tehnologije, možete vidjeti one koji se fokusiraju na dinamičko pokretanje CUDA drajvera, a za programera, smrad je jasan. U isto vrijeme, trčeći uokolo sa SSE naredbama iz trenutnih procesora (samo pokušajte s 4 kopije algoritma odjednom), programer je sam odgovoran za detalje: kombinirajte podatke po četiri, ne zaboravite na pregled i počnite pisati na niska razina, zapravo, kao u asembleru.

Iz brkova gorespomenutog cvili još jedan poštovani visnovok. Razgaluzhennya je razlog za pad produktivnosti od moćnih sila. Shkidlivimi ê manje od tí razgaluzhennya, na kojem se potoci razilaze u sredini jednog bazena osnovnih potoka. U nekim slučajevima, tokovi su se proširili u sredini jednog bloka, ali u različitim warp bazenima, ili u sredini različitih blokova, bez ikakvog učinka.

3.1.2. Interakcija između tokova

U trenutku pisanja ovog članka, je li interakcija između niti (sinkronizacija i razmjena podataka) bila moguća samo u sredini bloka. Zbog toga je nemoguće međusobno organizirati tokove različitih blokova, koji su korozivni s manje mogućnosti dokumentiranja.

Iako nedokumentirane mogućnosti, ne preporuča se njima uljepšavati. Razlog je taj što se smrad temelji na specifičnim hardverskim značajkama drugog sustava.

Sinkronizacija svih zadataka u sredini bloka kontrolirana je pozivom funkcije __synchtreads. Razmjena novca moguća je kroz memoriju koja je podijeljena pa je tako važna za sve zadatke u sredini bloka.

3.2. Memorija

CUDA ima šest vrsta memorije (slika 7). Ce registar, lokalna, globalna, distribuirana, konstantna i teksturna memorija.

Na tako veliki broj potaknule su specifičnosti video kartice i prva priznanja, a također i trgovci na malo da izgrade yakomoga sustav jeftinije, žrtvujući na različite načine bilo univerzalni ili švedski.

Riža. 7. Vidi memoriju CUDA

3.2.0. Registar

Koliko god može, prevodilac pokušava smjestiti sve lokalno promijenjene funkcije u registre. Pristup takvim promjenama ograničen je na maksimalnu brzinu. Streaming arhitektura ima 8192 32-bitna registra dostupna po multiprocesoru. Da odredite koliko je registara dostupno za jednu nit, trebate podijeliti broj (8192) po veličini bloka (broj niti za novu).

S pozamašnim brojem od 64 protoka, ukupno 128 registara će ući u blok (postoje neki objektivni kriteriji, ali 64 bi trebalo biti u sredini za bogate poslove). Stvarno, 128 registara nvcc se uopće ne vidi. Zovite VIN, ne dajte više od 40, ali izbrišite promijenjene u lokalnu memoriju. Dakle, čini se da se na jednom multiprocesoru može osvojiti nekoliko blokova. Kompajler će pokušati maksimizirati broj blokova koji se mogu obraditi odjednom. Za veću učinkovitost, trebate uzeti manje od 32 registra. Zatim, teoretski, možete pokrenuti 4 bloka (8 warp-ív, dakle 64 niti u jednom bloku) na jednom multiprocesoru. Međutim, ovdje je važnije voditi računa o zajedničkoj memoriji koja je zauzeta nitima, budući da jedan blok zauzima svu memoriju koja se dijeli, dva takva bloka ne mogu biti zauzeta multiprocesorom u isto vrijeme.

3.2.1. Lokalna memorija

U slučaju, ako lokalni podaci procedura zauzimaju previše prostora, ili prevodilac to ne može izračunati zadnji put, možete ih smjestiti u lokalnu memoriju. Kome možete prihvatiti, primjerice, dane pokazatelje vrsta različitih vrsta proširenja.

Fizički, lokalna memorija analogna je globalnoj memoriji i radi s tíêyu i swidkíst. U vrijeme pisanja članka nije bilo poznatih mehanizama koji bi omogućili eksplicitno dohvaćanje prevoditelja korištenjem lokalne memorije za određene promjene. Važno je provjeriti lokalnu memoriju, a ne koristiti je zovsím (div. odjeljak 4 "Preporuke za optimizaciju").

3.2.2. globalna memorija

CUDA dokumentacija jedno je od glavnih postignućaTehnologija za induciranje mogućnosti dovoljnog adresiranja globalne memorije. Dakle, možete čitati iz bilo koje sredine memorije, a možete pisati na isti način na određenoj sredini (na GPU-u to ne zvuči tako).

Prote za svestranost u vrijeme da se dovede do plakati swidkistyu. Globalna memorija nije predmemorirana. Vaughn pratsyuê još više povílno, kílkíst zvernení u globalnu memoriju síd u bilo kojem trenutku minimizuvati.

Globalna memorija neophodna je za spremanje rezultata robotskih programa prije njihovog uređivanja na glavnom računalu (na zadanoj DRAM memoriji). Razlog tome je što je memorija globalna - jedina vrsta memorije koja se može zabilježiti.

Promjene, izražene kvalifikatorom __globalno__, smještene u sjećanje svijeta. Globalna memorija također se može pregledavati dinamički pozivanjem funkcije cudaMalloc(void* mem, int size) na glavnom računalu. Dodat ću ovu funkciju, nije je moguće pozvati. Zvuči kao da se glavni program može pobrinuti za memoriju, ono što radi na CPU-u. Podatke s glavnog računala možete nadjačati klikom na funkciju cudaMemcpy:

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

Dakle, možete sami pokrenuti obrnuti postupak:

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

Ovaj wiki također radi s hosta.

Kada radite s globalnom memorijom, važno je zapamtiti pravila spajanja. Glavna ideja je da se treći dio memorije pohranjuje do zadnje srednje memorije, štoviše, 4,8 ili 16 bajtova. Ovime je prva nit kriva za traženje adrese, vibrirajući na kordonu, očito 4,8 ili 16 bajtova. Adrese koje rotira cudaMalloc su najmanje 256 bajtova preko granice.

3.2.3. Zapamtite što se dijeli

Memorija koja se dijeli ne može se predmemorirati, ali je memorija brza. Njenu preporuča se vicorate kao provjeru predmemorije. Ukupno je dostupno 16 KB zajedničke memorije po multiprocesoru. Dijeljenjem broja s brojem dana u bloku, uzimamo maksimalnu količinu memorije koja se dijeli, dostupna za jedan stream (jer je planirano osvojiti ga neovisno o svim streamovima).

Zapamtite rižu koja je podijeljena, one koje se adresiraju isto za sve vođe srednjeg bloka (slika 7). Očito je da možete osvojiti samo jedan blok za razmjenu podataka između streamova.

Zajamčeno je da će sat vremena blok na multiprocesoru biti pohranjen u memoriji. Međutim, budući da se blok mijenja na multiprocesoru, nije zajamčeno da će umjesto njega biti spremljen stari blok. Zato nije dovoljno pokušati sinkronizirati zadatke između blokova, ostavljajući ih u memoriji kao podatke i oslanjajući se na njihovu uštedu.

Promjene, izražene kvalifikatorom __shared__, smještaju se u memoriju da se dijele.

shared_float mem_shared;

Sljedeći put, uvjeri, što je memorija, što se dijeli, za blok sam. Za ovo je potrebno hakirati baš kao i cache cache, prateći pretraživanje različitih elemenata niza, na primjer:

float x = mem_shared;

De threadIdx.x – indeks x niti u sredini bloka.

3.2.4. Stalna memorija

Konstantna memorija je predmemorirana, kao što se vidi na sl. 4. Predmemorija se koristi u jednoj instanci jednog multiprocesora, što je ujedno i glavna zadaća srednjeg bloka. Na glavnom računalu možete pisati u konstantnu memoriju pozivanjem funkcije cudaMemcpyToSymbol. Dodat ću stalnu memoriju koja je dostupna samo za čitanje.

Stalna memorija je prikladnija za vikoristan. Možete rozmíschuvati u niy daní biti-koji tip koji ih čitati za pomoć jednostavne privlačnosti.

#definiraj N 100

Konstanta__int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

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

// __global__ znači da je device_kernel kernel, pa se može pokrenuti na GPU-u

Global__void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ZABAVA! stalna memorija dostupna je samo za čitanje

Budući da se predmemorija koristi za stalnu memoriju, pristup joj je siguran. Jedina, ali još uvijek velika, mala količina stalne memorije je u činjenici da ona postaje manja od 64 Kbajta (za cijeli privitak). Zašto je tako očito da u kontekstualnu memoriju ima smisla spremiti samo mali broj podataka koji su često pobjednici.

3.2.5. Memorija teksture

Memorija teksture je predmemorirana (slika 4). Za skin multiprocesor postoji samo jedna predmemorija, tako da je cijela predmemorija najveća za cijeli srednji blok.

Naziv teksturne memorije (i, nažalost, funkcionalnost) smanjen je da bi se razumjeli "tekstura" i "teksturiranje". Teksturiranje - proces primjene tekstura (samo slika) na poligon tijekom procesa rasterizacije. Memorija teksture optimizirana je za 2D podatke i može biti moguća:

shvidka vibirka vrijednost fiksne rozmíru (bajt, riječ, subwiyne ili četiri riječi) iz niza s jednim ili dva svijeta;

normalizirano adresiranje float brojevima u intervalima. Potom ih je moguće odabrati, koristeći normalízovanu adresiranje. Rezultirajuća vrijednost bit će riječ tipa float4, s intervalom;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // očito memorija GPU-a

// Podesite parametre stabla teksture

Texture.addressMode = cudaAddressModeWrap; // način rada Zamotati

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //najbliža vrijednost

texture.normalized = lažno; // nemoj izvrtati normalizirano adresiranje

CudaBindTexture(0, tekstura, gpu_memory, N ) // promijenimo memoriju u teksturu


cudaMemcpy(gpu_memory, cpu_buffer, N*sizeof(uint 4), cudaMemcpyHostToDevice ); // kopiraj podatke uGPU

// __global__ znači da je device_kernel kernel, pa ga treba paralelizirati

Global__void device_kernel()

uint4 a = tex1Dfetch(tekstura,0); // na ovaj način možete odabrati više podataka!

uint4 b = tex1Dfetch(tekstura,1);

int c = a.x*b.y;

...

3.3. Jednostavna guzica

Kao jednostavan dodatak, predlaže se da pogledate cppIntegration program iz CUDA SDK. Vaughn demonstrira tijek rada CUDA-e, kao i nvcc (specijalni C++ prevodilac za Nvidiju) verziju MS Visual Studija, koja će pojednostaviti razvoj programa na CUDA-i.

4.1. Pravilno izvedite poraz svog šefa

Nisu svi zadaci prikladni za SIMD arhitekturu. Na primjer, vaš zadatak za koga nije priložen, moguće je, nije moguće koristiti GPU. Pa ipak, čvrsto ste prekršili GPU trikove, bilo je potrebno razbiti algoritam na takve dijelove, kako bi smrad mogao učinkovito poraziti SIMD stil. Potrebno je - promijeniti algoritam za poboljšanje vašeg zadatka, osmisliti novi - onaj koji bi bio dobar za SIMD. Kao primjer drugačijeg područja GPU-a, možete implementirati piramidalno preklapanje elemenata u nizu.

4.2. Odaberite vrstu memorije

Smjestite svoje podatke u teksturnu ili stalnu memoriju, tako da se svi zadaci jednog bloka pretvaraju u jedan memorijski utor ili blizu naslaganih utora. Ta dva podatka mogu se učinkovito obraditi pomoću dodatnih funkcija text2Dfetch i text2D. Memorija teksture posebno je optimizirana za dva svijeta.

Osvojite globalnu memoriju iz memorije koja se dijeli, jer su svi zadaci nesustavno raspoređeni na različite, dalekosežne, jedne vrste jedne memorijske lokacije (s različitim adresama ili koordinatama, poput 2D / 3D podataka).

globalna memorija => memorija koja je podijeljena

syncthreads();

Prikupite podatke u memoriji

syncthreads();

globalna memorija<= разделяемая память

4.3. Podsjeti lichnike na sjećanje

Oznaka prevoditelja --ptxas-options=-v omogućuje vam da točno kažete koje riječi i koju memoriju (registri, koji su distribuirani, lokalni, konstantni) u victoristu. Budući da kompajler koristi lokalnu memoriju, znate za to. Analiza podataka o broju i vrstama memorije koje pobjeđuju može vam uvelike pomoći u optimizaciji programa.

4.4. Pokušajte minimizirati broj registara i memorija koje se dijele

Što je veća jezgra warp registra, ili memorija koja je podijeljena, manje protoka (primarni warp-iv) može istovremeno pobijediti na multiprocesoru, jer razmjenjuju se resursi multiprocesora. Dakle, malo povećanje popunjenosti registara, ali memorija koje se dijele, može dovesti do smanjenja produktivnosti za faktor dva u nekim slučajevima - čak i kroz one koji su sada čak dvostruko manje warp-in na jednom na multiprocesoru.

4.5. Sjećanje na ono što se dijeli, lokalni zamjenik.

Kao da je Nvidia kompajler uzrokovao oštećenje podataka u lokalnoj memoriji 'yat, scho razdelyaetsya (zajednička memorija).

Većinu vremena kompajler se može promijeniti u lokalnoj memoriji, jer se neće često koristiti. Na primjer, baterija devi akumulira vrijednosti, rozrakhovuyuschos s tsiklí. Kao veliki ciklus za obsyagi kod (iako ne za jedan sat vykonannya!), prevodilac može staviti vašu bateriju u lokalnu memoriju, tk. Vín vykoristovuêtsya vrlo rijetko, a registri su malo. Trošak produktivnosti ponekad može biti značajan.

Pa, ako ga stvarno rijetko mijenjate - bolje ga smjestite u globalnu memoriju.

Ako želite da prevoditelj automatski alocira takve promjene u lokalnoj memoriji, to može biti razumno, ali ne baš. Nije lako znati na određenom mjestu s nadolazećim izmjenama programa, jer se često mijenja češće. Prevodilac može ili ne mora prenijeti takvu promjenu u registarsku memoriju. Ako se modifikator __global__ eksplicitno navede, programer će biti brutalnije svjestan.

4.6. Rigthation ciklusi

Rotirajući ciklusi standardna su metoda za povećanje produktivnosti u bogatim sustavima. Bit yoga je da na skin iteraciji možete pobjeđivati ​​sve više i više, mijenjajući na taj način broj iteracija, a to znači i broj mentalnih prijelaza, kako bi procesor pobjeđivao.

Os je kako je moguće otvoriti ciklus značenja sumi niza (na primjer, cijeli):

int a[N]; intsum;

za (int i=0;i

Zrozumilo, ciklusi se mogu ispaliti i ručno (kao što je prikazano gore), ali praksa je neproduktivna. Bolje je prilagoditi C++ predloške za više funkcija koje trebate znati.

šablona

klasa NizZbroj

Uređaj__ static T exec (const T * arr) ( return arr + ArraySumm (arr+1); )

šablona

klasa NizZbroj<0,T>

Device__ static T exec(const T* arr) ( return 0; )

za (int i=0;i

zbroj+= Zbroj niza<4,int>::izvrši(a);

Označite jednu kvačicu kao značajku nvcc prevoditelja. Prevodilac uvijek treba dopustiti da se funkcije tipa __device__ zatvore (kako biste bili sigurni, upotrijebite posebnu direktivu __noinline__).

Otzhe, možete biti inspirirani činjenicom da se stražnjica, slična šiljastoj stvari, ljuti na jednostavan niz operatora, a zašto ne djelovati za učinkovitost koda napisanog rukom. Međutim, za divlji tip (ne nvcc) nije moguće da bilo tko ima plijen, jer inline nije ništa više od instrukcija prevoditelja, i možete je ignorirati. Nije zajamčeno da će vaše funkcije biti vraćene.

4.7. Pregledajte podatke i odaberite 16 bajtova

Pregledajte strukture podataka prema 16-bajtnom kordonu. U tom slučaju, kompajler može napisati posebne upute za njih, koje će pobijediti podatke jednom za 16 bajtova.

Ako je struktura kredita 8 b ili manje, možete je promijeniti za 8 b. Alternativno, možete odlučiti kombinirati dvije promjene od 8 bajtova odjednom kombinirajući dvije promjene od 8 bajtova u strukturu iza dodatne unije ili ih navesti na popisu. Sa sljedećim treba postupati pažljivo, kompajler može pohraniti podatke u lokalnu memoriju, ali ne i u registar.

4.8. Bankovni sukobi u memoriji, koji su podijeljeni

Memorija koja se distribuira organizirana je u 16 (sve u svemu!) banaka memorije s 4-bajtnim krokom. Tijekom sat vremena skup warp niti na multiprocesoru podijeljen je u dvije polovice (kao što je warp-size = 32) od po 16 niti, što omogućuje pristup memoriji karticom.

Zadaci u različitim polovicama warpa nisu u sukobu s podijeljenim sjećanjima. Kroz zavdannya jedna polovica warp bazena bit će reducirana na iste banke memorije, okrivljujući kolaps i, kao rezultat, pad produktivnosti. Zavdannya u granicama jedne polovice osnove može narasti do različitih sela sjećanja, koja su podijeljena, s pjesmom krokodila.

Optimalna veličina je 4, 12, 28, ..., 2 n-4 bajta (slika 8).



Riža. 8. Optimalno pristajanje.

Chi nije optimalna veličina - 1, 8, 16, 32, ..., 2^n bajtova (slika 9).



Riža. 9. Neoptimalno pristajanje

4.9. Minimiziranje pokretnih podataka Host<=>uređaj

Pokušajte prenijeti međurezultate na host za obradu za dodatni CPU. Ako ne implementirate cijeli algoritam, uzmite dio glavnog dijela na GPU, ostavljajući CPU manje nego što je potrebno.

5. CPU/GPU prijenosna matematička biblioteka

Autor ovog članka napisao je biblioteku MGML_MATH, koja se može prenijeti, za rad s jednostavnim svemirskim objektima, kod kao praktičan i na proširenje i na host.

Knjižnica MGML_MATH može se koristiti kao okvir za pisanje CPU/GPU prijenosnih (ili hibridnih) sustava za razvoj fizičkih, grafičkih i drugih svemirskih zadataka. Glavna prednost je da se jedan te isti kod može podešavati i na CPU-u i na GPU-u, a ako se prezentira biblioteci, može se instalirati i na drugoj strani.

6 . Književnost

Chris Kaspersky. Tehnika optimizacije programa. Učinkovit oporavak memorije. - St. Petersburg: BHV-Petersburg, 2003. - 464 str.: il.

Vodič za programiranje CUDA 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

Vodič za programiranje CUDA 1.1. stranica 14-15

Vodič za programiranje CUDA 1.1. stranica 48

Iz Darwinove teorije evolucije, prvi čovjekoliki mavpa (poput
buti točnije - homo antecessor, person-perepadnik) pretvarao se da je bog
imamo. Centri za brojanje bagatotona s tisuću i više radio lampi,
ono što zauzima zidine grada, promijenjeno pivkilogramskim novčanicama, yaki, do govora,
nemojte prvo žrtvovati produktivnost. Pretpotopni drukarski strojevi su se promijenili
za druge, uvijek je dobro za ono za što je dobro (navit na tijelu osobe)
bogato funkcionalne gospodarske zgrade. Giganti raptom procesore odlučili su zazidati
grafička jezgra "kamena". I video kartice su počele ne samo prikazivati ​​sliku
prihvatiti FPS i kvalitetu grafike, ali i izvršiti sve izračune. Tako
više posla! O tehnologiji bogatog strujanja, izračunajte je uz pomoć GPU-a, i to će biti spomenuto.

Zašto GPU?

Tsikavo, zašto su prenijeli cijelo brojanje míts na grafiku
adapter? Kao što vidite, procesor je u modi i malo je vjerojatno da će se odreći svoje topline
Mistečko. Ale, GPU ima par koziriva na rukavu odjednom sa jokerom, taj rukav.
vistachaê. Suvremeni središnji procesor zatvorske kazne za ograničenje maksimuma
produktivnost tijekom obrade cilijarnih podataka i podataka s plutajućim
Komi, pogotovo bez ometanja bilo kakve paralelne obrade informacija. Isti
sat vremena arhitektura video kartice omogućuje brz i bez problema "paralelizam"
prikupljanje podataka. S jedne strane, postoji izgled poligona (za okvir 3D transportne trake),
z ínshoy – obrada tekstura piksela. Vidi se da je „sretan
breakdown” u središtu kartice. Osim toga, memorijski robot i video procesor
optimalna, niža veza "RAM-cache-processor". U tom trenutku, ako ste usamljeni
video kartice počinju se pretvarati jednim GPU stream procesorom, inače
usamljenost je istovremeno upletena u drugoga i, u principu, lako se može dosegnuti
učinkovitost grafičkog procesora, zajedno s povećanjem propusnosti sabirnice,
međutim, za ovaj vantazhennya transporter može
sve vrste pametnih prijelaza i razgaluzhen. središnji procesor kroz svoj
svestranost za vlastite potrebe procesora, predmemorija, pohrana
informacija.

Vchení cholovíki zamyslilis schodo roboti GPU u paralelnim izračunima koji
matematičari su razvili teoriju da postoji mnogo znanstvenih istraživanja koja su na mnogo načina slična
obrada 3D grafike. Mnogi stručnjaci smatraju koji je glavni čimbenik
razvoj GPGPU (Računanje opće namjene na GPU-u – univerzalno
rozrahunki uz pomoć video kartice) uveden je 2003. u projekt Brook GPU.

Tvorci projekta sa Sveučilišta Stanford imali su problema
problem: hardver i softver zmusiti grafički adapter vibracija
razne planove. Smrdim veyshlo. Vikoristovuyuchi univerzalni mov C,
američki znanstvenici zmusili pratsyuvati GPU poput procesora, prilagođen za
paralelni rez. Nakon što je Brook pokazao cijeli niz projekata za VGA-projekte,
poput knjižnice Accelerator, knjižnice Brahma, sustava
metaprogramiranje GPU++ i drugi.

CUDA!

Percepcija perspektive razvoja bila je neugodna AMDі NVIDIA
ući u Brook GPU kao pit bull. Ako izostavite marketinšku politiku, onda
nakon što ste sve ispravno implementirali, možete zatvoriti barem u grafičkom sektoru
tržište i brojanje (divite se posebnim kartama za brojanje koje
poslužitelji Tesla sa stotinama multiprocesora), mijenjajući nazive za sve CPU-e.

Zvichayno, "FPS Volodymyrs" su se uzdigli do kamena posrnule kože za svoje.
šavova, ali glavno načelo postalo je nepromjenjivo - izvršite izračun
od strane GPU-a. Pogledajmo odmah "zelenu" tehnologiju. CUDA
(Compute Unified Device Architecture).

Robot naše "heroine" radi na sigurnom API-ju, štoviše, jednom za dvoje.
Prvi je high-core, CUDA Runtime, s funkcijama, poput
više jednostavnih jednakih se razbija i prenosi na niži API - CUDA Driver. Tako
da izraz "vrlo vjerodostojan" stagnira do natezanja. Sva moć znanja
sama kod vozača, i dobiti je za pomoć knjižnicama, ljubazno stvorena
trgovci na malo NVIDIA: CUBLAS
FFT (istraživanje izgleda Fur'e algoritma). Pa, prijeđimo na praktičnost
dijelovi materijala.

Terminologija CUDA

NVIDIA Radi s vlastitim dodjelama za CUDA API. tenis
v_dr_znyayutsya víd vyznachen, scho zastosovuyutsya za rad sa središnjim procesorom.

Potik (nit)- Prikupiti podatke koje je potrebno prikupiti (ne
Vymagaê veliki resursív píd sat orobki).

iskriviti se- Grupa od 32 toka. Podaci se prikupljaju samo
warps, aka warp - minimalni obsyag podataka.

blok- Sukupníst teče (víd 64 do 512) ili sukupníst
warpiv (tip 2 do 16).

Sitka (mreža)- Tse sukupníst blokív. Takav podíl danih
zastosovuêtsya vikljuchno podvischennya produktivnost. Dakle, koji je broj
multiprocesor velik, tada su blokovi paralelno spojeni. Yakscho w
nisam štedio karticu (preporuka trgovaca za sklopive ruže
adapter nije niži od GeForce 8800 (GTS 320 MB), tada se blokovi podataka obrađuju
sekvencijalno.

Također NVIDIA predstaviti tako razumljivo, kao zrno, domaćin (domaćin)
і uređaj.

Pratsyuemo!

Za opći rad sa CUDA-om potrebno vam je:

1. Upoznajte prirodu GPU shader jezgri, krhotine su bit programiranja
pogaê imaju jednake rozpodílí navantazhennya između njih.
2. Programirati sredinu C-a, popraviti neke aspekte.

Rozrobniki NVIDIA razbio "punjenja" video kartice kílka
inače, donji mi zvan bachiti. Pa zašto ne želiš slučajno vidjeti sve
suptilna arhitektura. Roberemo Budov "kamen" G80 legendarni GeForce 8800
GTX.

Jezgru shadera čini osam TPC (Texture Processor Cluster) klastera
procesori tekstura (npr. GeForce GTX 280- 15 jezgri 8800 GTS
njihovih šest, y 8600 - Chotiri, itd.). Ti, u srcu, zbroji u dva
strujni višeprocesori (streaming multiprocessor – far SM). SM (ukupno šest
16) presavijeni s prednje strane (zadatak čitanja i dešifriranja uputa) i
back end (kraj instrukcija) cjevovoda, kao i osam skalarnih SP (shader
procesor) i dvije SFU (Super Functional Units). Za skin beat (single
sat) front end izabrati warp i raditi jogu. Shchob usí warp teče
(Pretpostavljam, njihova 32 komada) pokazalo se da je potrebno 32/8 = 4 ciklusa na kraju cjevovoda.

Kožni multiprocesor može se nazvati zajedničkom memorijom.
Njezino proširenje postaje 16 kilobajta i daje programeru potpunu slobodu
uradi sam. Razpodílya yak želite :). Zajednička memorija osigurava strujanje poziva
jedan blok nije prepoznat za rad s pixel shaderima.

Također SM može pretvoriti u GDDR. Za koje su “prišili” 8 kilobajta
cache memorija, koja sprema sve naslove za robote (na primjer, nabrajanje
konstante).

Multiprocesor max 8192 registra. Broj aktivnih blokova ne može biti
više od osam, a broj zavoja - ne više od 768/32 = 24. Može se vidjeti da je G80
Možete obraditi maksimalno 32*16*24 = 12,288 streamova po satu. Ne mogu ne
ispravite ove brojke prilikom optimizacije programa nadal (na jednu zdjelu vag
- Razmír blok, na ínshiy - broj tokova). Ravnoteža parametara se može mijenjati
tome je data važna uloga NVIDIA preporučiti vicoristing blokove
sí 128 ili 256 tokova. Block z 512 tokovi su neučinkoviti, krhotine mogu
kreće zatrims. Vrakhovuychi sve tanke GPU video kartice plus
loše vještine programiranja, možete stvoriti produktivniji
zasib za paralelne proračune. Prije govora, o programiranju...

programiranje

Za "kreativnost" odjednom iz CUDA-e potrebno je GeForce video kartica nije niža
osma serija. W

Za službenu stranicu potrebna su tri softverska paketa: upravljački program
podržava CUDA (za skin OS - vlastiti) bez posredničkog CUDA SDK paketa (ostalo
beta verzija) i dodatne biblioteke (CUDA toolkit). Tehnološka podrška
Operativni sustavi Windows (XP i Vista), Linux i Mac OS X.
odabirom Vista Ultimate Edition x64 (reći ću da je sustav bio
upravo nevjerojatno). U vrijeme pisanja ovi su redovi relevantni za robote
drajver za ForceWare 177.35. Kao set vicorističkih alata
programski paket Borland C++ 6 Builder
moj C).

Ljudi se, koliko znam, lako naviknu na novu sredinu. Manje potrebno
zapamtiti glavne parametre. _global_ ključna riječ (prije funkcije)
Pokazuje da funkcija treba biti prije kernela. í̈ viklikatime središnji
procesor, a cijeli robot će biti na GPU-u. Wiklik _global_ čini više
konkretni detalji, ali širenje same mreže, proširenje bloka i jezgre će biti
zapeo. Na primjer, red _global_ void saxpy_parallel<<>>, od X –
veličina mreže, a Y - veličina bloka, koja određuje broj parametara.

Simbol _uređaj_ znači da funkciju poziva grafička jezgra, ali
slijedite ove upute. Ova funkcija je pohranjena u memoriji multiprocesora,
otzhe, otrimati njezinu adresu je nemoguće. Prefiks _host_ znači da wiki jest
da obrobka prođe manje za sudbinu CPU-a. Potrebno je jamčiti za ono što _globalno_ i
_uređaj_ ne može pozvati jedan od jedan i ne može pozvati sam sebe.

Također jezik za CUDA može imati niske funkcije za rad s video memorijom: cudafree
(promjena memorije između GDDR-a i RAM-a), cudamemcpy i cudamemcpy2D (kopija
memorije između GDDR-a i RAM-a) i cudamalloc (prikaz memorije).

Svi kodni programi moraju biti kompajlirani sa strane CUDA API-ja. Na klipu se uzima
kod, uključujući dodjele za središnji procesor, i poddaetsya
standardna kompilacija i drugi kod, dodjele za grafički adapter,
prepisati u PTX (vrlo pogodan asembler) za
pokazujući moguća pomilovanja. Nakon svih ovih “plesova” ostaje ostatak
prijevod (emitiranje) naredbi u razumijevanju za GPU/CPU mov.

Brojčanik za vjenčanje

Gotovo svi aspekti programiranja opisani su u dokumentaciji tj
zajedno s vozačem i dva programa, kao i na stranicama trgovaca. Rozmir
ne čitajte članke, da biste ih opisali (čitanje se može dodati
malo truda i samostalnog rada na gradivu).

CUDA SDK Browser je proširen posebno za početnike. Be-yaky bazhayuchy može
osjetite snagu paralelnih izračuna na vlastitoj koži (bolja ponovna provjera na
stabilnost - robot se primjenjuje bez artefakata i vills). Dodatak može
veliki broj show mini-programa (61 "test"). Na kožu dosvídu ê
Dokumentacija izvješća o programskom kodu plus PDF datoteke. Vidi se da ljudi
prisutni sa svojim prozorima u pregledniku, baveći se ozbiljnim poslom.
Ovdje također možete promijeniti brzinu robota i procesora i video kartice za sat obrade
danich. Na primjer, skeniranje bogatih nizova s ​​video karticom GeForce 8800
GT 512 MB s blokom od 256 niti kruži za 0,17109 milisekundi.
Tehnologija ne prepoznaje SLI tandeme, pa imate trio koji puše,
uključite funkciju "uparivanja" prije robota, inače će CUDA raditi samo jedan
uređaj. dvostruka jezgra AMD Athlon 64X2(frekvencija jezgre 3000 MHz) isto odobrenje
prolaze za 2,761528 milisekundi. Navedite da je G92 veći niži od 16 puta
shvidshe "kamen" AMD! Yak Bachish, sustav je daleko od toga da bude ekstreman
tandem s nevoljenim u masama operativnim sustavom koji prikazuje loše
rezultate.

Krím preglednik ísnuê niske korisnyh suspílstvo programa. Adobe
prilagodio svoje proizvode novoj tehnologiji. Sada je Photoshop CS4 ažuriran
svjetski pobjednički resursi grafičkih adaptera (potrebno je osigurati posebne
uključiti). S takvim programima, kao što su Badaboom media converter i RapiHD, možete
Mogućnost dekodiranja videa u MPEG-2 format. Za obradu zvuk je loš
pidide besplatni uslužni program Accelero. Broj softvera, izoštreni CUDA API,
naravno, nastavi rasti.

A u to vrijeme...

I dok čitate ovaj materijal, vrijedni radnici iz procesora se brinu
proširiti svoje tehnologije kako bi promovirali GPU umjesto CPU-a. 3 strane AMD brkovi
razumio: imaju veličanstven dosvíd, nabutiy ujedno iz ATI.

Stvaranje "mikrouređaja", Fusion, sastoji se od niza jezgri
kodnog naziva Bulldozer i videočip RV710 (Kong). Vaš će međusobni odnos biti
zdíysnyuvatsya za rahunok razrezheny guma HyperTransport. Pogled ugar
broj jezgri i frekvencijske karakteristike AMD planira stvoriti cijelu cijenu
hijerarhija "kamenja". Također je planirana konverzija procesora kao za prijenosna računala (Falcon),
i za multimedijske gadgete (Bobcat). Štoviše, sama stagnacija tehnologije
prijenosne gospodarske zgrade bit će prvi zadatak za Kanađane. Z odbor za razvoj
Paralelno nabrajanje stagnacije takvih "kamintsiv" može biti još popularnije.

Intel Tri na sat za vaš Larrabee. Proizvodi AMD,
yakscho da se ne napijete, da se pojavite na policama trgovina kao 2009 - na klipu
rock iz 2010. A odluka protivnika da vidi svjetlo Božje manje je vjerojatna u dvoje
sudbina.

Larrabee ima veliki broj (čitaj - stotine) jezgri. Na klipu
Pa, možete vidjeti robu, otplaćenu za 8 - 64 jezgre. Smrad više liči na Pentium, ale
do mlijeka je jako premoreno. Skin kernel može imati 256 kilobajta predmemorije drugog jednakog
(Uz godinu dana joge, rast će se povećati). Vzaymosv'yazok
1024-bitna dvosmjerna prstenasta sabirnica. Čini se da je Intel "dijete"
dobra praksa s DirectX i Open GL API (za "Yabluchnikov"), ništa od toga
softverski unos nije potreban.

I zašto sam sav tse rozpoviv? Očito je da se Larrabee i Fusion ne mogu vidjeti
fiksni, stacionarni procesori s tržišta
video kartice. Za igrače i ekstremne ljude, između snova, kao prije, kojih se treba riješiti
bogati CPU i tandem s vrhunskim VGA naljepnicama. Ale oni koji navit
procesorske tvrtke da prijeđu na paralelnu naplatu za načela,
analogno GPGPU, govoriti već bogato o čemu. Zokrema o onima koji jesu
tehnologija, poput CUDA-e, može imati pravo koristiti i, možda, biti
već popularan.

Mali životopis

Paralelna naplata uz pomoć video kartice samo je dobar alat
u rukama robotskog programera. Teško chi procesori na choli íz Mooreov zakon
Stop. Tvrtke NVIDIA leći ići još jedan dugačak put
uklizavanje u masu vašeg API-ja (isto se može reći i za dijete ATI/AMD).
Kakav ćeš biti, pokaži budućnost. Također, CUDA će se vratiti :).

p.s. Za sve programere i ljude koji su zapeli, preporučam da pogledaju
sljedeća "virtualna hipoteka":

NVIDIA službena web stranica i web stranica
GPGPU.com. svi
informacije su dane - moj engleski, ali hvala, volio bih da nije
Kineski. Zato samo tako nastavi! Siguran sam da vam je autor htio malo pomoći
sijede inicijacije CUDA znanja!

Í dodjele za prevođenje u host-code (head, key code) i device-code (hardware code) (datoteke s ekstenzijama.cu) u objektne datoteke, pribor u procesu kompajliranja konačnog programa ili biblioteke u bilo kojem programskom okruženju, na primjer NetBeans.

U CUDA arhitekturi pobjeđuju memorijski model grida, klasterski model tokova i SIMD instrukcije. Ne postoje samo grafičke kartice visokih performansi, već i drugi znanstveni računi za nVidia video kartice. Rezultati ovog istraživanja naširoko se koriste za CUDA-u u različitim područjima, uključujući astrofiziku, proračune biologije i kemije, modeliranje dinamike rijeka, elektromagnetske interakcije, računalnu tomografiju, seizmičku analizu i više. CUDA ima mogućnost spajanja na programe koji koriste OpenGL i Direct3D. CUDA je višeplatformski softver za operativne sustave kao što su Linux, Mac OS X i Windows.

Dana 22. ožujka 2010. nVidia je izdala CUDA Toolkit 3.0, koji predstavlja oživljavanje OpenCL podrške.

Vlasništvo

CUDA platforma prvi put se pojavila na tržištu izlaskom osme generacije NVIDIA čipa G80 i postala prisutna u svim nadolazećim serijama grafičkih čipova, koji se nalaze u obiteljima brzorastućih GeForce, Quadro i NVidia Tesla.

Prvi u nizu koji podržava CUDA SDK, G8x, mali je 32-bitni vektorski procesor s jednom preciznošću koji podržava CUDA SDK kao API (CUDA podržava dvostruki filmski tip Ci, proteo preciznost smanjena je na 32- bit pomičnog zareza). Niži GT200 procesori mogu podržavati 64-bitnu preciznost (samo za SFU), ali performanse su znatno veće, niže za 32-bitnu preciznost (preko njih postoje samo dva SFU-a po skin streaming multiprocesoru, a skalarni procesori - svi). Grafički procesor organizira hardversko bogatstvo toka, što vam omogućuje iskorištavanje svih resursa grafičkog procesora. Na taj način se pojavljuje mogućnost prebacivanja funkcija fizičkog prečaca na grafički (primjer implementacije - nVidia PhysX). Također, postoje široke mogućnosti korištenja grafičke postavke računala za sklopivi negrafički izračun: na primjer, u izračunu biologije i drugim znanstvenim trikovima.

Perevagi

Naizmjence s tradicionalnim pristupom organiziranju izračuna globalnog priznanja za dodatne mogućnosti grafičkih API-ja, CUDA arhitektura također može pobijediti u ovoj galeriji:

Razmjena

• Sve funkcije, koje su prikazane na proširenju, ne podržavaju rekurziju (u verziji CUDA Toolkit 3.1, podržavaju indikatore te rekurzije) i mogu raditi na drugim razmjenama

GPU podešavanja i grafička poboljšanja

Pretvorba dodataka u obliku Nvidia-omogućene verzije deklarirane nove podrške za CUDA tehnologiju usmjerena je na službenu Nvidia web stranicu: CUDA-Enabled GPU Products (engleski).

Naime, u ovom trenutku na tržištu hardvera za osobna računala podrška CUDA tehnologiji bit će osigurana naprednim perifernim uređajima:

Verzija specifikacije	GPU	Video kartice
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, G94b, G94b, G92b, 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/30 /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, GTX140 500
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, GForce GTX 680MX, GeForce GTX 675MX, GeFor 6 GeForce GT 645M, GeForce GT 640M
3.5	GK110

Nvidia GeForce za stolna računala 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 9400mGPU GeForce 9300mGPU GeForce 8800 GTS 512 GeForce 8800 GT GeForce 8600 GTS GeForce 8600 GT GeForce 8500 GT GeForce 8400GS

Nvidia GeForce za mobilna računala 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 9650M GS GeForce 9600M GT GeForce 9600M GS GeForce 9500M GS GeForce 9500M G GeForce 9300M GS GeForce 9300M G GeForce 9200M GS GeForce 9100MG GeForce 8800M GTS GeForce 8700M GT GeForce 8600M GT GeForce 8600M GS GeForce 8400M GT GeForce 8400M GS

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

Nvidia Quadro za stolna računala Quadro 6000 Quadro 5000 Quadro 4000 Quadro 2000 Quadro 600 QuadroFX 5800 QuadroFX 5600 QuadroFX4800 Quadro FX 4700X2 QuadroFX4600 QuadroFX 3700 QuadroFX 1700 QuadroFX 570 QuadroFX470 Quadro FX 380 niskog profila QuadroFX 370 Quadro FX 370 niskog profila Quadro CX Quadro NVS450 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 za mobilna računala Quadro 5010M Quadro 5000M Quadro 4000M Quadro 3000M Quadro 2000M Quadro 1000M QuadroFX 3800M QuadroFX 3700M QuadroFX 3600M QuadroFX 2800M QuadroFX 2700M QuadroFX 1800M QuadroFX 1700M QuadroFX 1600M QuadroFX 880M QuadroFX 770M QuadroFX 570M QuadroFX 380M QuadroFX 370M QuadroFX 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

• Modeli Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 omogućuju izvođenje izračuna na GPU-u s promjenjivom točnošću.

Značajke i specifikacije različitih verzija

Podrška za značajke (nenavedene značajke su podržano za sve računalne mogućnosti)	Računalna sposobnost (verzija)
	1.0	1.1	1.2	1.3	2.x
32-bitne riječi u globalnoj memoriji	bok	Tako
vrijednosti pomičnog zareza u globalnoj memoriji
Cjelobrojne atomske funkcije koje djeluju na 32-bitne riječi u zajedničkoj memoriji	bok		Tako
atomicExch() radi na 32-bitnom vrijednosti pomičnog zareza u zajedničkoj memoriji
Cjelobrojne atomske funkcije koje djeluju na 64-bitne riječi u globalnoj memoriji
Warp glasovne funkcije
Operacije s pomičnim zarezom dvostruke preciznosti	bok			Tako
Atomske funkcije koje rade na 64-bitu cjelobrojne vrijednosti u zajedničkoj memoriji	bok				Tako
Djeluje atomsko zbrajanje s pomičnim zarezom 32-bitne riječi u globalnoj i zajedničkoj memoriji
_glasački listić()
_threadfence_system()
_syncthreads_count(), _syncthreads_and(), _syncthreads_or()
Površinske funkcije
3D mreža bloka niti

Tehničke specifikacije	Računalna sposobnost (verzija)
	1.0	1.1	1.2	1.3	2.x
Maksimalna dimenzija mreže blokova niti	2				3
Maksimalna x-, y- ili z-dimenzija mreže blokova niti	65535
Maksimalna dimenzija bloka navoja	3
Maksimalna x- ili y-dimenzija bloka	512				1024
Maksimalna z-dimenzija bloka	64
Maksimalan broj niti po bloku	512				1024
Veličina osnove	32
Maksimalan broj rezidentnih blokova po multiprocesoru	8
Maksimalni broj rezidentnih warps multiprocesora	24		32		48
Maksimalan broj rezidentnih niti po multiprocesoru	768		1024		1536
Broj 32-bitnih registara za multiprocesor	8K		16K		32K
Maksimalna količina zajedničke memorije po multiprocesoru	16 KB				48 KB
Broj zajedničkih memorijskih banaka	16				32
Količina lokalne memorije po niti	16 KB				512 KB
Konstantna veličina memorije	64 KB
Radni skup predmemorije po multiprocesoru za stalnu memoriju	8 KB
Radni skup predmemorije po multiprocesoru za memoriju teksture	Ovisno o uređaju, između 6 KB i 8 KB
Maksimalna širina za 1D teksturu	8192				32768
Maksimalna širina za 1D teksturu referenca vezana za linearnu memoriju	2 27
Maksimalna širina i broj slojeva za referencu 1D slojevite teksture	8192x512				16384x2048
Maksimalna širina i visina za 2D referenca teksture vezana za linearna memorija ili CUDA polje	65536 x 32768				65536 x 65535
Najveća širina, visina i broj slojeva za 2D slojevitu referencu teksture	8192 x 8192 x 512				16384 x 16384 x 2048
Maksimalna širina, visina i dubina za referencu 3D teksture vezanu za linearnu memorije ili CUDA polja	2048x2048x2048
Maksimalan broj tekstura koje može biti vezan za kernel	128
Maksimalna širina za 1D površinu referenca vezana za CUDA polje	Ne podržan				8192
Maksimalna širina i visina za 2D površinska referenca vezana za CUDA polje					8192 x 8192
Maksimalan broj površina koje može biti vezan za kernel					8
Maksimalan broj uputa po zrno	2 miliona

stražnjica

CudaArray* cu_array; tekstura< float , 2 >tex; // Dodijeli niz cudaMalloc( & cu_array, cudaCreateChannelDesc< float>(), širina Visina); // Kopiraj slikovne podatke u niz cudaMemcpy( cu_array, image, width* height, cudaMemcpyHostToDevice) ; // Povežite polje s teksturom cudaBindTexture(tex, cu_array) ; // Pokreni kernel dim3 blockDim(16, 16, 1); dim3 gridDim(širina/blokDim.x, visina/blokDim.y, 1); zrno<<< gridDim, blockDim, 0 >>> (d_podaci, širina, visina); cudaUnbindTexture(tex); __global__ void kernel(float * odata, int visina, int širina) ( unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float texfetch( tex, x, y);odata[y* širina+ x] = c;

Uvezi pycuda.driver poput drv import numpy drv.init() dev = drv.Device(0) ctx=dev.make_context() mod=drv.SourceModule( """ __globalno__ 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 (a) multiply_them( drv .Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

CUDA kao predmet na sveučilištima

Dolazeći početkom 2009. godine, softverski model CUDA postavljen je na 269 sveučilišta diljem svijeta. U Rusiji se prvi tečajevi o CUDA-i predaju na Politehničkom sveučilištu u Sankt Peterburgu, Državnom sveučilištu u Jaroslavlju nazvanom po. P. G. Demidov, Moskovski, Nižnji Novgorod, St. Petersburg, Tverskoye, Kazan, Novosibirsky, Novosibirsky Holden Tehničko sveučilište Permskog UNIVITITIES, the MITENENTY UNIVITITITITIONALY ONISTITITIONALS. Bauman, RKhTU im. Mendeljev, Međuregionalni centar za superračunala Ruske akademije znanosti, . Osim toga, početkom 2009. godine najavljen je početak rada prvog znanstvenog i rasvjetnog centra u Rusiji "Paralelni proračun", koji je uređen u gradu Dubni.

U Ukrajini se tečajevi o CUDA-i čitaju na Kijevskom institutu za analizu sustava.

Posilannya

Službeni resursi

• CUDA zona (ruski) - službena CUDA web stranica
• CUDA GPU Computing (engleski) - službeni web forumi posvećeni CUDA izračunima

Neslužbeni izvori

Tomov hardver
• Dmitro Čekanov. nVidia CUDA: troškovi na video karticama ili smrt CPU-a? . Tom's Hardware (22. rujna 2008.).
• Dmitro Čekanov. nVidia CUDA: GPU testiranje za mainstream tržište. Tom's Hardware (utorak, 19. 2009.).
iXBT.com
• Oleksij Berillo. NVIDIA CUDA - negrafičko računalstvo na grafičkim procesorima. 1. dio. iXBT.com (23. rujna 2008.). Arhivirano prošlog 4. veljače 2012. Revidirano 20. rujna 2009.
• Oleksij Berillo. NVIDIA CUDA - negrafičko računalstvo na grafičkim procesorima. 2. dio. iXBT.com (22. srpnja 2008.). - Primijenite NVIDIA CUDA zakrpu. Arhivirano prošlog 4. veljače 2012. Revidirano 20. rujna 2009.
Ostali resursi
• Boreskov Oleksij Viktorovič. Osnove CUDA-e (20. rujna 2009.). Arhivirano prošlog 4. veljače 2012. Revidirano 20. rujna 2009.
• Volodimir Frolov. Uvod u CUDA tehnologiju. Časopis Merezhevy "Računalna grafika i multimedija" (19. prosinca 2008.). Arhivirano prošlog 4. veljače 2012. Revidirano 28. lipnja 2009.
• Igor Oskolkov. NVIDIA CUDA pristupačna je ulaznica u svijet velikih brojeva. Računala (30. travnja 2009.). Preuzeto 3. svibnja 2009.
• Volodimir Frolov. Uvod u CUDA tehnologiju (1. rujna 2009.). Arhivirano prošlog 4. veljače 2012. Revidirano 3. travnja 2010.
• GPGPU.ru. Wikoristannya video kartice za izračun
• . Paralelni centar za izračun

Bilješke

div. također

Nvidia
Grafički procesori	Rano NV1 NV2 Obitelj RIVA TNT TNT2 Obitelj