Stencil puffer a cuda-ban

Stencil puffer vagy stencil puffer általában grafika (OpenGL, DirectX), hogy elfedje néhány pixel a képen a hívás a pixel shader egyes területeken a kép. A szöveg kifejezetten hangsúlyozza, hogy a stencil tesztet a pixel shader előhívása előtt kell végrehajtani, és így azokon a helyeken, ahol a kép hiányzik, a pixel shader-eket egyáltalán nem hívják meg, és nem végeznek további munkát.

Fizikai értelemben a stencil a GPU-n tárolódik ugyanabban a pufferben, ahol a mélység tárolódik és különböző formátumú lehet. Például a legszélesebb körben használt D3DFMT_D24S8 formátum azt jelenti, hogy a hátsó pufferben 24 bitet osztanak le mélységre és 8 bitet a sablonra. Ebben a cikkben egyszerűsítést használunk és azt feltételezzük, hogy a stencil puffer minden pixelhez (vagy streamhez) csak egy bitet tárol. Ha a bit 1, akkor a pixel (stream) aktív. Ha 0, inaktív. Ezzel megtakaríthat egy kis memóriát és egyszerűsítené a prezentációt.

A Stencil Tesztet gyakran használják a gondolatok megalkotására:

1. ábra. A stencil puffer szükséges a reflektumok maszkolásához azokon a helyeken, ahol valójában nem léteznek (mint a jobb oldali képen).

1. Tisztítsa meg a stencil puffert nullákkal.

2. A rekordot a stencil pufferbe tesszük és rajzolunk egy síkot, amelyhez képest megfontoljuk a reflexiót. Mindig írunk egyet. Kiderül, hogy a tükörünk bináris képe a maszk pufferben van tárolva (azaz ahol tükör van tárolt egységekben, és ahol nincs tükör nullák).

3. Minden geometriát tükrözzenek a síkhoz képest egy speciális mátrix segítségével, és rajzolják le, beleértve a stencil tesztet is. Így, ha a képen tükör van, akkor egy visszaverődés fog megjelenni. És ahol nincs, semmi sem fog változni.

A szoftver megvalósítása a CUDA-n

Sajnos, a kudban, mint a többi "számítási" technológiában (DX11 CS, OpenCL), a stencil-teszt mechanizmus egyszerűen nem áll rendelkezésre. Ez a dolog ugyanakkor nagyon hasznos, különösen akkor, ha a számításokat több (gyakran meglehetősen kicsi) magból (kernelből) hosszú csővezetékként valósítjuk meg. Tegyük fel, hogy van N szál.

Például, ez a helyzet akkor következik be, ha a sugárkövetést az útvonalon hajtja végre. Néhány jelenet esetén körülbelül 5 visszaverődés mélysége esetén az áramlás kevesebb mint 10% -a aktív lesz az utolsó szinten.

Annak érdekében, hogy ne végezzen munkát az inaktív szálaknál, valószínűleg minden pufferben zászlót helyez el, és ellenőrizni fogja, hogy ez a jelzés 0, majd semmit sem csinál.

uintactiveFlag = a_flags [tid];

Ebben a cikkben javasoljuk, hogy egyetlen bitet tároljunk a stencil pufferben egy adatfolyamonként, és elkerüljük az adatok tömeges adatátvitelét a buszon (vagy legalábbis jelentősen csökkentsük őket, hatékonyan használva a gyorsítótárat).

Tehát egy stencil buffer méretét pontosan (N / 32) * sizeof (int) byte-ba állítjuk be. És összekapcsoljuk a textúrát.

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

Maga a textúra valamilyen fejlécben (.h fájlban) a következőképpen kerül bejelentésre:

struktúra stencil_tex;

Ezután ugyanabban a fájlban deklarálunk egy ilyen kiegészítő tömböt:

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

