Stencil tampon pe cuda

Stencil tampon sau tampon de stencil este de obicei folosit în grafică (OpenGL, DirectX) pentru a masca câțiva pixeli pe imaginea apelului pixelului pentru anumite zone ale imaginii. Textul evidențiază în mod special faptul că testul de stencil este efectuat înainte de a fi apelat shaderul pixelilor și astfel, în acele locuri în care lipsește imaginea, shaderul pixelului nu va fi apelat deloc și nu va mai fi efectuată nicio extra lucrare.







Din punct de vedere fizic, stencilul este stocat pe GPU-ul în același buffer, unde adâncimea este stocată și poate fi de diferite formate. De exemplu, formatul D3DFMT_D24S8 utilizat cel mai frecvent înseamnă că 24 biți sunt alocați în tamponul de spate la o adâncime și 8 biți la șablon. În acest articol, vom folosi simplificarea și presupunem că tamponul stencil stochează pentru fiecare pixel (sau flux) doar un bit. Dacă bitul este 1, atunci pixelul (fluxul) este activ. Dacă este 0, este inactivă. Acest lucru va salva o mică memorie și va simplifica prezentarea.

Stencil Test este adesea folosit pentru a construi reflecții în acest fel:

Stencil tampon pe cuda

Figura 1. Tamponul de stencil este necesar pentru a masca reflexiile în acele locuri unde nu există de fapt (ca în imaginea din dreapta).

1. Ștergeți tamponul de stencil cu zerouri.

2. Includeți înregistrarea în tamponul de stencil și trageți un plan în el, în raport cu care vom lua în considerare reflecția. Întotdeauna scriem una. Se pare că imaginea binară a oglinzii noastre este stocată în tampon de mască (adică acolo unde este o oglindă vor fi stocate unități și unde nu există nicio oglindă - zerouri).

3. Reflectați toată geometria față de plan folosind o matrice specială și trageți-o, inclusiv un test de stencil. Astfel, acolo unde era o oglindă în imagine, se va produce o reflecție. Și unde nu există, nimic nu se va schimba.

Implementarea de software pe CUDA

Din păcate, în kud, ca și în celelalte tehnologii "compute" (DX11 CS, OpenCL), mecanismul de testare a șablonului nu este disponibil. În același timp, acest lucru este foarte util, mai ales dacă calculele dvs. sunt implementate ca o conductă lungă de mai multe nuclee (adesea destul de mici). Să presupunem că aveți N fire.

Stencil tampon pe cuda

De exemplu, această situație apare atunci când se realizează trasarea razei pe cale. La o adâncime de reflexie de aproximativ 5, în unele scene, mai puțin de 10% din fluxuri vor fi active la ultimul nivel.

Pentru a nu efectua lucrări pentru fire inactive, probabil veți pune un steag în orice tampon și veți verifica dacă acest steag este 0, apoi nu faceți nimic.

uintactiveFlag = a_flags [tid];

În acest articol, se recomandă să se păstreze un singur bit în tamponul stencil pentru 1 flux și să se evite transferul de date în masă pe magistrala (sau cel puțin să le reducă în mod semnificativ, utilizând eficient cache-ul).







Așadar, am creat o dimensiune a buffer-ului stencil exact pe (N / 32) * sizeof (int) byte. Și le legăm textura.

cudaBindTexture (0, stencil_tex, m_stencilBuffer, N * sizeof (int) / 32);

Textura în sine este declarată în unele antet (fișier .h) după cum urmează:

textură stencil_tex;

În continuare, în același fișier, declarăm o astfel de matrice auxiliară:

0x00000001, 0x00000002, 0x00000004, 0x00000008, 0x00000010, 0x00000020, 0x00000040, 0x00000080,

În această matrice sunt stocate măști, cu care vom face logică pentru a obține rapid fluxul de biți dorit. Adică, obțineți exact bitul al cărui număr este egal cu numărul firului din urzeală. Iată cum arată testul de stencil:

dacă {! (tex1Dfetch (stencil_tex, (tid) >> 5) g_stencilMask [(tid) 0x1f]))

Pentru nucleele care au citit doar tamponul de stencil, macroul trebuie aplicat la începutul kernel-ului după cum urmează:

__global__void my_kernel (...)

uinttid = blocDim.x * blocIdx.x + threadIdx.x;

În practică (GTX560), un astfel de test de stencil este cu aproximativ 20-25% mai rapid decât o verificare simplă a formularului:

uintactiveFlag = a_flags [tid];

Deci, rămâne doar să scrieți tamponul de stencil. Mai întâi, citiți valoarea pentru tot în warp din bufferul stealth la variabila activWarp; Apoi fiecare fir primește bitul său din această variabilă utilizând o logică și o stochează în variabila activă. La sfârșitul kernelului, vom colecta din toate variabilele active pentru o valoare warp dată într-un uint de 32 de biți, iar urzeala zero a firului va scrie rezultatul înapoi în memorie.

// (tid 0x1f) la fel ca (tid% 32)

__global__void my_kernel2 (..., uint * a_stencilBuffer)

uinttid = blocDim.x * blocIdx.x + threadIdx.x;

uint activWarp = a_stencilBuffer [tid >> 5];

dacă (activeWarp == 0) // toate firele în warp inactive

// fiecare fire va stoca un anumit bit din grupul de 32 fire

uint active = activeWarp g_stencilMask [tid0x1f];

Funcționează așa. În cazul în care urzeala a terminat activitatea (de exemplu, este la sfârșitul în timp ce ciclul sau a fost reclined testul matrita), acesta fură porțiunea următoare de lucru pentru sine, creșterea globală contra 32. Counter dintr-o dată indică cât de mult operație este lăsată liberă.

La firele G80 la fel de persistente nu vor fi implementate, din cauza absenței operațiilor atomice. Dar puteți face o bucla ca "pentru (int i = 0; i<8;i++) doMyWork(i);” для того, чтобы увеличить количество работы, выполняемое одним warp-ом. На GT200 в некоторых случаях, использование persistent threads давало прирост производительности до 2 раз.

Testarea șablonului

De fapt, pentru nevoile de raytracing, un astfel de tampon de stencil a apărut cu succes. Dacă vă îngropați în vid, pe GTX560 este posibil să primiți aproximativ 4 miliarde de apeluri pe nuclee pe secundă (adică 4 miliarde de apeluri goale pe secundă). Odată cu creșterea adâncimii urmei, performanța practic nu a căzut (mai degrabă, ea a căzut în funcție de cât de multe obiecte reflectate vedem). Testele au fost făcute special într-o etapă cât mai simplă și mai difuză, unde nu există nici o reflecție. La o adâncime de peste 2, toate firele sunt inactive. Din nefericire, nu toate nucleele din raytracerul meu ar putea fi aruncate de stencil, deci cu o adâncime mai mare a reflexiilor, chiar și pentru o scenă difuză, FPS cade. Dinamica FPS este după cum urmează:

Pentru o scenă în oglindă. 30, 25, 23,7, 20, 19,4, 18,8 (Figura 2)

Pentru o scenă difuză. 40, 37, 34, 32, 30, 29,5

Pentru comparație, într-o scenă oglindă mai complexă:

Pentru oglinda 2: 32, 23, 18.6, 16.1, 14.4 (figura 3)

Stencil tampon pe cuda

Figura 2. O scenă simplă, mai mică de 100 de triunghiuri.

Stencil tampon pe cuda

Figura 3. O scenă puțin mai complicată,

23 de mii primitivi.







Articole similare

Trimiteți-le prietenilor: