nvidia cuda technologija. NVidia CUDA: GPU mokesčiai ir procesoriaus mirtis? Duomenų migracijos sumažinimo pagrindinio kompiuterio įrenginys

Tai inshi. Tačiau „CUDA scan“ derinio paieška, pamačius tik 2 straipsnius, niekaip neatitinka nuskaitymo algoritmo GPU – tačiau tai yra vienas iš pagrindinių algoritmų. Taigi, stipriai atsikvėpę, pažvelkime į „Udacity“ kursą – paralelinio programavimo įvadas ir aš drįstu parašyti naują straipsnių ciklą apie CUDA. Dar kartą pasakysiu, kad serialas bus pagrįstas sava eiga, ir net jei turėsi valandą, jogą išgyvensi turtingiau. Šiuo metu planuojami šie straipsniai:
1 dalis: nuostata.
2 dalis: GPU aparatinė įranga ir lygiagrečios komunikacijos modeliai.
3 dalis: Pagrindiniai GPU algoritmai: sumažinimas, nuskaitymas ir histograma.
4 dalis: Pagrindiniai GPU algoritmai: kompaktiškas, segmentuotas nuskaitymas, rūšiavimas. Praktinis kai kurių algoritmų įgyvendinimas.
5 dalis: GPU programinės įrangos optimizavimas.
6 dalis. Vėlesnių algoritmų lygiagretinimo pavyzdys.
7 dalis: Papildomos paralelinio programavimo temos, dinaminis paralelizmas.

Zatrimka prieš statybos leidimą

Visų pirma, ar galite įdėti odą prieš GPU užšaldymą, kad atliktumėte savo užduotis – ir tokiems tikslams geras GPU, jei reikia jį užšaldyti? Vidpovidі sіd reiškia 2 sąvokas:
Zatrimka(latency) – valanda, per kurią atliekama viena instrukcija/operacija.
Pralaidumas- instrukcijų/operacijų skaičius, kuris skaičiuojamas per valandą.
Paprastas pavyzdys: gal lengvasis automobilis, kurio greitis 90 km/met ir 4 asmenys, ir autobusas, kurio galia 60 km/met ir 20 žmonių. Jei operacijai priimti 1 asmens judėjimą 1 kilometre, tai lengvojo automobilio vėlavimas yra 3600/90 = 40 s - kelioms sekundėms 1 žmogus yra pusiaukelėje iki 1 kilometro, automobilio pralaidumas yra 4/ 40 = 0,1 operacijos per sekundę; magistralės vėlavimas - 3600/60=60s, magistralės pralaidumas - 20/60=0,3(3) operacijos/sek.
Taigi iš procesoriaus - tse automobilio, GPU - magistralės: gali būti didelis užsikimšimas, bet ir didelis pralaidumo pastatas. Kalbant apie jūsų konkrečios odos operacijos laiką, tai nėra taip svarbu, kaip šių operacijų skaičius per sekundę – tiesiog pažiūrėkite į GPU.

Pagrindinis CUDA supratimas ir sąlygos

• Priedas (įrenginys)- GPU. „Underdog“ vaidmuo yra pergalingas - tie, kurie netgi gali naudoti centrinį procesorių, mažiau dirbs.
• Priegloba (host)- CPU. Vikonu vaidina vaidmenį - paleidžia užduotį priede, mato atmintį priede, perkelia atmintį / iš priedo. Taigi, CUDA skanduotė sako, kad tai yra priedas, todėl šeimininkas gali prisiminti savo atmintį.
• Branduolys- Užduotis, kurią šeimininkas pradeda ūkiniame pastate.

Pagrindiniai CUDA programų etapai

1. Šeimininkas priede mato reikiamą atminties kiekį.
2. Priegloba nukopijuos duomenis iš mano atminties į mano atmintį.
3. Šeimininkas pradeda dainuojančius branduolius ūkiniame pastate.
4. Pritvirtinkite branduolius.
5. Prisegsiu rezultatų kopiją iš atminties į savo atmintį.

Branduoliai

1. Vidurinės šerdys gali atpažinti „identifikatorių“ arba, paprasčiau, regis, srauto, kuris yra užkrečiamai pergalingas, padėtį - pergalinga padėtimi pasiekiame, kad viena šerdis yra praktiškai įmanoma su skirtingu džinsinio pūdymu sraute, kurį jis veikia. Prieš kalbą toks lygiagrečių skaičiavimų organizavimas vadinamas SIMD (Single Instruction Multiple Data) – jei keli procesoriai vienu metu atlieka tą pačią operaciją, nors ir skirtingais duomenimis.
2. Kai kuriais atvejais branduolio kodus reikia pakeisti skirtingus sinchronizavimo būdus.

• Galinėje pusėje nustatyti vadinamojo tinklelio (tinklelio) matmenys 3D koordinatėmis: tinklelis_x, tinklelis_y, tinklelis_z. Dėl to tinklas susidaro iš tinklelis_x*tinklelis_y*tinklelis_z blokai.
• Tada bloko dydis pakeičiamas 3D koordinatėmis: blokas_x, blokas_y, blokas_z. Dėl to blokas bus sulankstytas blokas_x*blokas_y*blokas_z srautai. Tėve, prašau grid_x*grid_y*grid_z*block_x*block_y*block_z srautai. Svarbi pagarba – didžiausias srautų skaičius viename bloke yra apribotas ir deponuojamas kiekvienam GPU modeliui – tipinės reikšmės yra 512 (senesni modeliai) ir 1024 (naujesni modeliai).
• Galimas vidurinės šerdies pakeitimas threadIdxі blockIdx su laukais x, y, z- smirdi, kad atkeršytų 3D koordinates srautui bloke ir blokui svetainėje aišku. Taip pat galimi pakeitimai blockDimі tinklelisDim su tais pačiais laukais - pakeiskite bloko dydį ir tinklelis bus aiškus.

Rašome programą CUDA

pagrindinis.cpp

#įtraukti #įtraukti #įtraukti #įtraukti #įtraukti #įtraukti #įtraukti #įtraukti #include "openMP.hpp" #include "CUDA_wrappers.hpp" #include "common/image_helpers.hpp" naudojant vardų erdvę cv; naudojant vardų sritį std; int main(int argc, char** argv) ( naudojant vardų erdvę 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(sistemos_laikrodis::now() - start); cout<<"OpenMP time (ms):" << duration.count() << endl; memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols); RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols); return 0; }

paruoštiImagePointers

šabloną void readyImagePointers(const char * const inputImageFileName, cv::Mat& inputImage, T1** inputImageArray, cv::Mat& outputImage, T2** outputImageArray, const int outputImageType) ( naudojant vardų erdvę imm 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); *outputImageArray = (T2*)outputImage.ptr (0); }

openMP.hpp

#įtraukti #įtraukti #įtraukti void RGBtoGrayscaleOpenMP(uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) ( #pragma omp paralelė sutraukimui(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; } } }

Priedo tekstas

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 numerio paleisties atmintis įjungimo , sizeof(uchar4) * numPixels)); GPU checkCudaErrors(cudaMemcpy(d_imageRGBA, h_imageRGBA, sizeof(uchar4) * pikselių skaičius, cudaMemcpyHostToDevice));

Priedo tekstas

#įtraukti #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__) šablonas negaliojantis patikrinimas (T err, const char * const func, const char * const failas, const int eilutė) ( if (err! = cudaSuccess) ( std::cerr<< "CUDA error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString(err) << " " << func << std::endl; exit(1); } }

Priedo tekstas

dim3 bloko dydis; dim3 tinklelio dydis; int threadNum; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); sriegio skaičius = 1024; bloko dydis = dim3 (sriegio skaičius, 1, 1); tinklelio dydis = dim3 (stulpelių skaičius / gijosSkaičius + 1, eilučių skaičius, 1); cudaEventRecord(pradžia); rgba_to_greyscale_simple<<>>(d_imageRGBA, d_imageGray, numRows, numCols); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); plaukiojimo milisekundės = 0; cudaEventElapsedTime(&milisekundės, pradžia, pabaiga); std::out<< "CUDA time simple (ms): " << milliseconds << std::endl;

rgba_to_greyscale_simple

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 .x | >=eilučių skaičius) grąžinimas, const int poslinkis = y*numCols+x; const uchar4 pikselis = d_imageRGBA; 0,114f*pikselis.z;

Priedo tekstas

checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(nesigned char) * pikselių skaičius, cudaMemcpyDeviceToHost)); cudaFree(d_imageGray); cudaFree(d_imageRGBA);

Priedo tekstas

Procesorius: Intel Core(TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4 GB, 1600 MHz.
OS: OS X 10.9.5.
Kompiliatorius: g++ (GCC) 4.9.2 20141029.
CUDA kompiliatorius: Cuda kompiliavimo įrankiai, 6.0, V6.0.1 leidimas.
Palaikoma OpenMP versija: OpenMP 4.0.

rgba_to_greyscale_optimized

32 WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x;for (int i=loop_start, j=0; j

Priedo tekstas

gijosNum=128; const int elemsPerThread = 16; bloko dydis = dim3 (sriegio skaičius, 1, 1); tinklelio dydis = dim3(SkaičiusCools/(TreadNum*elementsPer Thread) + 1, eilučių skaičius, 1); cudaEventRecord(pradžia); rgba_to_greyscale_optimized<<>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); milisekundės = 0; cudaEventElapsedTime(&milisekundės, pradžia, pabaiga); std::out<< "CUDA time optimized (ms): " << milliseconds << std::endl;

Žymos: pridėti žymų

- Mažo programavimo sąsajų rinkinys ( API) igor ir kitų labai produktyvių multimedijos priedų kūrimui. Įjunkite didelio našumo apdailą 2D- І 3D- grafika, garsas ir įžanga.

Direct3D (D3D) - sąsaja triviworlds peržiūrai primityvai(geometriniai kūnai). Įeikite .

OpenGL(anglų kalba. Atidarykite grafikos biblioteką, Pažodžiui – atvira grafinė biblioteka) – specifikacija, apibrėžianti nepriklausomą kelių platformų programuojamos programinės įrangos sąsajos tipą, skirtą rašyti priedus, kurie gali būti naudojami kuriant dvimatę ir trimatę kompiuterinę grafiką. Apima daugiau nei 250 funkcijų, leidžiančių piešti sulankstomas trijų pasaulių scenas iš paprastų primityvų. Pergalė už vaizdo kalnų kūrimą, virtualią realybę, vizualizaciją moksliniuose tyrimuose. Ant platformos Windows konkuruoti su .

OpenCL(anglų kalba. Atidarykite „Computing Language“., pažodžiui - vodkrita mov skaičiavimas) - sistema(Programinės įrangos sistemos karkasas), skirtas rašyti kompiuterių programas, sujungtas su lygiagrečiais įkrovimais įvairiose grafikose ( GPU) ir ( ). U karkasas OpenCLįveskite mov programavimo ir sąsajos programavimo priedus ( API). OpenCL užtikrinti lygiagretumą nurodymų ir duomenų lygmeniu bei technologijų diegimu GPGPU.

GPGPU(trumpa anglų kalba) Bendrosios paskirties G raphics P ranking Units, pažodžiui - GPU visuotinis atpažinimas) - vaizdo plokštės grafinio procesoriaus naudojimo pasauliniams skaičiavimams technika, kaip skamba atlikti.

