În același specificatorii __host__ __device__ și pot fi utilizate împreună (ceea ce înseamnă că funcția corespunzătoare poate fi realizată pe GPU și CPU - codul corespunzător va fi generat automat de compilator pentru ambele platforme). Specificatorii __global__ și __host__ nu pot fi utilizați împreună.
Specificatorul __global__ denotă un kernel, iar funcția corespunzătoare trebuie să returneze o valoare de tip void.
__global__ void myKernel (float * a, float * b, float * c)
int index = threadIdx.x;
Funcțiile care sunt executate pe GPU (__device__ și __global__) sunt supuse următoarelor restricții:
Următorii specificatori sunt utilizați pentru a specifica destinația de plasare în memoria GPU: __ device__, __constant__ și __shared__. O serie de restricții se impun și asupra utilizării acestora:
Următoarele variabile speciale sunt adăugate la limba:
- gridDim este dimensiunea grilei (are tipul dim3)
- blockDim - dimensiunea blocului (are tipul dim3)
- blockIdx - indicele blocului curent din grilă (are tipul uint3)
- threadIdx - indicele firului curent din bloc (are tipul uint3)
- warpSize este dimensiunea urzei (este de tip int)
Limbajul adăugat 1 vector / 2/3/4-dimensional al tipurilor de bază - char1, char2, char3, char4, uchar1, uchar2, uchar3, uchar4, short1, short2, short3, short4, ushort1, ushort2, ushort3, ushort4, INT1 , INT2, INT3, INT4, uint1, uint2, uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, float1, float2, float3, float2 și double2.
Referința la componentele vectorului trece prin numele - x, y, z și w. Pentru a crea valori vectoriale ale unui anumit tip, utilizați funcția make_
int2 a = make_int2 (1, 7);
float3 u = make_float3 (1, 2, 3.4f);
Rețineți că pentru aceste tipuri (spre deosebire de limbajul shader GLSL / Cg / HLSL), operațiile componentelor vectoriale nu sunt acceptate. nu puteți adăuga doar doi vectori folosind operatorul "+" - acest lucru trebuie făcut în mod explicit pentru fiecare componentă.
De asemenea, pentru a seta dimensiunea, tastați dim3 se bazează pe tipul uint3, dar are un constructor normal care inițializează toate unitățile nesignificate.
Directiva privind apelul kernelului
Pentru a rula nucleul de pe GPU, se utilizează următoarea construcție:
kernelName <<
De asemenea, funcția __syncthreads este adăugată în limba C, care sincronizează toate firele blocului. Comanda de la acesta va fi returnată numai când toate firele acestui bloc apelează această funcție. Ie când tot codul care merge înainte de acest apel este deja executat (și, prin urmare, puteți să vă bazați în siguranță pe rezultatele acestuia). Această funcție este foarte convenabilă pentru organizarea unei activități fără conflicte cu memoria partajată.
CUDA susține, de asemenea, toate funcțiile matematice din biblioteca standard C, dar din punct de vedere al performanței este mai bine să folosești analogii flotați (mai degrabă decât dublu) - de exemplu sinf. În plus, oferă CUDA un set suplimentar de funcții matematice (__sinf, __powf etc.) oferă o acuratețe mai mică, dar performanțe semnificativ mai mare decât sinf, powf etc.
API gazdă CUDA
API-ul CUDA pentru CPU (gazdă) acționează în două forme - API-ul CUDA drievr de nivel CUDA și API-ul CUDA runtime (implementat prin driverul API CUDA). În aplicația dvs., puteți utiliza doar una dintre ele, apoi vom lua în considerare API-ul CUDA runtime, deoarece este mai simplu și mai convenabil.
Toate funcțiile API ale driverului CUDA încep cu prefixul cu și toate funcțiile CUDA API runtime încep cu prefixul cuda. Fiecare dintre aceste API oferă un nucleu set de funcții de bază, cum ar fi prin toate dispozitivele disponibile (GPU), care lucrează cu contexte și fire, lucru cu memorie GPU, interacțiunea cu OpenGL și D3D (sprijinit numai versiunea nouă de DirectX).
API-ul CUDA runtime nu necesită inițializare explicită - aceasta se produce automat la primul apel al oricărei funcții. Un punct important în lucrul cu CUDA este că multe funcții API sunt asincron, adică controlul este returnat chiar înainte de finalizarea efectivă a operațiunii solicitate.
Operațiile asincrone includ:
- pornind kernel-ul
- Funcțiile de copiere de memorie ale căror nume se termină în Async
- funcția de copiere a dispozitivului de memorie <-> dispozitiv
- funcții de inițializare a memoriei.
CUDA susține sincronizarea prin fluxuri (fluxuri) - fiecare fir specifică o secvență de operații efectuate într-o ordine strict definită. În acest caz, ordinea execuției operațiunilor între diferite fluxuri nu este definită strict și se poate schimba.
Fiecare funcție API CUDA (cu excepția inițierii kernelului) returnează o valoare de tip cudaError_t. Dacă funcția reușește, cudaSuccess este returnat, altfel codul de eroare este returnat.
Pentru a obține o descriere a erorii sub forma unui șir prin codul său, puteți utiliza funcția cudaGetErrorString:
char * cudaGetErrorString (codul cudaError_t);
De asemenea, puteți obține ultimul cod de eroare utilizând funcția cudaGetLastError:
Vă rugăm să rețineți că, din cauza executarea asincronă a multor apeluri pentru a obține codul de eroare este mai bine să utilizați cudaThreadSynchronize funcția, care este în așteptare pentru finalizarea GPU toate cererile transmise și returnează o eroare în cazul în care una dintre aceste cereri a dus la eroarea.
Lucrul cu memoria în CUDA
Cel mai simplu mod de a aloca și memoria liberă (este vorba exclusiv despre memoria GPU, cu numai memorie liniară, un alt tip - CUDA-matrice vor fi discutate în articolul următor) este de a utiliza funcțiile cudaMalloc, cudeMallocPitch și cudaFree.
float * devPtr; // indicatorul pentru memoria dispozitivului
// alocați memoria liniară pentru 256 de flotoare
cudaMalloc ((void **) devPtr, 256 * sizeof (float));
Pentru desemnarea capacității CUDA utilizează conceptul de Compute Capabilitatea, exprimată printr-o pereche de numere - MAJOR.MINOR. Primul număr se referă la versiunea de arhitectură la nivel mondial, al doilea - o mica schimbare. Deoarece GeForce GPU 8800 Ultra / GTX / GTS sunt egale Compute Capability 1.0, GPU GeForce 8800 GT / GS și GeForce 9600 GT sunt egale Compute Capability 1.1, GPU GeForce GTX 260 si GeForce GTX 280 au egal Compute Capability 1.3.
Compute Capacitatea 1.1 suportă operații atomice de cuvinte pe 32 de biți în memoria globală, Compute Capacitatea 1.2 suportă operații atomice în partajată memorie și operații atomice de cuvinte pe 64 de biți în memoria globală, Compute Capacitatea 1.3 suportă operații asupra numerelor de tip dublu.
Mai jos este codul sursă al unui program simplu care conține toate GPU-urile disponibile și principalele caracteristici ale acestora.
int principal (int argc, char * argv [])
printf ("Dispozitive găsite \ n", deviceCount);
pentru (dispozitiv int = 0; dispozitiv cudaGetDeviceProperties ( devProp, dispozitiv); printf ("Dispozitiv. \ n", dispozitiv); printf ("Captura de calcul .... \ n", devProp.major, devProp.minor); printf ("Nume.% s \ n", devProp.name); printf ("Memoria globală totală \ n", devProp.totalGlobalMem); printf ("Memorie partajată pe bloc :. \ n", devProp.sharedMemPerBlock); printf ("Registrele per bloc ...", devProp.regsPerBlock); printf ("Warp size ... n", devProp.warpSize); printf ("Fire maxime per bloc ...", devProp.maxThreadsPerBlock); printf ("Memorie totală constantă \ n", devProp.totalConstMem); Să examinăm câteva exemple simple de utilizare a CUDA, demonstrând tehnicile de bază ale lucrului cu acesta. Cel mai simplu exemplu ar fi o creștere simplă a fiecărui element dintr-o matrice unidimensională pe unitate - programul incr.cu. __global__ void incKernel (date float *) Kernelul este cel mai simplu aranjat - fiecare fir corespunde unui fir, blocurile și grila sunt unidimensionale. Kernelul (funcția incKernel) primește numai un pointer la o matrice cu date din memoria globală. Sarcina nucleului este de a determina exact ce element corespunde threadului dat de threadIdx și blockIdx și de a-l crește exact. Deoarece ambele blocuri și grilă sunt unidimensionale, numărul firului va fi definit ca numărul blocului înmulțit cu numărul de fire în bloc, plus numărul firului din interiorul blocului, adică blockIdx.x * blockDim.x + threadIdx.x. Funcția principală este oarecum mai complicată: trebuie să pregătească o matrice cu date în memoria CPU, apoi, folosind cudaMalloc, să aloce memoria pentru o copie a matricei cu date în memoria globală (DRAM GPU). Apoi, datele sunt copiate de către funcția cudaMemcpy din memoria CPU în memoria globală GPU. După ce datele sunt copiate în memoria globală, puteți porni kernel-ul pentru prelucrarea datelor și, după apelul său, copiați rezultatele calculului înapoi din memoria globală a GPU-ului în memoria procesorului. Multiplicarea a două matrice este cea mai simplă abordare Următorul exemplu va fi mai greu (și mai relevant) - vom analiza utilizarea CUDA pentru a multiplica două matrice N * N pătrate. Să presupunem că avem două matrici pătrate A și B de dimensiunea N * N (vom presupune că N este un multiplu de 16). Cea mai simplă versiune utilizează un fir pentru fiecare element al matricei rezultate C, în timp ce firul extrage toate elementele necesare din memoria globală și produce calculele necesare. Elementul ci, j al produsului a două matrici A și B este calculat de următorul fragment pseudo-cod: Astfel, pentru a calcula un element al produsului matricelor, este necesar să se efectueze operații aritmetice 2 * N și citiri 2 * N din memoria globală. Este clar că în acest caz principalul factor limitator este viteza de acces la memoria globală, care este foarte scăzută. Modul în care firele individuale sunt grupate în blocuri nu este importantă și nu are un efect semnificativ asupra performanței, care este foarte scăzut în acest caz. Următoarea este lista programelor corespunzătoare. Înmulțirea a două matrice utilizând memoria partajată Puteți îmbunătăți semnificativ performanța programului nostru utilizând memoria partajată. Pentru a face acest lucru, împărțim matricea rezultată în submatricule 16 * 16, fiecare bloc va fi calculat de un bloc. Rețineți că pentru a calcula o astfel de submatrică sunt necesare doar "benzi" mici ale matricelor A și B. Din nefericire, copierea completă a acestor "benzi" în memoria partajată este aproape imposibilă din cauza cantității destul de mici de memorie partajată. Prin urmare, putem proceda într-un alt mod - am împărțit aceste "benzi" în 16 * 16 matrici și calculam submatricea produsului matrice în etape N / 16. În acest scop, observăm că calculul elementului ci, j poate fi rescris după cum urmează, folosind partiționarea benzilor în submatricule pătrate pentru pas în 0..N / 16: c [i] [j] + = a [i] [k + pas * 16] * b [k + Rețineți că pentru fiecare valoare pas, valorile din matricile A și B sunt luate din două submulțimi 16 * 16. De fapt, benzile din Fig. 4 sunt pur și simplu împărțite în submăsuri pătrată și la fiecare valoare de pas corespunde unei astfel de submatrice A și a unei submatrice B. La fiecare pas vom încărca în memoria partajată pe un 16 * 16 submatrice A și unul de 16 * 16 submatricii B. În continuare, vom calcula valoarea corespunzătoare a acestora pentru elemente ale produsului, apoi încărcați următoarele 16 * 16 submatricele, etc. În acest caz, la fiecare etapă, un fir atacă exact un element din fiecare dintre matricile A și B și calculează suma corespunzătoare a termenilor. La sfârșitul tuturor calculelor, elementul este scris în matricea finală. Vă rugăm să rețineți că, după descărcarea de elementele A și B au nevoie pentru a sincroniza fire prin apelarea __synchronize la momentul la începutul așezări ar fi fost deja încărcate toate elementele necesare (descărcabile rămase fire de bloc). În mod similar, la sfârșitul tratamentului au fost adăugate submatricile și au nevoie de sincronizare înainte de a descărca aceste, de asemenea, (pentru a se asigura că actuala 16 * 16 submatricea nu mai este necesară și încărcare nouă posibilă este). Codul sursă corespunzător este prezentat mai jos.
Totul despre managerii de fotbal
Stiri despre proiect
La ce lucrez acum pe proiecte
Planuri pentru proiecte, ce trebuie făcut și doar planuri
Evoluții care au făcut sau au făcut odată, precum și gânduri și note.
Dezvoltare web
Dezvoltarea de programe sau programe, funcții sau blocuri.
Diverse note, gânduri despre proiecte.
Tot ce nu este în secțiunile anterioare
Știri, fapte din lume.
Revizuirea programelor care sunt interesante în opinia mea
Revizuirea în lumea fierului
Diferite articole legate de programare, dezvoltare de aplicații.Noul tag
KarpOlya - note ale profesorului de biologie și chimie
Manager de fotbal online
Articole similare
Trimiteți-le prietenilor: