Шаблон буфер или шаблон буфер обикновено се използва в графика (OpenGL, DirectX), с цел да се прикрие някои от пикселите в пиксел шейдър izobrazheniivyzovy за някои части на изображението. Текстът изрично посочи, че тест шаблон се извършва преди повикване пиксел шейдър и поради това в тези места, където няма изображение, пиксел шейдър като цяло няма да се обади и нищо допълнително, за да се извърши.

Физически, шаблон се съхранява на графичния процесор в същия буфер, който съхранява дълбочината и се предлагат в различни формати. Например, най-широко използваният формат D3DFMT_D24S8 означава, че 24 бита са определени в задния буфер и дълбочина от 8 бита на шаблон. В тази статия, ние ще използваме опростяване и да приемем, че буферните магазини шаблони за всеки пиксел (или поток), само един бит. Ако битът е 1, а след това на пиксела (поток) е активен. Ако 0, не е активна. Това ще спести малко памет и опростяване на презентацията.

Шаблон Тест често се използва за изграждане на отраженията по този начин:

Фигура 1. Шаблон буфер е необходимо да се прикрие отраженията в тези места, където те наистина не са (както е на фиг. Вдясно).

1. Изчистване на шаблон буфер с нули.

2. Включване на запис в шаблон буфер и привлече в своя самолет, за която не се поема размисъл. Напиши винаги yedinichku. Оказва се, че съхраняваните бинарен имидж на нашия огледала в буфера маска (тоест, когато има огледало единица трябва да се съхраняват, а когато огледалото е не - нула).

3. отразява геометрия спрямо равнината със специална матрица, за да я направи, включително тест шаблон. По този начин, от която е огледало, ще бъдат показани на размисъл. А там, където тя не съществува, нищо няма да се промени.

изпълнение на софтуер за CUDA

За съжаление, Куде, както и в "изчислителни технологии" всички останали (DX11 CS, OpenCL) тест механизъм шаблон, просто не е на разположение. В същото време, това нещо е много полезна, особено ако вашите изчисления са реализирани като тръбопровод с дължина от няколко (често доста малки) ядра (ядки). Да речем, че имаме N нишки.

Например, такава ситуация се появили се при изпълнението на проследяване на Куде лъч. На дълбочина около 5 отражения на някои етапи, е по-малко от 10% потоци да бъдат активни в последния ниво.

За да не вършат работа за неактивни потоци вероятно ще zavedet флаг във всеки буфер и ще провери дали на знамето е 0, а след това не се прави нищо.

uintactiveFlag = a_flags [три пъти дневно];

В тази статия ние предлагаме да се съхранява в буфера шаблон 1 само един поток малко и да се избегне масово трансфер на данни над автобуса (или поне да ги намали значително, ефективно използване на кеша).

Така че, ние започваме шаблон размер на буфера точно (N / 32) * sizeof (инт) байта. И се свързва с него текстура.

cudaBindTexture (0, stencil_tex, m_stencilBuffer, N * sizeof (междинно съединение) / 32);

Самият текстурата е обявен в някои cheder (.h файл), както следва:

текстура stencil_tex;

Освен това, в същия файл декларират под-масив:

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

Този масив съхранява на маската, което ние ще направим логичното за да получите бързо правилния поток на бита. Това е да се получи точно един бит, броят на които е равен на броя на потока вътре в основата-а. Ето как ще изглежда тест шаблон:

ако (! (tex1Dfetch (stencil_tex, (три пъти дневно) >> 5) g_stencilMask [(три пъти дневно) 0x1f])) \

Ядрото за тези, които само четат шаблон буфер, използвайте макро, за да бъде в началото на ядрото, както следва:

__global__void my_kernel (...)

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

На практика (GTX560) тест шаблон около 20-25% по-бързо от обикновен тип проверка проверка:

uintactiveFlag = a_flags [три пъти дневно];

По този начин, остава да се реализира само влизане в шаблон буфер. Първо, прочете стойността само за основа, и от стелт буфер променлива activeWarp; След това, всяка нишка излиза от тази променлива с помощта на малко логика и го съхранява в променлива активна. В края на ядрото, която събираме от всички активни променливи за основата-и ценностите назад към 32-битова uint една и нула деформират-поток и записва резултата обратно в паметта.

// (три пъти дневно 0x1f) същото като (три пъти дневно% 32)

__global__void my_kernel2 (..., uint * a_stencilBuffer)

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

uint activeWarp = a_stencilBuffer [три пъти дневно >> 5];

ако (activeWarp == 0) // всички нишки в основата неактивен

// всеки конци ще се съхранява това е специално малко от група от 32 нишки

uint активен = activeWarp g_stencilMask [tid0x1f];

Тя работи по следния начин. Когато основата е приключил работата си (например, че е в края, докато цикълът или е излегнала тест шаблон), той краде следващата част от работата за себе си, за увеличаване на общото брояч 32. Counter наведнъж това показва колко много работа е останало свободно.

В G80 е толкова устойчиви нишки, за да осъзнаят, това не е така, поради липса на атомни операции. Но вие може просто да направи една линия на формата "за (INT I = 0; аз<8;i++) doMyWork(i);” для того, чтобы увеличить количество работы, выполняемое одним warp-ом. На GT200 в некоторых случаях, использование persistent threads давало прирост производительности до 2 раз.

Тестване за шаблони

Всъщност за нуждите на рейтрейсинг, шаблон буфер мина доста добре. Ако погребе празнотата, на GTX560 тя може да получи около 4 милиарда Kernel повикване в секунда (т.е. 4 милиарда празни разговори в секунда). Чрез увеличаване на дълбочината и следа от изпълнението почти не попада (или по-скоро попада в съответствие с колко всъщност отразява на обектите, които виждаме). Тестовете бяха проведени на специално възможно най-прости отразява етап и напълно дифундират, където няма отражение на всички. На дълбочина проследяване> = 2 всички нишки са неактивни. За съжаление не всички от ядрото в моя raytracer може да наклони шаблон, така че с увеличаване на дълбочината отражения дори дифузно сцена FPS пада. динамика FPS е както следва:

За огледало сцена. 30, 25, 23.7, 20, 19.4, 18.8 (фиг. 2)

За да се разпространи на сцената. 40, 37, 34, 32, 30, 29,5

За сравнение, по-сложен огледално на сцената:

За огледало етап 2: 32, 23, 18.6, 16.1, 14.4 (Фигура 3)

Фигура 2. проста сцена, по-малко от 100 триъгълници.

Фигура 3. малко по-сложна сцена,

23 хил. Примитиви.

Свързани статии

• Ние работим с шаблона в Unity3D shaderlab