šešėliuotojas(angl. šešėliuotojas) - programa, skirta susintetintuose vaizduose sukelti šešėlius, išryškinti trivialią vaizdo objekto liekamųjų parametrų grafiką. Paprastai tai apima tam tikrą lankstymo laipsnį, molio ir rožių šviesos aprašymą, tekstūros perdengimą, sulankstymą, šešėliavimą, paviršiaus poslinkį ir tolesnio apdorojimo efektus. Paprastų geometrinių formų pagalba galima vizualizuoti sulankstomus paviršius.

perteikimas(angl. perteikimas) - vizualizacija, kompiuterinėje grafikoje, modelio vaizdo paėmimo procesas programinės įrangos pagalba.

SDK(trumpa anglų kalba) Programinės įrangos kūrimo rinkinys) – Programinės įrangos saugos kūrimo įrankių rinkinys.

CPU(trumpa anglų kalba) Centrinis apdorojimo blokas, Pažodžiui - centrinis / pagrindinis / galvos skaičiavimo priedas) - centrinis (mikro); priedas, kurio vikonu mašinos instrukcijos; aparatinės įrangos saugumo dalis, kuri yra atsakinga už operacijų (operacinės sistemos ir taikomosios programinės įrangos užduočių) apskaičiavimą ir koordinuoja visų priedų darbą.

GPU(trumpa anglų kalba) Grafikos apdorojimo blokas, Pažodžiui - grafinio skaičiavimo priedas) - grafinis procesorius; okremy priedai ar žaidimų konsolės, kurių vikonu grafinis atvaizdavimas (vizualizacija). Šiuolaikiniai grafikos procesoriai gali efektyviai apdoroti ir tikroviškai pateikti kompiuterinę grafiką. Grafikos procesorius šiandieniniuose vaizdo adapteriuose zastosovuetsya kaip trivivirnoy grafikos tvirtinimo elementas, proteo yogo gali būti susuktas kai kuriais režimais ir skaičiavimui ( GPGPU).

problemų CPU

Ilgą laiką tradicinių produktyvumo padidėjimą labiausiai lėmė vėlesnis laikrodžio dažnio padidėjimas (apie 80% paties laikrodžio dažnio pradžios produktyvumo), viena valanda padidėjus laikrodžių skaičiui. tranzistoriai viename kristale. Tačiau toliau, takto dažnio padidėjimas (kai takto dažnis didesnis nei 3,8 GHz, lustai tiesiog perkaista!) susiduria su daugybe esminių fizinių kliūčių (technologinio proceso šukės galėjo priartėti prie atomo išsiplėtimo): , O silicio atomo dydis yra maždaug 0,543 nm):

Pirma, norėdami pakeisti kristalų plėtimąsi ir padidinti laikrodžio dažnį, padidinkite tranzistorių srautą. Tse vede iki sumažėjusios įtampos padidėjimo ir šilumos praradimo padidėjimo;

Kitu būdu pagrindinio laikrodžio dažnio didinimas dažnai varomas per strigimus perėjimo į atmintį metu, todėl prieiga prie atminties neatitinka didėjančių laikrodžio dažnių;

Trečia, kai kurie tradicinių nuoseklių architektūrų papildymai tampa neveiksmingi dėl didėjančio laikrodžio dažnio per vadinamąją „von Neumanno aukštąją mokyklą“ – našumo sumažėjimą dėl nuoseklaus skaičiavimo srauto. Kuriems yra varžinis-talpinis signalo perdavimo trukdymas arba su papildoma siaura erdve, arba padidinus laikrodžio dažnį.

Rosvitokas GPU

Tuo pačiu metu vystosi GPU:

…

Lapų kritimas 2008 m - Intel pristatė 4 branduolių liniją Intel Core i7, kurios yra pagrįstos naujos kartos mikroarchitektūra Nehalem. Procesoriai veikia 26-32 GHz taktiniu dažniu. Vikonani už 45 nm proceso technologiją.

Krūtinė 2008 m – rozpochalis 4 branduolių pristatymai AMD Phenom II 940(Kodinis pavadinimas - Deneb). Gaminamas 3 GHz dažniu, pagamintas 45 nm proceso technologijai.

…

„Traven“ 2009 m - įmonė AMD pristatė grafikos procesoriaus versiją ATI Radeon HD 4890 nuo šerdies laikrodžio dažnis padidėjo nuo 850 MHz iki 1 GHz. Tse pirma grafinis procesorius, veikiantis 1 GHz dažniu. Lusto padermių skaičius ir dažnio padidėjimas padidėjo nuo 1,36 iki 1,6 teraflopo. Procesorius, skirtas pakeisti 800 (!) skaičiuojančių branduolių, palaiko vaizdo atmintį GDDR5, „DirectX 10.1“., ATI CrossFireX ir visos kitos technologijos, maitinančios šiuolaikinius vaizdo plokščių modelius. Lustų preparatai 55 nm technologijos pagrindu.

Pagrindinės galios GPU

Vidminnimi ryžiai GPU(pataisyta iš ) є:

- architektūra, maksimaliai nukreipta į tekstūrų lankstumo didinimą ir grafinių objektų lankstymą;

- tipinis didžiausias slėgis GPU turtingesnis, žemesnis ;

– specialiosios konvejerio architektūros lyderiai, GPU efektyvesnis apdorojant grafinę informaciją, mažesnis.

„Žanro krizė“

„Žanro krizė“ už brandino iki 2005 m., – atsirado jie patys. Ale, nepaisant technologijų plėtros, produktyvumo augimas pastebimai sumažėjo. Vandens valandos produktyvumas GPU toliau augti. Taip, iki 2003 m. ir revoliucinė idėja išsikristalizavo - vikoristovuvat už mіts grafinio skaičiavimo poreikius. Grafiniai procesoriai tapo aktyviai naudojami „negrafiniams“ skaičiavimams (fizikos modeliavimas, signalų apdorojimas, skaičiavimo matematika/geometrija, operacijos su duomenų bazėmis, skaičiavimo biologija, skaičiavimo ekonomika, informatika).

Pagrindinė problema yra ta, kad nebuvo standartinės sąsajos programavimui GPU. Rozrobniki vikoristovuvali OpenGL arba Direct3D bet buvo per sunku. Korporacija NVIDIA(viena didžiausių grafikos, medijos ir komunikacijos procesorių, taip pat besmiginių medijos procesorių kolekcijų; įkurta 1993 m.) užsiėmė vieno rankinio standarto kūrimu – ir pristatė technologiją. CUDA.

Kaip tai prasidėjo

2006 m - NVIDIA demonstruodamas CUDA™; skaičiavimo revoliucijos pradžia GPU.

2007 m - NVIDIA išleidimo architektūra CUDA(Pochatovo versija CUDA SDK bula pristatyta 2007 m. vasario 15 d.); žurnalo nominacija „Geriausia naujovė“. Populiarusis mokslas ir "Vybіr chitachіv" vіd vidannya HPCWire.

2008 m – technologija NVIDIA CUDA laimėjo nominaciją „Techninė Perevaga“. PC žurnalas.

Kas yra CUDA

CUDA(trumpa anglų kalba) Apskaičiuokite vieningą įrenginių architektūrą, Pažodžiui - vieningas ūkinių pastatų architektūros skaičiavimas) - architektūra (programinės ir techninės įrangos paveldėjimas), leidžianti vibruoti GPU slapto prisipažinimo apskaičiavimas, jo GPU Praktiškai atlieka alinančio procesoriaus vaidmenį.

Technologijos NVIDIA CUDA™- mano programavimo kūrimo esmė C, leidžianti mažmenininkams kurti programas, skirtas lankstymo skaičiavimo užduotims trumpiau nei valandą, grafinių procesorių intensyvumui apskaičiuoti. Pasaulis jau praktikuoja milijonus GPU su parama CUDA, tie tūkstančiai programuotojų jau naudoja įrankius (nemokamai!) CUDA greitesniems papildymams ir daugiausiai išteklių reikalaujančių užduočių – vaizdo kodavimo, garso įrašymo, naftos ir dujų tyrimų, gaminių modeliavimo, medicininio vaizdo gavimo ir mokslinių tyrimų – plėtrai.

CUDA suteikia mažmenininkui galimybę organizuoti teismo tyrimą su prieiga prie instrukcijų rinkinio grafiniam pagreitinimui ir atminties cherubavimui, organizuoti pagal naują sulankstomą lygiagretų skaičiavimą. Grafinis priskoryuvach іz pіdtrimkoy CUDA tampa laidine programuota architektūra, panašia į šiandieninę. Vis dėlto mes suteikiame užsakovui žemą, žemą ir aukštą prieigą prie nuosavybės, roblyach CUDA būtinas rimtų aukštos kokybės įrankių, tokių kaip kompiliatoriai, skaičiuotuvai, matematinės bibliotekos, programinės įrangos platformos, pagrindas.

Uralsky, pagrindinis technologijų specialistas NVIDIA, porіvnyuyuchi GPUі sakyk taip: - Tse poshlyahovik. Vіn їzdit zavzhdі kad skіz, bet ne per greitai. BET GPU- Tse sportinis automobilis. Jūs tiesiog niekur neikite ant nešvaraus brangaus vyno, o uždenkite garną ir parodysite visą savo švediškumą, apie kurį pošliachovikas net nesvajojo!

Technologijos galimybė CUDA

CUDA technologija

Vladimiras Frolovas,[apsaugotas el. paštas]

Abstraktus

Straipsnyje pasakojama apie CUDA technologiją, kuri leidžia programuotojui nulaužti vaizdo plokštes kaip konkretesnį skaičių. Nvidia sukurti įrankiai leidžia rašyti grafikos procesoriaus (GPU) programas C++ pogrupiuose. Tai padės programuotojui naudoti atspalvius ir suprasti robotų grafikos dujotiekio procesą. Straipsnyje pristatomas programavimo su CUDA alternatyvomis taikymas bei įvairūs optimizavimo metodai.

1. Įvadas

Skaičiavimo technologijų vystymasis sparčiai tęsiasi dešimtis metų. Grindų dangos yra švelnios, tačiau tuo pat metu procesorių pardavėjai praktiškai perėjo į vadinamąją „silicio kurčiųjų kutą“. Neįsivaizduojamas laikrodžio dažnio padidėjimas tapo neįmanomas dėl daugelio rimtų technologinių priežasčių.

Štai kodėl visi šiuolaikinių skaičiavimo sistemų algoritmai yra skirti procesorių ir branduolių skaičiaus didinimui, o ne vieno procesoriaus dažnio didinimui. Centrinio procesoriaus (CPU) branduolių skaičius pažangiose sistemose yra daugiau nei 8.