Ebben a tömbben vannak tárolt maszkok, amelyekkel logikusan fogunk eljárni annak érdekében, hogy gyorsan elérje a kívánt bitfolyamot. Ez azt jelenti, hogy pontosan kapjuk meg a számot, amelynek száma megegyezik a láncban lévő cérna számával. Így néz ki a stencil teszt:

ha (! (tex1Dfetch (stencil_tex, (tid) >> 5) g_stencilMask [(tid) 0x1f])) \

Azokat a rendszermagokat, amelyek csak a stencil puffereket olvassák el, a makrót a rendszermag elején kell alkalmazni a következőképpen:

__global__void my_kernel (...)

uinttid = blockDim.x * blockIdx.x + threadIdx.x;

A gyakorlatban (GTX560) ez a stencil teszt kb. 20-25% -kal gyorsabb, mint a formanyomtatvány egyszerű ellenőrzése:

uintactiveFlag = a_flags [tid];

Így csak a stencil pufferre kell írni. Először olvassa el az összes értéket a lopakodó pufferből az aktívWarp változóba; Ezután mindegyik szál logikusan veszi fel ebből a változóból a bitjét és tárolja az aktív változóban. A rendszermag végére egy adott láncértékre aktív összes változót visszaadjuk egy 32 bites Uint-be, és a nulla fonalhúzási láng visszaírja az eredményt a memóriába.

// (sz 0x1f) ugyanaz, mint (tid% 32)

__global__void my_kernel2 (..., uint * a_stencilBuffer)

uinttid = blockDim.x * blockIdx.x + threadIdx.x;

uint activeWarp = a_stencilBuffer [tid >> 5];

ha (aktívWarp == 0) // mindegyik szál láncban inaktív

// mindegyik szál külön bitet tárol a 32 szálcsoportból

uint aktív = aktívWarp g_stencilMask [tid0x1f];

Ez így működik. Amikor a lánc- befejezte munkáját (például ez a végén, míg a ciklus vagy volt döntve stencil teszt), hogy ellopja a következő része a munka maga, a növekvő globális számláló 32. Counter egyszerre azt jelzi, hogy mennyi művelet szabadon marad.

A G80-ban csak annyira tartós szálak kerülnek végrehajtásra, az atomos műveletek hiánya miatt. De csak egy hurkot készíthet, mint a "for (int i = 0; i<8;i++) doMyWork(i);” для того, чтобы увеличить количество работы, выполняемое одним warp-ом. На GT200 в некоторых случаях, использование persistent threads давало прирост производительности до 2 раз.

Stencil tesztelése

Valójában a raytracing igényeihez hasonlóan egy ilyen stencil puffer nagyon sikeresnek bizonyult. Ha eltemetsz magadat az ürességben, a GTX560-on kb. 4 milliárd hívás érkezhet másodpercre a másodpercre (vagyis 4 milliárd üres hívás másodpercenként). A nyomkövetés mélységében a teljesítmény gyakorlatilag nem esett (inkább azt, hogy mennyi visszavert objektumot látunk). A teszteket a lehető legegyszerűbb tükröződő színpadon végezték el, és teljesen diffúzak voltak, ahol egyáltalán nincsenek tükröződések. A> 2-es mélységben minden szál inaktív. Sajnos a raytracerben levő összes magot nem lehet eldobni a sablon, így a tükröződés mélységével még a diffúz jelenet miatt is az FPS leesik. Az FPS dinamikája a következő:

Tükörképhez. 30, 25, 23,7, 20, 19,4, 18,8 (2. ábra)

Egy diffúz jelenet miatt. 40, 37, 34, 32, 30, 29,5

Összehasonlításképpen egy összetettebb tükörképen:

A tükörkép 2: 32, 23, 18.6, 16.1, 14.4 (3. ábra)

2. ábra Egy egyszerű jelenet, kevesebb mint 100 háromszög.

3. ábra: Egy kicsit bonyolultabb jelenet,

23 ezer primitív.

Kapcsolódó cikkek

• A stencil használata az unity3d shaderlab-ban