Kita priežastis – akivaizdžiai mažas roboto greitis ir operacinė atmintis. Tarsi procesorius nedirbtų greitai, siaurose erdvėse, kaip rodo praktika, tai ne aritmetiniai veiksmai, o veikiau artimi atminties talpyklos praleidimams.

Pasigrožėkite ir stebėkite GPU (grafikos procesoriaus bloko) grafikos procesorius, ten, lygiagretumo keliu, jie buvo turtingi anksčiau. Dabartinėse vaizdo plokštėse, pavyzdžiui, GF8800GTX, procesorių skaičius gali siekti 128. Tokių sistemų produktyvumas, tinkamai suprogramavus, gali būti nemenkas (1 pav.).

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

Jei pirmosios vaizdo plokštės tik pasirodytų parduodamos, smarvė būtų paprasta (kartu su centriniu procesoriumi) aukštesnės specializacijos ūkiniuose pastatuose, skirtuose procesoriui pašalinti iš dviejų pasaulių duomenų vizualizacijos. Plėtojant žaidimų pramonę ir atsiradus tokiems nereikšmingiems žaidimams kaip Doom (2 pav.) ir Wolfenstein 3D (3 pav.), vinilo 3D vizualizavimo paklausa.

Kūdikiai 2.3. Žaidimai Doom ir Wolfenstein 3D

Pirmąsias „Voodoo“ vaizdo plokštes (1996 m.) ir iki 2001 m. „3Dfx“ kompanija sukūrė prieš pat 2001 m. GPU, fiksuodama tik įvesties duomenų operacijų rinkinį.

Programuotojai neturėjo pasirinkimo vizualizacijos algoritmų, o padidėjus lankstumui, atsirado šešėliai – nedidelės programos, kurios vaizdo plokšte vizualizuoja odos viršūnę ir odos pikselį. Jų užduotys apėmė transformacijas virš smailių ir šešėlių paryškinimą taškuose, pavyzdžiui, Phong modeliui.

Norėdamas kartais šešėliai atimdavo net stiprų vystymąsi, supratę, kad aukštosiose mokyklose buvo sukurta trivimerių transformacijų ir rosterizacijos smarvė. Tuo pačiu metu, kai GPU vystosi daugiafunkcinėse turtingų procesorių sistemose, filmų šešėliai yra perpildyti aukštų specializacijų.

Galima palyginti su mano FORTRAN su tuo, kad smirda, kaip ir FORTRAN, buvo pirmieji, tačiau jie buvo pripažinti vyrishennya tik vienos rūšies užduotimis. Shader'iai mažai naudingi tobulinant bet kokias kitas užduotis, krim trivialias transformacijas ir rastracijas, kaip FORTRAN, netinkamos užduoties užbaigimui, nesusijusios su skaitiniais rozrahunki.

Šiandien yra tendencija naudoti netradicines vaizdo plokštes, skirtas vizualizacijai kvantinės mechanikos spintose, gabalų žvalgybai, fizinėms rekonstrukcijoms, kriptografijai, fiziškai taisyklingai vizualizacijai, rekonstrukcijai iš nuotraukų, identifikavimui. Qi zavdannya be rankų vibravo ties grafinių API (DirectX, OpenGL) ribomis, qi API šukės buvo sukurtos kitų zastosuvanų.

Bendrojo GPU programavimo kūrimas (General Programming on GPU, GPGPU) logiškai paskatino technologijas, skirtas platesnei gamybos sričiai, mažesnei tvarkaraščiui. Dėl to Nvidia sukūrė technologiją Compute Unified Device Architecture (arba trumpiau CUDA), o konkuruojanti ATI sukūrė STREAM technologiją.

Reikėtų pažymėti, kad šio straipsnio rašymo metu STREAM technologija tvirtai palaikė CUDA plėtrą, ir to čia nematyti. Mes orientuojamės į CUDA – GPGPU technologiją, kuri leidžia rašyti programas keliuose C++ filmuose.

2. Pagrindinis skirtumas tarp procesoriaus ir GPU

Trumpai pažvelkime į regionų skirtumus ir specifines centrinio procesoriaus bei vaizdo plokštės savybes.

2.1. Galimybės

CPU krūva priedų, kad būtų galima vykdyti pagrindinį planą ir išnaudoti atmintį, kuri yra tinkamai sprendžiama. CPU programas galima atsisiųsti nepertraukiamai iki tiesinės arba vienalytės atminties vidurio.

Dėl GPU tai nėra blogai. Kaip žinote, perskaičius šį straipsnį, CUDA gali matyti 6 atminties rūšis. Skaityti galima nuo bet kokio vidurio, fiziškai pasiekiamo, bet užsirašyti – ne per vidurį. Priežastis slypi tame, kad GPU bet kuriuo atveju yra specifinis priedas, jį atpažįstame konkretiems tikslams. Tse obezhennya zaprovadzhennja už zbіlshennya svydkostі robotai singhnyh algogorіvіv і zvіzhennya vartosti і obladannya.

2.2. Švediškas atminties kodas

Ta pati daugiau skaičiavimo sistemų problema yra ta, kad atmintis yra efektyvesnė nei procesorius. CPU įsilaužėliai tam tikra prasme pažeidžia talpyklas. Dažniausiai atminties perkėlimas atliekamas superoperacinėje arba talpykloje, kuri veikia procesoriaus dažniu. Tse leidžia jums sutaupyti valandą nuo mirties iki mirties, kuri dažniausiai būna pergalinga, ir geriausiu įmanomu būdu sugadinti procesorių.

Pagarbiai programuotojui talpyklos praktiškai aiškios. Kaip ir skaitant, taip ir rašant duomenys ne vieną kartą nukeliami į operatyvinę atmintį, o eina per talpyklas. Leisk man, zokrema, greitai perskaityti dienos prasmę po įrašo.

GPU (čia galite naudoti 8-os serijos GF vaizdo plokštes) talpyklos taip pat svarbios, tačiau mechanizmas nėra toks sunkus kaip CPU. Pirmuoju būdu išgryninimas iš visiško atminties tipų manija, bet kitu būdu talpyklos praktikuojamos tik skaitant.

GPU yra daugiau nei pakankamai laiko nepamiršti sumokėti už papildomus lygiagrečius skaičiavimus. Šiuo metu vienas zavdannya patikrina duomenis, pratsyut іnshі, pasiruošęs skaičiuoti. Tai vienas iš pagrindinių CUDA principų, leidžiantis gerokai padidinti visos sistemos produktyvumą.

3. CUDA šerdis

3.1. Srauto modelis

Skaitmeninė CUDA architektūra pagrįsta koncepcijaviena komanda anoniminiams duomenims(Single Instruction Multiple Data, SIMD) daugiaprocesorius.

Galima naudoti SIMD koncepciją, kad viena instrukcija leidžia vienu metu rinkti anoniminius duomenis. Pavyzdžiui, komanda addps Pentium 3 procesoriuje ir naujesniuose Pentium modeliuose leidžia vienu metu pridėti 4 vieno tikslumo slankiojo kablelio skaičius.

Daugiaprocesorius yra kelių branduolių SIMD procesorius, kurio visuose branduoliuose yra tik viena instrukcija. Kelių procesorių šerdies oda nėra skaliarinė, taigi. jis nepalaiko vektorinių operacijų grynai.

Prieš Timą, kaip tęsti, pristatysime keletą susitikimų. Svarbu tai, kad pagal priedą ir šios statistikos pagrindą, tai nėra tie, prieš kuriuos skambino dauguma programuotojų. Mes naudosime tokias sąvokas, kad išvengtume skirtumų nuo CUDA dokumentacijos.

Pagal įrenginį (įrenginį) mūsų straipsnyje yra pagrįstas vaizdo adapteris, palaikantis CUDA tvarkyklę arba kitas priedų specializacijas, programavimo programas, kurios naudoja CUDA (pvz., NVIDIA Tesla), priskyrimus. Mūsų straipsnyje matome, kad GPU yra mažiau panašus į loginį turtą, būdingą konkrečioms diegimo detalėms.

Pagrindinis kompiuteris (host) yra pagrindinėje kompiuterio operatyvinėje atmintyje esančios programos pavadinimas, pagrindinis CPU ir pagrindinių roboto su priedu funkcijų nepaisymas.

Tiesą sakant, ta jūsų programos dalis, kuri veikia CPUšeimininkas, ir tavo vaizdo plokštė - priedas. Logiškai mąstant, galite taikyti kaip kelių procesorių rinkinį (maži 4) ir CUDA tvarkyklę.

Ryžiai. 4. Priedas

Tarkime, kad savo plėtinyje norime paleisti procedūrą N gijų (todėl norime lygiagretinti robotą). Kalbant apie CUDA dokumentaciją, pavadinkime procedūrą branduoliu.

CUDA architektūros ypatumas – blokinis sietas, neribojamas gausiais srauto priedais (5 pav.). CUDA tvarkyklė savarankiškai paskirsto išteklius ir kuria gijas.

Ryžiai. 5. Srauto organizavimas

Ant pav. 5. branduolys yra pažymėtas kaip branduolys. Visos gijos, kurios atsitrenkia į šerdį, yra sujungtos į blokus (Block), o blokai pagal savo prigimtį yra sujungti į tinklelį (Grid).

Kaip matyti iš 5 paveikslo, srautams nustatyti naudojami dviejų pasaulių indeksai. CUDA rozrobniks suteikė galimybę dirbti su trivum, dviejų pasaulių arba paprastais (vieno pasaulio) indeksais, priklausomai nuo programuotojo patogumo.

Laukiniame tipe indeksas yra trivialus pagal vektorius. Dermaliniam siūlui bus pateikta: gijos indeksas threadIdx bloko viduryje ir bloko indeksas blockIdx tinklelio viduryje. Paleidžiant, visos gijos bus atnaujintos be jokių indeksų. Tiesą sakant, per qi indeksą programuotojas kontroliuoja valdymą, o tai reiškia, kad dalis jo yra apdorojama odos prakaitu.

Tiekimo įrodymai, kodėl mažmenininkai išsivežė tokią organizaciją, nėra banalūs. Atrodo, kad viena iš priežasčių yra ta, kad vienas blokas garantuoja pergalę ant vieno Pridėsiu multiprocesorių, bet vienas multiprocesorius gali laimėti šprotą įvairių blokų. Kitos priežastys, kodėl reikia išsiaiškinti, buvo nurodytos straipsnio valandą.

Užduočių (srautų) blokas daugiaprocesoriuje įveikiamas dalimis arba telkiniais, vadinamais deformacija. Metmenų išplėtimas šiuo metu vaizdo plokštėse iš CUDA palaikymo yra iki 32 srautų. Komandos metmenų telkinio viduryje nustatomos į SIMD stilius, ty. visi siūlai metmenų viduryje vienu metu gali turėti tik vieną nurodymą.

Štai kitas įspėjimas. Architektūrose šių straipsnių rašymo metu procesorių skaičius vieno daugiaprocesoriaus viduryje yra 8, o ne 32. Aišku, kad ne visas metimas per vieną valandą sumušamas, jis padalintas į 4 dalis, jie yra plakami nuosekliai (nes procesoriai) .

Ale, visų pirma, CUDA pardavėjai nereguliuoja metmenų dydžio. Jų robotuose smarvė nustato metmenų dydžio parametrą, o ne skaičių 32. Kitaip, loginiu požiūriu, pats metmenys yra minimalus srautų baseinas, apie kurį galima sakyti, kad visi srautai karo vidurys skaičiuojami iš karto - išspręsti sistema nebus pažeista.

3.1.1. Druskos pašalinimas

Na, jūs kaltinate maistą: jei tą pačią akimirką visi srautai metmenų viduryje rašo būtent tą instrukciją, tai kaip jus galima išvalyti? Net jei programos kodas bus išvalytas, instrukcijos bus kitokios. Čia pradeda veikti standartinis SIMD programavimo sprendimas (6 pav.).

Ryžiai. 6. Derinimo organizavimas SIMD

Nagi kitas kodas:

jei (sąlyga)B;

SISD (Single Instruction Single Data) atveju operatorius A yra vikonuuojamas, jį iš naujo patikrina protas, tada operatoriai B ir D yra vikonuojami (taigi, protas yra tikras).

Dabar turėkime 10 srautų, parašytų SIMD stiliumi. Visuose 10 srautų sugendame operatoriui A, tada patikriname psichinę būseną ir pasirodo, kad 9 iš 10 srautų tai tiesa, o viename – blogai.

Supratau, kad negalime paleisti 9 gijų operatoriaus B vykdymui, o vienos operatoriaus C vykdymui, todėl vieną valandą visoms gijomis galima paleisti tik vieną instrukciją. Prie šitos vapadkos reikia taip pridurti: puodas „įvaromas“ ant pakaušio, kad buvo išvalytas, kad vynas nedavė jokios duoklės, ir paimti 9 upeliai, kurie liko. išeiti. „Įveskime“ 9 gijas, kurios įveikia operatorių B, o vieną giją perduodame operatoriui C. Kitos gijos vėl sujungiamos ir įveikia operatorių D visus iš karto.

Rezultatas yra bendras: procesorių resursai ne tik švaistomi tuščioms perdirbimo kovoms nutrūkusiuose srautuose, bet ir daug turtingesni, dėl OBIGI gіlki pralaimėjimo mums bus nepatogu.

Tačiau ne viskas taip blogai, kaip matote iš pirmo žvilgsnio. Didelis technologijų pranašumas yra tai, kad galite pamatyti tuos, kurie sutelkia dėmesį į dinamišką CUDA tvarkyklės valdymą, o programuotojui smarvė yra aišku. Tuo pačiu metu, lakstydamas su SSE komandomis iš šiuolaikinių procesorių (tik išbandykite 4 algoritmo kopijas iš karto), pats programuotojas kaltas dėl smulkmenų: sujunkite duomenis iš keturių, nepamirškite apie atranką ir pradėkite rašyti. žemu lygiu, iš tikrųjų, kaip surinkime.

Iš ūsų minėtos vaikštančios dar vienas pagarbus visnovokas. Razgaluzhennya yra galingų jėgų našumo kritimo priežastis. Shkidlivimi є mažiau nei tі razgaluzhennya, ant kurių upeliai skiriasi viename telkinio metmenų srautų viduryje. Kai kuriais atvejais srautai pasklido vieno bloko viduryje, bet skirtinguose metmenų telkiniuose arba skirtingų blokų viduryje, nesukeldami jokio poveikio.

3.1.2. Sąveika tarp srautų

Rašant šį straipsnį, ar sąveika tarp gijų (sinchronizavimas ir keitimasis duomenimis) buvo įmanoma tik bloko viduryje. Štai kodėl neįmanoma tarpusavyje organizuoti skirtingų blokų srautų, nes jie yra koroziniai ir turi mažiau dokumentavimo galimybių.

Nors ir nedokumentuotos galimybės, tačiau jomis puošti nerekomenduojama. Priežastis ta, kad smarvė kyla dėl specifinių kitos sistemos aparatinės įrangos savybių.

Visų bloko viduryje atliekamų užduočių sinchronizavimą valdo __synchtreads funkcijos iškvietimas. Pinigų keitimas įmanomas per atmintį, kuri yra padalinta, todėl ji yra tokia svarbi visoms užduotims bloko viduryje.

3.2. Atmintis

CUDA turi šešių tipų atmintį (7 pav.). C registras, vietinė, globali, paskirstyta, pastovi ir tekstūrinė atmintis.

Tokį didelį skaičių pakerta vaizdo plokštės specifika ir pirmieji pripažinimai, taip pat prekybininkai yakomog sistemą statyti pigiau, įvairiais būdais aukodami arba universalumą, arba swidkistyu.

Ryžiai. 7. Žiūrėkite CUDA atmintį

3.2.0. Registras

Kiek gali, kompiliatorius stengiasi visas lokaliai pakeistas funkcijas patalpinti į registrus. Prieiga prie tokių pakeitimų ribojama iki didžiausio greičio. Srautinio perdavimo architektūroje yra 8192 32 bitų registrai, galimi vienam daugiaprocesoriui. Norėdami nustatyti, kiek registrų yra vienai gijai, turite padalyti skaičių (8192) pagal bloko dydį (naujai gijų skaičių).

Esant dideliam 64 srautų skaičiui, į bloką pateks iš viso 128 registrai (yra keletas objektyvių kriterijų, tačiau 64 turėtų būti per vidurį, jei norite atlikti turtingus darbus). Tikrai, 128 registrų nvcc visai nesimato. Skambinkite VIN, daugiau nei 40 neduokite, o pakeistus ištrinkite į vietinę atmintį. Taigi atrodo, kad ant vieno daugiaprocesoriaus galima laimėti šprotą blokų. Kompiliatorius stengsis maksimaliai padidinti blokų, kuriuos galima apdoroti vienu metu, skaičių. Siekiant didesnio efektyvumo, reikia paimti mažiau nei 32 registrus. Tada teoriškai viename daugiaprocesoriuje galite paleisti 4 blokus (8 warp-іv, taigi 64 gijos viename bloke). Tačiau čia svarbiau rūpintis bendra atmintimi, kuri yra užimta gijomis, nes vienas blokas užima visą bendrinamą atmintį, dviejų tokių blokų negali užimti keli procesoriai vienu metu.

3.2.1. Vietinė atmintis

Jei vietiniai procedūrų duomenys užima per daug vietos arba kompiliatorius negali jų apskaičiuoti paskutinį kartą, galite juos patalpinti į vietinę atmintį. Kam galite priimti, pavyzdžiui, pateiktus skirtingų išplėtimo tipų rodiklius.

Fiziškai vietinė atmintis yra pasaulinės atminties analogas ir veikia su tієyu ir swidkіst. Straipsnio rašymo metu nebuvo žinomų mechanizmų, kurie leistų aiškiai nuskaityti kompiliatorių naudojant vietinę atmintį konkretiems pakeitimams atlikti. Svarbu patikrinti vietinę atmintį, o ne vikoristovuvat її zovsіm (4 skyrius „Rekomendacijos optimizavimui“).

3.2.2. pasaulinė atmintis

CUDA dokumentacija yra vienas iš pagrindinių laimėjimųTechnologija, suteikianti galimybę pakankamai adresuoti pasaulinę atmintį. Taigi galite skaityti iš bet kurio atminties vidurio ir taip pat galite rašyti tam tikrame viduryje (GPU tai neskamba).

Prote už universalumą laikais turi būti pareikštas verkti swidkistyu. Pasaulinė atmintis nėra talpykloje. Vaughn pratsyuє dar daugiau povіlno, kіlkіst zvernenі į pasaulinę atmintį sіd bet kuriuo metu minimizuvati.

Visuotinė atmintis yra būtina norint išsaugoti robotų programų rezultatus prieš redaguojant juos pagrindiniame kompiuteryje (numatytojoje DRAM atmintyje). To priežastis yra ta, kad atmintis yra globali – vienintelė atmintis, kurią galima įrašyti.

Pakeitimai, išreikšti kvalifikatoriumi __global__, įrašyti į pasaulio atmintį. Visuotinę atmintį taip pat galima peržiūrėti dinamiškai, iškvietus pagrindinio kompiuterio funkciją cudaMalloc(void* mem, int size). Pridėsiu šią funkciją, iškviesti negalima. Atrodo, kad pagrindinė programa gali pasirūpinti atmintimi, tuo, kas veikia CPU. Galite nepaisyti duomenų iš pagrindinio kompiuterio spustelėdami funkciją cudaMemcpy:

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

Taigi galite patys pradėti atvirkštinę procedūrą:

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

Ši wiki taip pat veikia iš pagrindinio kompiuterio.

Dirbant su globalia atmintimi, svarbu atsiminti susijungimo taisykles. Pagrindinė idėja yra ta, kad trečioji atminties dalis yra saugoma iki paskutinės vidurinės atminties, be to, 4,8 arba 16 baitų. Dėl to pirmoji gija kalta dėl adreso trūkimo, vibruojant kordone, aišku, 4,8 arba 16 baitų. CudaMalloc pasukami adresai yra mažiausiai 256 baitai už sienos.

3.2.3. Prisiminkite, kas yra padalinta

Dalijama atmintis nėra talpinama, tačiau atmintis yra greita. Її rekomenduojama vicorate kaip talpyklos patikrinimą. Iš viso yra 16 KB bendrinamos atminties vienam kelių procesoriui. Padalinę skaičių iš dienų skaičiaus bloke, imame didžiausią padalytą atminties kiekį, prieinamą vienam srautui (kadangi planuojama jį laimėti nepriklausomai nuo visų srautų).

Prisiminkite ryžius, kurie yra padalinti, tuos, kurie yra adresuojami vienodai visiems vidurinio bloko lyderiams (7 pav.). Akivaizdu, kad galite laimėti tik vieną bloką duomenų mainams tarp srautų.

Garantuojama, kad valandą daugiaprocesoriaus blokas bus saugomas atmintyje. Tačiau kadangi blokas pakeičiamas daugiaprocesoriuje, negarantuojama, kad vietoj jo bus išsaugotas senas blokas. Štai kodėl neužtenka bandyti sinchronizuoti užduotis tarp blokų, paliekant jas atmintyje kaip duomenis ir pasikliaujant jų santaupomis.

Pakeitimai, išreikšti su __shared__ kvalifikatoriumi, įrašomi į atmintį, kad jie yra bendrinami.

share_float mem_shared;

Kitą kartą nuraminkite, kokia atmintis, kas dalijama, vien blokui. Tam reikia nulaužti kaip talpyklos talpyklą, atliekant įvairių masyvo elementų paiešką, pavyzdžiui:

float x = atmintis_shared;

De threadIdx.x – gijos indeksas x bloko viduryje.

3.2.4. Nuolatinė atmintis

Nuolatinė atmintis saugoma talpykloje, kaip parodyta Fig. 4. Talpykla naudojama viename kelių procesorių egzemplioriuje, o tai taip pat yra pagrindinė vidurinio bloko užduotis. Pagrindiniame kompiuteryje galite rašyti į nuolatinę atmintį iškviesdami funkciją cudaMemcpyToSymbol. Pridėsiu nuolatinę atmintį, kurią galima tik skaityti.

Nuolatinė atmintis vikoristanui patogesnė. Galite rozmіschuvati in niy danі be-kokio tipo, kad skaityti їх už paprastą traukos pagalbą.

#define N 100

Constant__int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

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

// __global__ reiškia, kad įrenginio_kernel yra branduolys, todėl jį galima paleisti GPU

Global__void device_kernel()

int a = gpu_buferis;

int b = gpu_buferis + gpu_buferis;

// gpu_buffer = a; VAKARĖLIS! nuolatinė atmintis yra prieinama tik skaitymui

Kadangi talpykla naudojama nuolatinei atminčiai, prieiga prie jos yra saugi. Vienintelis, bet vis tiek puikus, mažas pastovios atminties kiekis yra tai, kad ji tampa mažesnė nei 64 Kbaitai (visam priedui). Kodėl taip akivaizdu, kad kontekstinėje atmintyje prasminga išsaugoti tik nedidelį skaičių duomenų, kurie dažnai būna pergalingi.

3.2.5. Tekstūros atmintis

Tekstūros atmintis yra talpykloje (4 pav.). „Skin“ daugiaprocesoriuje yra tik viena talpykla, todėl visa talpykla yra didžiausia visam viduriniam blokui.

Tekstūrinės atminties pavadinimas (ir, deja, funkcionalumas) sumažintas, kad būtų suprantama „tekstūra“ ir „tekstūra“. Tekstūravimas – tekstūrų (tik paveikslėlių) pritaikymo daugiakampiui procesas rastravimo proceso metu. Tekstūros atmintis yra optimizuota 2D duomenims ir gali būti:

shvidka vibirka fiksuoto rozmіru reikšmė (baitas, žodis, subwiyne arba keturių žodžių) iš vieno ar dviejų pasaulių masyvo;

normalizuotas adresavimas slankiaisiais skaičiais intervalais. Potim їх galima rinktis, vikoristuuu normalіzovanu kreipiantis. Gauta reikšmė bus float4 tipo žodis su intervalu ;

CudaMalloc((galioja**) &gpu_memory, N*dydis (uint4 )); // matyt GPU atmintis

// Sureguliuokite tekstūros medžio parametrus

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

Texture.addressMode = cudaAddressModeWrap;

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

tekstūra.normalizuota = klaidinga; // nesukite normalizuoto adresavimo

CudaBindTexture(0, tekstūra, gpu_memory, N ) // pakeiskime atmintį į tekstūrą


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

// __global__ reiškia, kad įrenginio_kernel yra branduolys, todėl jį reikia lygiagretinti

Global__void device_kernel()

uint4 a = tex1Dfetch(tekstūra,0); // tokiu būdu galite pasirinkti daugiau duomenų!

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

int c = a.x*b.y;

...

3.3. Paprastas užpakalis

Kaip paprastas užpakalis, siūloma pažvelgti į cppIntegration programą iš CUDA SDK. Vaughn demonstruoja CUDA darbo eigą, taip pat nvcc (specialus C++ kompiliatorius, skirtas Nvidia) MS Visual Studio versiją, kuri supaprastins programų kūrimą CUDA.

4.1. Tinkamai įvykdykite savo viršininko pralaimėjimą

Ne visos užduotys tinka SIMD architektūrai. Pavyzdžiui, jūsų užduotis, kuriai nėra pridėta, tai įmanoma, negalima naudoti GPU. Ir vis dėlto, jūs griežtai pažeidėte GPU triukus, reikėjo suskaidyti algoritmą į tokias dalis, kad smarvė galėtų veiksmingai nugalėti SIMD stilių. Būtina – pakeisti savo užduoties tobulinimo algoritmą, sugalvoti naują – tokį, kuris tiktų SIMD. Kaip skirtingos GPU srities pavyzdį, galite įgyvendinti masyvo elementų piramidinį lankstymą.

4.2. Pasirinkite atminties tipą

Įdėkite savo duomenis į tekstūros arba nuolatinę atmintį, kad visos vieno bloko užduotys būtų konvertuojamos į vieną atminties lizdą arba arti sukrautų lizdų. Du duomenys gali būti efektyviai apdorojami naudojant papildomas text2Dfetch ir text2D funkcijas. Tekstūros atmintis yra specialiai optimizuota dviem pasauliams.

Užkariaukite pasaulinę atmintį iš padalintos atminties, nes visos užduotys nesistemingai paskirstomos į skirtingas, toli siekiančias vienos rūšies atminties vietas (su skirtingais adresais ar koordinatėmis, pvz., 2D / 3D duomenys).

pasaulinė atmintis => atmintis, kuri yra padalinta

syncthreads();

Surinkite duomenis atmintyje

syncthreads();

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

4.3. Priminkite lichnikams apie atmintį

Kompiliatoriaus vėliavėlė --ptxas-options=-v leidžia tiksliai pasakyti, kurie žodžiai ir kokia atmintis (registrai, kurie yra paskirstyti, vietiniai, pastovūs) victorist. Kadangi kompiliatorius naudoja vietinę atmintį, jūs apie tai žinote. Duomenų apie laimėjusių atminties skaičių ir tipus analizė gali labai padėti optimizuoti programą.

4.4. Stenkitės kuo labiau sumažinti bendrinamų registrų ir prisiminimų skaičių

Kuo didesnė deformacijos registro šerdis arba atmintis, kuri yra padalinta, tuo mažiau srautų (pirminis iškrypimas-iv) vienu metu gali laimėti daugiaprocesoriuje, nes keičiamasi kelių procesorių ištekliais. Todėl nedidelis registrų užimtumo padidėjimas, tačiau plečiamos atminties, produktyvumas gali sumažėti du kartus – per tuos, kurie dabar net dvigubai mažiau deformuojasi vienu metu daugiaprocesoriuje. .

4.5. Atmintis, ką dalija, vietos deputatas.

Tarsi Nvidia kompiliatorius prarado duomenis vietinėje atmintyje 'yat, scho razdelyaetsya (bendra atmintis).

Dažniausiai kompiliatorius gali būti pakeistas vietinėje atmintyje, nes jis nebus paimamas labai dažnai. Pavyzdžiui, akumuliatorius devi kaupia vertes, rozrakhovuyuschos s tsiklі. Kaip puikus obsyagi kodo ciklas (nors ir ne valandą vykonannya!), Kompiliatorius gali įdėti jūsų bateriją į vietinę atmintį, tk. Vіn vykoristovuєtsya labai retai, o registrų nedaug. Produktyvumo kaina kartais gali būti didelė.

Na, o jei jūs tikrai retai jį keičiate nugalėtoju, geriau įdėkite jį į pasaulinę atmintį.

Jei norite, kad kompiliatorius automatiškai paskirstytų tokius pakeitimus vietinėje atmintyje, tai gali būti protinga, bet ne iš tikrųjų. Tai nėra lengva sužinoti tam tikroje vietoje su būsimomis programų modifikacijomis, nes ji dažnai keičiama dažniau. Kompiliatorius gali arba negali perkelti tokio pakeitimo į registro atmintį. Jei __global__ modifikatorius bus nurodytas aiškiai, programuotojas bus žiauriau informuotas.

4.6. Teisingumo ciklai

Sukimosi ciklai yra standartinis būdas padidinti produktyvumą turtingose ​​sistemose. Jogo esmė ta, kad odos iteracijoje jūs galite laimėti vis daugiau ir daugiau, taip keisdami iteracijų skaičių, o tai reiškia mentalinių perėjimų skaičių, kad procesorius galėtų laimėti.

Ašis yra tai, kaip galima atidaryti prasmės ciklą sumi masyvu (pavyzdžiui, visą):

int a[N]; intsum;

už (int i=0;i

Zrozumilo, ciklus galima paleisti ir rankiniu būdu (kaip parodyta aukščiau), tačiau praktika yra neproduktyvi. Geriau pakoreguoti C++ šablonus, kad gautumėte daugiau funkcijų, kurias reikia žinoti.

šabloną

klasė ArraySumm

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

šabloną

klasė ArraySumm<0,T>

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

už (int i=0;i

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

Pažymėkite vieną varnelę kaip nvcc kompiliatoriaus funkciją. Kompiliatorius visada turi leisti uždaryti __device__ tipo funkcijas (kad įsitikintumėte, naudokite specialią direktyvą __noinline__).

Otzhe, jus gali įkvėpti faktas, kad užpakalis, panašus į smailią daiktą, užsidega nuo paprastos operatorių sekos, o kodėl gi nepasielgus dėl ranka parašyto kodo efektyvumo. Tačiau laukiniam tipui (ne nvcc) niekas negali turėti grobio, nes inline yra tik kompiliatoriaus įvestis, ir jūs galite į tai nepaisyti. Negarantuojama, kad jūsų funkcijos bus atkurtos.

4.7. Naršykite duomenis ir pasirinkite 16 baitų

Naršykite duomenų struktūras pagal 16 baitų kordoną. Tokiu atveju kompiliatorius gali paimti jiems specialias instrukcijas, kad jie galėtų paimti duomenų kiekį po 16 baitų.

Jei paskolos struktūra yra 8 b ir mažesnė, galite ją pakeisti 8 b. Arba galite pasirinkti du 8 baitų pakeitimus vienu metu, sujungdami du 8 baitų pakeitimus į struktūrą, esančią už papildomos sąjungos, arba įtraukdami juos į sąrašą. Su šiais dalykais reikia elgtis atsargiai, kompiliatorius gali saugoti duomenis vietinėje atmintyje, bet ne registre.

4.8. Bankų konfliktai atmintyje, kurie yra padalinti

Platinama atmintis suskirstyta į 16 (viską!) atminties bankų su 4 baitų Croc. Valandai kelių procesorių metmenų gijų telkinys yra padalintas į dvi dalis (kaip metmenų dydis = 32) po 16 gijų, o tai leidžia pasiekti atmintį kortele.

Užduotys skirtingose ​​metmenų pusėse neprieštarauja dalijamiems prisiminimams. Per zavdannya vienos pusės metmenų baseinas bus sumažintas iki tų pačių atminties bankų, kaltinant žlugimą ir dėl to sumažėjusį produktyvumą. Zavdannya vienos pusės metmenų ribose gali išaugti iki skirtingų atminties kaimų, kurie yra padalinti, su daina kroko.

Optimalus dydis yra 4, 12, 28, ..., 2 n-4 baitai (8 pav.).



Ryžiai. 8. Optimalus prigludimas.

Chi nėra optimalaus dydžio – 1, 8, 16, 32, ..., 2^n baitų (9 pav.).



Ryžiai. 9. Neoptimalus prigludimas

4.9. Judančių duomenų sumažinimas Host<=>prietaisas

Pabandykite perkelti tarpinius rezultatus į pagrindinį kompiuterį, kad būtų galima apdoroti papildomą procesorių. Jei nediegiate viso algoritmo, dalį pagrindinės dalies paimkite GPU, palikdami CPU mažiau užduočių.

5. CPU/GPU nešiojama matematikos biblioteka

Šio straipsnio autorius parašė MGML_MATH biblioteką, kurią galima perkelti, darbui su paprastais erdviais objektais, kodą kaip praktišką tiek plėtinyje, tiek pagrindiniame kompiuteryje.

MGML_MATH biblioteka gali būti naudojama kaip pagrindas CPU/GPU nešiojamoms (arba hibridinėms) sistemoms rašyti fizinėms, grafinėms ir kitoms erdvės užduotims kurti. Pagrindinis privalumas yra tas, kad vienas ir tas pats kodas gali būti koreguojamas tiek CPU, tiek GPU, o tuo pačiu metu iš kitos pusės jį būtų galima įdiegti bibliotekai.

6 . Literatūra

Krisas Kasperskis. Programos optimizavimo technika. Veiksmingas atminties atkūrimas. - Sankt Peterburgas: BHV-Petersburg, 2003. - 464 p.: il.

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

CUDA programavimo vadovas 1.1. 14-15 psl

CUDA programavimo vadovas 1.1. 48 puslapis

Iš Darvino evoliucijos teorijos, pirmoji į žmogų panaši mavpa (pvz
buti tiksliau – homo antecessor, person-perepadnik) apsimetė dievu
mes turime. Bagatotono skaičiavimo centrai su tūkstančiu ir daugiau radijo lempų,
kas užima miesto sienas, pakeista pivkilograminiais užrašais, jaki, į kalbą,
pirmiausia neaukokite produktyvumo. Pasikeitė antediluvinės drukarų mašinos
kitiems visada tinka tam, kam naudinga (navit ant žmogaus kūno)
gausiai funkcionalūs ūkiniai pastatai. „Giganti raptom“ procesoriai nusprendė pasitraukti
grafinė „akmens“ šerdis. Ir vaizdo plokštės pradėjo ne tik rodyti paveikslėlį
priimti FPS ir grafikos kokybę, bet taip pat atlikti visus skaičiavimus. Taigi
daugiau darbo! Apie turtingo srautinio perdavimo technologiją apskaičiuokite ją GPU pagalba, ir bus paminėta.

Kodėl GPU?

Tsikavo, kodėl jie perkėlė visą skaičiavimo mіt į grafiką
adapteris? Kaip matote, procesorius yra madingas, ir vargu ar jis atsisakys savo šilumos
Mistechko. Ale, GPU turi prie rankovės iš karto porą kozirivų su jokeriu, ta rankovė
vistachaє. Šiuolaikinis centrinis laisvės atėmimo procesorius už maksimumo apribojimą
produktyvumas apdorojant ciliarinius duomenis ir duomenis su slankiuoju
Komi, ypač nesikišant į jokį lygiagretų informacijos apdorojimą. Tas pats
valandą vaizdo plokštės architektūra leidžia greitai ir be problemų „lygiagretumą“
duomenų rinkimas. Iš vienos pusės matyti daugiakampiai (3D konvejerio rėmui),
z іnshoy – pikselių tekstūrų apdorojimas. Matyti, kad tai „laimingas
gedimas“ kortelės šerdyje. Be to, atminties robotas ir vaizdo procesorius
optimali, apatinė nuoroda "RAM-cache-procesorius". Tuo metu, jei esate vienišas
vaizdo plokštes pradeda konvertuoti vienas GPU srauto procesorius, kitaip
Vienatvė tuo pačiu įsipainioja į kitą ir iš principo ją galima lengvai pasiekti
grafikos procesoriaus efektyvumas kartu su magistralės pralaidumo didinimu,
tačiau šiam vantazhennya konvejeriui gali
visokie gudrūs perėjimai ir razgaluzhen. centrinis procesorius per savo
universalumas jūsų procesoriaus poreikiams, talpykla, saugykla
informacija.

Vchenі cholovіki zamyslilis schodo roboti GPU lygiagrečiai apskaičiuojant
matematikai sukūrė teoriją, kad yra daug panašių mokslinių tyrimų
3D grafikos apdorojimas. Daugelis ekspertų mano, kas yra pagrindinis veiksnys
plėtra GPGPU (Bendrosios paskirties skaičiavimas GPU – universalus
rozrahunki vaizdo plokštės pagalba) buvo pristatytas 2003 m. Brook GPU projekte.

Stenfordo universiteto projekto kūrėjams buvo sunku
problema: aparatinės ir programinės įrangos zmusiti grafinio adapterio vibracija
įvairių planų. Aš smirdau veyshlo. Vikoristovuyuchi universalus mov C,
Amerikos mokslininkai zmusili pratsyuvati GPU kaip procesorius, pritaikytas
lygiagretus pjūvis. Po to, kai Brookas pristatė visą eilę VGA projektų projektų,
pvz., Accelerator biblioteka, Brahma biblioteka, sistema
metaprogramavimas GPU++ ir kt.

CUDA!

Plėtros perspektyvų suvokimas buvo gėdingas AMDі NVIDIA
patekti į Brook GPU kaip pitbulis. Jei praleisite rinkodaros politiką, tada
viską teisingai įgyvendinę galite užsidaryti bent jau grafikos sektoriuje
turguje, o skaičiuojant (stebėkite specialiomis skaičiavimo kortomis, kurios
serveriai Tesla su šimtais kelių procesorių), pakeisdami visų procesorių pavadinimus.

Zvichayno, "FPS Volodymyrs" pakilo į kliūvančios odos akmenį saviesiems.
dygsnių, tačiau pagrindinis principas tapo nekintantis – atlikite skaičiavimą
GPU. Iš karto pažvelkime į „žaliąją“ technologiją. CUDA
(Apskaičiuokite vieningą įrenginių architektūrą).

Mūsų „herojės“ robotas dirba saugioje API, be to, vieną kartą dviem.
Pirmasis yra didelio branduolio, CUDA Runtime, su funkcijomis, pvz
paprastesni lygūs yra suskaidomi ir perkeliami į žemesnę API – CUDA tvarkyklę. Taigi
kad frazė „labai patikima“ yra sustingusi iki galo. Visa galia žinoti
pati prie vairuotojo, ir gauti її padėti bibliotekoms, maloniai sukurta
mažmenininkai NVIDIA: CUBLAS
FFT (tiria Fur'e algoritmo išvaizdą). Na, pereikime prie praktikos
medžiagos dalys.

Terminija CUDA

NVIDIA Veikia su savo CUDA API užduotimis. tenisas
v_dr_znyayutsya vіd vyznachen, mokyklų mainai zastosovuyutsya dirbti su centriniu procesoriumi.

Potik (gija)- Surinkite duomenis, kuriuos būtina rinkti (ne
Vymagaє puikus resursіv pіd valandą orobki).

metmenys- 32 srautų grupė. Duomenys renkami tik
warps, dar žinomas kaip warp – minimalūs obsyag duomenys.

blokas- Sukupnіst srautai (vіd 64 iki 512) arba sukupnіst
warpiv (nuo 2 iki 16 tipo).

Sitka (tinklelis)- Tse sukupnіst blokіv. Toks podіl danih
zastosovuєtsya vikljuchno podvischennya produktyvumą. Taigi, koks yra skaičius
daugiaprocesorius yra didelis, tada blokai sujungiami lygiagrečiai. Yakscho w
nepagailėjo kortelės (rekomenduoja mažmenininkams sulankstyti rožes
adapteris yra ne žemesnis nei GeForce 8800 (GTS 320 MB), tada duomenų blokai apdorojami
nuosekliai.

Taip pat NVIDIA pristatyti taip suprantama, kaip branduolys, šeimininkas (šeimininkas)
і prietaisas.

Pratsyuemo!

Bendram darbui su CUDA jums reikia:

1. Žinokite GPU shader branduolių prigimtį, shards yra programavimo esmė
pogaє turi lygias rozpodіlі navantazhennya tarp jų.
2. Užprogramuoti C vidurį, sutvarkyti kai kuriuos aspektus.

Rozrobnikai NVIDIA sulaužė kilka vaizdo plokštės "užpildus".
kitaip, žemesni mi vadinami bachiti. Taigi kodėl nenorite visko pamatyti
subtili architektūra. Rozberemo Budov "akmuo" G80 legendinis GeForce 8800
GTX.

Shader branduolį sudaro aštuonios TPC (Texture Processor Cluster) grupės
tekstūros procesoriai (pvz. GeForce GTX 280- 15 branduolių 8800 GTS
їх šeši, m 8600 - Chotiri ir kt.). Ti, savo širdyje, sudėkite į dvi dalis
streaming multiprocessors (streaming multiprocessor – toli SM). SM (iš viso šeši
16) sulankstytas iš priekio (instrukcijų skaitymo ir iššifravimo užduotis) ir
dujotiekių galas (instrukcijų pabaiga), taip pat aštuoni skaliariniai SP (shader
procesorius) ir du SFU (Super Functional Units). Odos plakimui (vienas
valanda) priekyje pasirinkite metmenis ir užsiimkite joga. Shchob usі metmenų srautai
(Spėju, їх 32 vnt.) Paaiškėjo, kad dujotiekio gale reikia 32/8 = 4 ciklai.

Odinis daugiaprocesorius gali būti vadinamas bendra atmintimi.
Її išplėtimas tapti 16 kilobaitų ir suteikti programuotojui visišką laisvę
pasidaryk pats. Razpodіlya jakas norite :). Bendrinama atmintis apsaugo srautinius skambučius
vienas blokas neatpažįstamas darbui su pikselių šešėliais.

Taip pat SM gali konvertuoti į GDDR. Kam „pasiuvo“ 8 kilobaitus
talpyklos atmintis, kuri išsaugo visas robotams skirtas antraštes (pavyzdžiui, surašymas
konstantos).

Daugiaprocesorių maks. 8192 registrų. Aktyvių blokų skaičius negali būti
daugiau nei aštuoni, o metmenų skaičius - ne daugiau 768/32 = 24. Matyti, kad G80
Per valandą galite apdoroti daugiausiai 32*16*24 = 12 288 srautus. Negaliu ne
pataisykite šiuos skaičius optimizuodami programą nadal (ant vieno dubenėlio vag
- Razmіr blokas, ant іnshiy - srautų skaičius). Galima keisti parametrų balansą
tam buvo skirtas svarbus vaidmuo NVIDIA rekomenduoti vikoristinius blokus
sі 128 arba 256 srautai. Blokas z 512 srautai yra neefektyvus, šukės gali
juda zatrims. Vrakhovuychi visos plonos GPU vaizdo plokštės plius
blogų programavimo įgūdžių, galite sukurti produktyvesnius
zasib lygiagrečiam skaičiavimui. Prieš kalbą apie programavimą ...

programavimas

„Kūrybiškumui“ iš karto iš CUDA tai būtina GeForce vaizdo plokštė ne žemesnė
aštuntoji serija. W

Oficialioje svetainėje reikia trijų programinės įrangos paketų: tvarkyklės
CUDA palaikymas (skirtas skin OS – nuosavas) be tarpinio CUDA SDK paketo (kitas
beta versija) ir papildomos bibliotekos (CUDA įrankių rinkinys). Technologijų palaikymas
„Windows“ operacinės sistemos (XP ir Vista), „Linux“ ir „Mac OS X“.
pasirinkus Vista Ultimate Edition x64 (sakau, kad sistema buvo
Tiesiog nuostabu). Rašymo metu šios eilutės yra svarbios robotams
ForceWare 177.35 tvarkyklė. Kaip vikoristų įrankių rinkinys
programinės įrangos paketas Borland C++ 6 Builder
mano C).

Žmonės, kaip žinau, lengvai pripranta prie naujos aplinkos. Mažiau reikalingas
įsiminti pagrindinius parametrus. _global_ raktinis žodis (prieš funkciją)
Rodo, kad funkcija turi būti prieš branduolį. ї viklikatime central
procesoriaus, o visas robotas bus ant GPU. Wiklik _global_ sukuria daugiau
konkrečios detalės, bet paties tinklo išplėtimas, bloko ir branduolio išplėtimas bus
įstrigo. Pavyzdžiui, eilutė _global_ void saxpy_parallel<<>>, de X –
tinklelio dydis, o Y – bloko dydis, kuris lemia parametrų skaičių.

Simbolis _įrenginys_ reiškia, kad funkciją iškviečia grafinė šerdis, bet
vadovaukitės šiomis instrukcijomis. Ši funkcija saugoma kelių procesorių atmintyje,
otzhe, otrimati її adresas neįmanomas. Priešdėlis _host_ reiškia, kad wiki yra
kad obrobka praeis mažiau dėl procesoriaus likimo. Būtina laiduoti už tai, ką _pasaulio_ i
_device_ negali iškviesti vieno iš vieno ir negali paskambinti sau.

Taip pat CUDA kalba gali turėti mažai funkcijų dirbant su vaizdo atmintimi: cudafree
(atminties keitimas tarp GDDR ir RAM), cudamemcpy ir cudamemcpy2D (kopija
atmintis tarp GDDR ir RAM) ir „cudamalloc“ (atminties rodinys).

Visos kodo programos turi būti sukompiliuotos iš CUDA API pusės. Ant burbuolės paimama
kodas, priskyrimai imtinai centriniam procesoriui ir poddaetsya
standartinis kompiliavimas ir kiti kodai, grafinio adapterio priskyrimai,
perrašyti į PTX (labai spėliojantis surinkėjas) už
parodydamas galimus atleidimus. Po visų šių „šokių“ lieka likutis
komandų vertimas (transliavimas) į suprasti GPU/CPU mov.

Skambinkite vestuvėms

Beveik visi programavimo aspektai yra aprašyti dokumentacijoje, kuri yra
kartu su vairuotoju ir dviem programomis, taip pat mažmenininkų svetainėje. Rozmiras
neskaityti straipsnių, juos aprašyti (skaitymas gali būti pridėtas
mažai pastangų ir dirbti su medžiaga savarankiškai).

CUDA SDK naršyklė buvo išplėsta specialiai pradedantiesiems. Be-yaky bazhayuchy gali
pajuskite lygiagrečių skaičiavimų galią savo kailiu (geresnis pakartotinis patikrinimas
stabilumas – robotas naudojamas be artefaktų ir dėmių). Priedas gali
daugybė šou mini programų (61 "testas"). Į odą dosvіdu є
Programos kodo ataskaitų dokumentacija ir PDF failai. Jūs matote, kad žmonės
pateikti su savo langais naršyklėje, užsiima rimtu darbu.
Čia taip pat galite pakeisti roboto ir procesoriaus bei vaizdo plokštės greitį apdorojimo valandai
Danich. Pavyzdžiui, turtingų masyvų nuskaitymas su vaizdo plokšte GeForce 8800
GT 512 MB su 256 gijų bloku cirkuliuoja per 0,17109 milisekundės.
Technologijos neatpažįsta SLI tandemų, todėl pučiate trijulę,
prieš robotą įjunkite „suporavimo“ funkciją, kitaip CUDA veiks tik vieną
prietaisas. dviejų branduolių AMD Athlon 64X2(pagrindinis dažnis 3000 MHz) tą patį patvirtinimą
praeina per 2,761528 milisekundės. Išsiaiškink, kad G92 yra didesnis nei 16 kartų
shvidshe "akmuo" AMD! Jak Bachish, sistema toli gražu nėra kraštutinė
tandemas su nemylimaisiais masinėse operacinėse sistemose, rodančiais blogus dalykus
rezultatus.

Krіm naršyklė іsnuє mažas korisnyh suspіlstvo programas. Adobe
savo gaminius pritaikė naujoms technologijoms. Dabar „Photoshop CS4“ atnaujinta
pasaulyje pergalingi grafinių adapterių ištekliai (būtina pateikti specialius
prijungti). Su tokiomis programomis, tokiomis kaip „Badaboom Media Converter“ ir „RapiHD“, galite
Galimybė iššifruoti vaizdo įrašą į MPEG-2 formatą. Dėl apdorojimo garsas yra blogas
pidide nemokama „Accelero“ programa. Programinės įrangos skaičius, patobulinta CUDA API,
aišku, augti toliau.

O tuo metu...

O kol jūs skaitote šią medžiagą, rūpestingi procesoriaus darbuotojai
išplėsti savo technologijas, kad būtų skatinamas GPU iš procesoriaus. 3 pusė AMDūsai
suprato: jie turi didingą dosvіd, nabutiy iš karto iz ATI.

„Mikroprietaisų“ kūrimas „Fusion“ susideda iš daugybės branduolių
kodiniu pavadinimu Buldozeris ir vaizdo lustas RV710 (Kong). Jūsų tarpusavio santykiai bus
zdіysnyuvatsya už rahunok razrezheny padangų HyperTransport. Pūdymo vaizdas
branduolių skaičius ir dažnio charakteristikos AMD planuoja sukurti visą kainą
„akmenų“ hierarchija. Taip pat planuojama konvertuoti procesorius kaip nešiojamiesiems kompiuteriams (Falcon),
ir daugialypės terpės programėlėms („Bobcat“). Be to, pats technologijų sąstingis
kilnojamieji ūkiniai pastatai bus pirmoji kanadiečių užduotis. Z plėtros komitetas
Lygiagretus tokio „kamintsivo“ sąstingio išvardijimas gali būti dar populiaresnis.

Intel Trys per valandą jūsų Larrabee. Produktai AMD,
yakscho neprisigerti, pasirodyti parduotuvių lentynose kaip 2009 - ant burbuolės
2010 roko. O priešininko sprendimas pamatyti Dievo šviesą mažiau tikėtinas dviese
likimas.

Larrabee turi daug (skaityti – šimtus) branduolių. Ant burbuolės
Na, prekes matote, atsipirko už 8 - 64 branduolius. Kvapas labiau panašus į Pentium, ale
melžti tai stipriai pervargiama. Skin branduolys gali turėti 256 kilobaitų talpyklą
(Su metais jogos augimas padidės). Vzaymosv'yazok
1024 bitų dvikryptė žiedinė magistralė. Panašu, kad „Intel“ yra „vaikas“
gera praktika naudojant „DirectX“ ir „Open GL“ API („Yabluchnikov“), nieko iš to nėra
programinės įrangos įvestis nereikalinga.

Ir kodėl aš visas tse rozpoviv? Akivaizdu, kad Larrabee ir Fusion nematyti
stacionarių procesorių iš rinkos
vaizdo plokštės. Žaidėjams ir ekstremaliems žmonėms, tarp svajonių, kaip ir anksčiau, atsikratyti
turtingo branduolio procesorius ir tandemas su aukščiausios klasės VGA lipdukais. Ale tie, kurie naviguoja
procesorių įmonės pereis prie lygiagretaus atsiskaitymo pagal principus,
analogiškai GPGPU, jau sodriai kalbėti apie ką. Zokrema apie tuos, kurie yra
technologija, kaip ir CUDA, gali turėti teisę naudoti ir galbūt būti
jau populiarus.

Mažas gyvenimo aprašymas

Lygiagretusis atsiskaitymas vaizdo plokštės pagalba yra tik gera priemonė
programuotojo roboto rankose. Vargu ar chi procesoriai pagal choli іz Moore'o dėsnį
sustabdyti. Įmonės NVIDIA atsigulkite, kad eitumėte dar vieną ilgą kelią
slysti į savo API masę (tą patį galima pasakyti ir apie vaiką ATI/AMD).
Koks būsi, parodyk ateitį. Taip pat CUDA grįš :).

P.S. Visiems programuotojams ir užkliuvusiems žmonėms rekomenduoju pasižiūrėti
kita "virtuali hipoteka":

NVIDIA oficiali svetainė ir svetainė
GPGPU.com. Visi
duota informacija - mano anglų kalba, bet ačiū, nenorėčiau
kinų. Taigi taip ir toliau! Esu tikras, kad autorius norėjo jums šiek tiek padėti
nuoširdžios CUDA žinių iniciacijos!

І priskyrimai, skirti versti pagrindinio kompiuterio kodą (galvos, rakto kodas) ir įrenginio kodą (aparatinės įrangos kodą) (failus su plėtiniais.cu) į objektų failus, priedus baigiamos programos ar bibliotekos sudarymo procese bet kurioje programavimo aplinkoje, pavyzdžiui, NetBeans.

CUDA architektūroje nugalėjo tinklo atminties modelis, srautų klasterio modelis ir SIMD instrukcijos. Yra ne tik didelio našumo vaizdo plokštės, bet ir kitos mokslinės sąskaitos už nVidia vaizdo plokštes. Šio tyrimo rezultatai plačiai naudojami CUDA įvairiose galuose, įskaitant astrofiziką, biologijos ir chemijos skaičiavimus, upių dinamikos modeliavimą, elektromagnetines sąveikas, kompiuterinę tomografiją, seisminę analizę ir kt. CUDA turi galimybę prisijungti prie programų, naudojančių OpenGL ir Direct3D. CUDA yra kelių platformų programinė įranga, skirta tokioms operacinėms sistemoms kaip Linux, Mac OS X ir Windows.

2010 m. kovo 22 d. nVidia išleido CUDA Toolkit 3.0, kuris yra OpenCL palaikymo atgaivinimas.

Nuosavybė

CUDA platforma pirmą kartą pasirodė rinkoje išleidus aštuntos kartos NVIDIA G80 lustą ir tapo visose būsimose grafikos lustų serijose, kurios yra greitesnių GeForce, Quadro ir NVidia Tesla šeimose.

Pirmoji serija, palaikanti CUDA SDK, G8x, yra mažas 32 bitų vieno tikslumo vektorinis procesorius, kuris palaiko CUDA SDK kaip API (CUDA palaiko dvigubo filmo Ci tipą, proteo tikslumas sumažintas iki 32 bitų slankiojo punktas). Žemesni GT200 procesoriai gali palaikyti 64 bitų tikslumą (tik SFU), tačiau našumas yra žymiai didesnis, o 32 bitų tikslumas mažesnis (per juos yra tik du SFU vienam odos srautiniam daugiaprocesoriui, o skaliariniai procesoriai - visi). Grafikos procesorius organizuoja aparatinės įrangos srauto turtingumą, kuris leidžia išnaudoti visus grafikos procesoriaus išteklius. Tokiu būdu atsiranda galimybė perkelti fizinio nuorodos funkcijas į grafinę (diegimo pavyzdys - nVidia PhysX). Taip pat yra plačios galimybės panaudoti grafinį kompiuterio nustatymą sulankstomam negrafiniam skaičiavimui: pavyzdžiui, skaičiuojant biologiją ir kitus mokslo triukus.

Perevagi

Be tradicinio požiūrio į visuotinio pripažinimo apskaičiavimo organizavimą dėl papildomų grafinių API galimybių, CUDA architektūra taip pat gali laimėti šioje galerijoje:

Mainai

• Visos funkcijos, kurios rodomos plėtinyje, nepalaiko rekursijos (CUDA Toolkit 3.1 versijoje palaiko rekursijos indikatorius) ir gali veikti kituose mainuose

GPU patobulinimai ir grafikos patobulinimai

Papildinio vertimas kaip Nvidia įgalinta įranga iš deklaruoto naujo CUDA technologijos palaikymo yra parodytas oficialioje Nvidia svetainėje: CUDA-Enabled GPU Products (anglų kalba).

Tiesą sakant, šią valandą kompiuterių aparatinės įrangos rinkoje CUDA technologijos palaikymas bus užtikrintas tobulėjančiais periferiniais įrenginiais:

Specifikacijos versija	GPU	Vaizdo plokštės
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, /3/18/30 /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 staliniams kompiuteriams 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 Vaizdo plokštė GeForce 8800 GTS 512 GeForce 8800 GT GeForce 8600 GTS GeForce 8600 GT GeForce 8500 GT GeForce 8400GS

Nvidia GeForce mobiliesiems kompiuteriams 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 9300MG 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 staliniams kompiuteriams 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 žemo profilio QuadroFX 370 Quadro FX 370 žemo profilio 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 IV modelis

Nvidia Quadro mobiliesiems kompiuteriams 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

• Modeliai Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 leidžia atlikti GPU skaičiavimus kintamu tikslumu.

Įvairių versijų savybės ir specifikacijos

Funkcijų palaikymas (neįtrauktos funkcijos yra palaikomos visos skaičiavimo galimybės)	Skaičiavimo galimybė (versija)
	1.0	1.1	1.2	1.3	2.x
32 bitų žodžiai pasaulinėje atmintyje	Sveiki	Taigi
slankiojo kablelio reikšmės pasaulinėje atmintyje
Veikiančios sveikųjų skaičių atominės funkcijos 32 bitų žodžiai bendrojoje atmintyje	Sveiki		Taigi
atomicExch(), veikiantis 32 bitų slankiojo kablelio reikšmės bendrojoje atmintyje
Veikiančios sveikųjų skaičių atominės funkcijos 64 bitų žodžiai pasaulinėje atmintyje
Metimo balsavimo funkcijos
Dvigubo tikslumo slankiojo kablelio operacijos	Sveiki			Taigi
Atominės funkcijos, veikiančios 64 bitų sveikųjų skaičių reikšmės bendrojoje atmintyje	Sveiki				Taigi
Slankaus kablelio atominis papildymas veikia 32 bitų žodžiai pasaulinėje ir bendrojoje atmintyje
_balso biuletenis()
_threadfence_system()
_syncthreads_count(), _syncthreads_and(), _syncthreads_or()
Paviršiaus funkcijos
3D siūlų bloko tinklelis

Techninės specifikacijos	Skaičiavimo galimybė (versija)
	1.0	1.1	1.2	1.3	2.x
Didžiausias sriegių blokų tinklelio matmuo	2				3
Didžiausias sriegių blokų tinklelio x, y arba z matmuo	65535
Maksimalus sriegio bloko matmuo	3
Maksimalus bloko x arba y matmuo	512				1024
Didžiausias bloko z matmuo	64
Maksimalus gijų skaičius bloke	512				1024
Metmenų dydis	32
Didžiausias nuolatinių blokų skaičius viename daugiaprocesoriuje	8
Didžiausias nuolatinių kelių procesorių deformacijų skaičius	24		32		48
Didžiausias nuolatinių gijų skaičius viename daugiaprocesoriuje	768		1024		1536
32 bitų kelių procesorių registrų skaičius	8K		16 tūkst		32 tūkst
Didžiausias bendrai naudojamos atminties kiekis vienam procesoriui	16KB				48KB
Bendrinamų atminties bankų skaičius	16				32
Vietinės atminties kiekis vienoje gijoje	16KB				512KB
Pastovus atminties dydis	64KB
Talpyklos darbo rinkinys, skirtas daugiaprocesoriui nuolatinei atminčiai	8KB
Talpyklos darbo rinkinys, skirtas daugiaprocesoriui tekstūros atminčiai	Priklauso nuo įrenginio, nuo 6 KB iki 8 KB
Maksimalus plotis 1D tekstūrai	8192				32768
Maksimalus plotis 1D tekstūrai nuoroda susieta su tiesine atmintimi	2 27
Maksimalus plotis ir sluoksnių skaičius 1D sluoksniuotos tekstūros nuorodai	8192x512				16384x2048
Maksimalus plotis ir aukštis 2D susieta tekstūros nuoroda linijinė atmintis arba CUDA masyvas	65536 x 32768				65536 x 65535
Didžiausias plotis, aukštis ir skaičius sluoksnių 2D sluoksniuotos tekstūros nuorodai	8192 x 8192 x 512				16384x16384x2048
Maksimalus plotis, aukštis ir gylis 3D tekstūros nuorodai, susietai su linijine atmintis arba CUDA masyvas	2048x2048x2048
Maksimalus tekstūrų skaičius, kuris gali būti susietas su branduoliu	128
Maksimalus plotis 1D paviršiui nuoroda susieta su CUDA masyvu	Ne palaikoma				8192
Maksimalus plotis ir aukštis 2D paviršiaus nuoroda, susieta su CUDA matrica					8192 x 8192
Maksimalus paviršių skaičius gali būti susietas su branduoliu					8
Maksimalus instrukcijų skaičius vienam branduolys	2 mln

Užpakalis

CudaArray* cu_array; tekstūra< float , 2 >tex; // Paskirstyti masyvą cudaMalloc( & cu_array, cudaCreateChannelDesc< float>(), plotis Aukštis); // Nukopijuoti vaizdo duomenis į masyvą cudaMemcpy( cu_array, image, plotis* aukštis, cudaMemcpyHostToDevice) ; // Susieti masyvą su tekstūra cudaBindTexture(teksas, cu_masyvas) ; // Paleisti branduolio dim3 blockDim(16, 16, 1); dim3 gridDim(width/blockDim.x, height/blockDim.y, 1); branduolys<<< gridDim, blockDim, 0 >>> (d_duomenys, plotis, aukštis); cudaUnbindTexture(tekstas); __global__ void branduolys(float * odata, int aukštis, int plotis) ( unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float texfetch ( tex, x, y);duomenys[y* plotis+ x] = c;

Importuoti pycuda.driver kaip drv importuoti 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 ("padauginti_juos") a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 (a) drvly .Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

CUDA kaip dalykas universitetuose

2009 m. pradžioje pasirodysiantis CUDA programinės įrangos modelis bus įdiegtas 269 universitetuose visame pasaulyje. Rusijoje pirmieji CUDA kursai skaitomi Sankt Peterburgo politechnikos universitete, Jaroslavlio valstybiniame universitete. P. G. Demidovas, Moskovskis, Nižnij Novgorodas, Sankt Peterburgas, Tverskojė, Kazanė, Novosibirskis, Novosibirsky Holden technikos universitetas Permsky UNIVITIES, MIENENTY UNIVITITICITIONALY INITY ONITYAS. Baumanas, RKhTU im. Mendelevas, Rusijos mokslų akademijos tarpregioninis superkompiuterių centras. Be to, 2009 m. gruodį buvo pranešta apie pirmojo Rusijoje mokslo ir apšvietimo centro „Paralelinis skaičiavimas“ darbo pradžią, kuris buvo surengtas Dubnos mieste, iki kurio datos buvo konsultuojamasi dėl 2009 m. lankstymo skaičiavimai ant GPU.

Ukrainoje CUDA kursai skaitomi Kijevo sistemos analizės institute.

Posilannya

Oficialūs ištekliai

• CUDA zona (rusų kalba) – oficiali CUDA svetainė
• CUDA GPU Computing (anglų k.) – oficialūs interneto forumai, skirti CUDA skaičiavimams

Neoficialūs ištekliai

Tomo aparatūra
• Dmitro Čekanovas. nVidia CUDA: vaizdo plokščių mokesčiai ar procesoriaus mirtis? . Tom's Hardware (2008 m. rugsėjo 22 d.).
• Dmitro Čekanovas. nVidia CUDA: GPU testavimas pagrindinei rinkai. Tom's Hardware (antradienis, 2009 m. 19 d.).
iXBT.com
• Oleksijus Berilas. NVIDIA CUDA – negrafinis skaičiavimas grafikos procesoriuose. 1 dalis. iXBT.com (2008 m. rugsėjo 23 d.). Suarchyvuota 2012 m. vasario 4 d. Patikslinta 2009 m. rugsėjo 20 d.
• Oleksijus Berilas. NVIDIA CUDA – negrafinis skaičiavimas grafikos procesoriuose. 2 dalis. iXBT.com (2008 m. liepos 22 d.). - Uždėkite NVIDIA CUDA pleistrą. Suarchyvuota 2012 m. vasario 4 d. Patikslinta 2009 m. rugsėjo 20 d.
Kiti ištekliai
• Boreskovas Oleksijus Viktorovičius. CUDA pagrindai (2009 m. rugsėjo 20 d.). Suarchyvuota 2012 m. vasario 4 d. Patikslinta 2009 m. rugsėjo 20 d.
• Vladimiras Frolovas.Įvadas į CUDA technologiją. Žurnalas Merezhevy „Kompiuterinė grafika ir multimedija“ (2008 m. gruodžio 19 d.). Suarchyvuota 2012 m. vasario 4 d. Peržiūrėta 2009 m. birželio 28 d.
• Igoris Oskolkovas. NVIDIA CUDA yra prieinamas bilietas į didžiulių skaičių pasaulį. Kompiuteriai (2009 m. balandžio 30 d.). Žiūrėta 2009 m. gegužės 3 d.
• Vladimiras Frolovas.Įvadas į CUDA technologiją (2009 m. rugsėjo 1 d.). Suarchyvuota 2012 m. vasario 4 d. Peržiūrėta 2010 m. balandžio 3 d.
• GPGPU.ru. Wikoristannya vaizdo plokštės skaičiavimams
• . Lygiagretusis skaičiavimo centras

Pastabos

Div. taip pat

Nvidia
Grafika procesoriai	Anksti NV1 NV2 RIVA šeima TNT TNT2 Šeima