Stencil buffer используется для маскировки отражений в тех местах где их на самом деле нет. Техника Stencil используется в OpenGL и DirectX. Перед применением алгоритма производится Stencil тест и в тех местах где нет изображения пиксельный шейдер не будет рендерится. Таким образом мы пресекаем лишнюю работу.
Stencil хранится в буфере с глубиной. Например в формате D3DFMT_D24S8 24 бита — биты глубины и 8 бит — Stencil. Для упрощения, далее мы будем считать, что Stencil хранится в последнем бите. Если этот бит = 1, то пиксель активен. Таким образом упрощенный алгоритм представляет следующие шаги:
- Обнуляем (заполняем нулями) Stencil buffer.
- Начинаем запись и рисуем в Stencil buffer плоскость, относительно которой будем считать отражение. Там где есть зеркало будут храниться единицы, а там где зеркала нет — нули.
- Отражаем всю геометрию относительно плоскости при помощи отдельной матрицы, и затем рисуем отражение, попутно запуская Stencil-тест.
Таким образом, там где на изображении находилось зеркало, будет выведено отражение. А там где его нет, ничего не изменится.
Программная реализация на CUDA
CUDA, к сожалению, механизм Stencil-теста отсутствует. Это очень полезный прием и как обойти это ограничение я расскажу в следующей статье, а сейчас рассмотрим детали реализации.
Итак, заводим стенсил буфер размером ровно на (N/32)*sizeof(int) байт. И привязываем к нему текстуру.
cudaMalloc((void**)&m_stencilBuffer, N*sizeof(int)/32);
cudaBindTexture(0, stencil_tex, m_stencilBuffer, N*sizeof(int)/32);
Сама текстура объявлена в каком-нибудь хедере (.h файл) следующим образом:
Texture<int, 1, cudaReadModeElementType> stencil_tex;
Далее, в том же файле объявим такой вспомогательный массив:
static __device__ int g_stencilMask[32] = {
0x00000001, 0x00000002, 0x00000004, 0x00000008, 0x00000010, 0x00000020, 0x00000040, 0x00000080,
0x00000100, 0x00000200, 0x00000400, 0x00000800, 0x00001000, 0x00002000, 0x00004000, 0x00008000,
0x00010000, 0x00020000, 0x00040000, 0x00080000, 0x00100000, 0x00200000, 0x00400000, 0x00800000,
0x01000000, 0x02000000, 0x04000000, 0x08000000, 0x10000000, 0x20000000, 0x40000000, 0x80000000
};
Для тех кернелов, которые только читают стенсил буфер, применять макрос следует в начале кернела следующим образом:
__global__ void my_kernel(…)
{
uint tid = blockDim.x * blockIdx.x + threadIdx.x;
STENCIL_TEST(tid);
// my code here
}
На практике (GTX560) такой стенсил тест примерно на 20-25% быстрее, чем простая проверка проверка вида:
uint activeFlag = a_flags[tid];
if(activeFlag==0)
return;
Однако, с учетом экономии памяти, профит определенно есть. Следует так же отметить, что на видеокартах с менее широкой шиной (например GTS450) ускорение может быть более существенным.
Итак, осталось реализовать лишь запись в стенсил-буфер. Сначала читаем значение для всего в warp-а из стелсил-буфера в переменную activeWarp; Затем каждый поток получает из этой переменной свой бит при помощи логического & и хранит его в переменной active. В конце кернела мы соберем из всех переменных active для данного warp-а значения обратно в один 32 разрядный uint, и нулевой поток warp-а запишет результат назад в память.
// (tid >> 5) same as (tid/32)
// (tid & 0x1f) same as (tid%32)
__global__ void my_kernel2(…,uint* a_stencilBuffer)
{
uint tid = blockDim.x * blockIdx.x + threadIdx.x;
uint activeWarp = a_stencilBuffer[tid >> 5];
if(activeWarp==0) // all threads in warp inactive
return;
// each threads will store it's particular bit from group of 32 threads
uint active = activeWarp & g_stencilMask[tid&0x1f];
if(!active)
goto WRITE_BACK_STENCIL_DATA;
// my code here
WRITE_BACK_STENCIL_DATA:
WriteStencilBit(tid, a_stencilBuffer, active);
}
Если поток неактивен, он сразу перейдет в конек кернела. Если по какой-либо причине вы внутри вашего кода решили, что этот поток должен быть неактивен, сделайте так:
if(want to kill thread)
{
active = 0;
goto WRITE_BACK_STENCIL_DATA;
}
В примере намеренно использована метка и оператор goto. Хоть это и является плохим стилем программирования, в данном случае это добавляет безопасности вашему коду. Дело в том, что вы обязаны гарантированно достичь кода функции WriteStencilBit. Если по какой-то причине внутри вашего кода вы решите сделать return, всё поломается (чуть позже обсудим почему). Вместо return надо ставить goto WRITE_BACK_STENCIL_DATA, чтобы перед выходом, все потоки из warp-a могли собрать данные, а нулевой поток (нулевой внутри warp-a) запишет их в стенсил-буфер. Собственно, функция WriteStencilBit выглядит следующим образом:
__device__ void WriteStencilBit(int tid, uint* a_stencilBuffer, uint value)
{
uint stencilMask = __ballot(value);
if((tid & 0x1f) == 0) // same as tid%32 == 0
a_stencilBuffer[tid >> 5] = stencilMask;
}
Функция __ballot() возвращает uint, где каждый i-ый бит равен 1 тогда и только тогда, когда то, что находится в ее аргументе не равно нулю. То есть она делает в точности то, что там нужно, сшивая обратно в uint флаги от разных потоков внутри warp-а.
Функция __ballot() принадлежит к так называемым “warp vote functions” и работает очень быстро. К сожалению, она доступна только для compute capability 2.0, то есть видеокарт с архитектурой Fermi. Важное замечание по её работе, следующий код будет неправильным:
__device__ void WriteWrongStencilBit(int tid, uint* a_stencilBuffer, uint value)
{
if((tid & 0x1f) == 0) // same as tid%32 == 0
a_stencilBuffer[tid >> 5] = __ballot(value);
}
Дело в том, что __ballot() будет всегда помещать 0 в те биты, чьи потоки замаскированы в данный момент. А все потоки с номером внутри варпа не равным нулю (1..31) будут замаскированны и не попадут внутрь оператора if, поэтому 1..31 биты результата функции __ballot() для такого кода всегда будут равны нулю. Отсюда правда следует интересный вывод. Если вы гарантированно пишете для видеокарт с архитектурой Fermi, то даже для кернелов которые пишут в стенсил буфе, вы можете убивать поток следующим образом:
if(want to kill thread)
return;
Таким образом, потоки, для которых вы сделали return будут замаскированы и __ballot() вернет в своем результате нули для соответствующих бит. Есть правда одна тонкость. По крайней мере для нулевого потока внутри warp-а вы так сделать не можете, иначе результат просто не запишется назад. Поэтому на самом деле можно делать только так
if(want to kill thread && (tid&0x1f!=0))
return;
Или пользуйтесь формой предложенной выше:
if(want to kill thread)
{
active = 0;
goto WRITE_BACK_STENCIL_DATA;
}
Особенности реализации для старой аппаратуры (G80-GT200)
Рассмотрим теперь, какие расширения должны быть сделаны, чтобы стенсил эффективно работал и на более старых GPU. На этих видеокартах не поддерживается функция __ballot(). Перепишем функцию WriteStencilBit в соответствие с теми возможностями, которые у нас имеются:
template<int CURR_BLOCK_SIZE>
__device__ inline void WriteStencilBit(int tid, uint* a_stencilBuffer, uint value)
{
#if COMPUTE_CAPABILITY >= COMPUTE_CAPABILITY_GF100
uint stencilMask = __ballot(value);
if((tid & 0x1f) == 0)
a_stencilBuffer[tid >> 5] = stencilMask;
#elif COMPUTE_CAPABILITY >= COMPUTE_CAPABILITY_GT200
if(__all(value==0))
{
if((tid & 0x1f) == 0)
a_stencilBuffer[tid >> 5] = 0;
}
else if(__all(value))
{
if((tid & 0x1f) == 0)
a_stencilBuffer[tid >> 5] = 0xffffffff;
}
else
{
__shared__ uint active_threads[CURR_BLOCK_SIZE/32];
uint* pAddr = active_threads + (threadIdx.x >> 5);
if((tid & 0x1f) == 0)
*pAddr = 0;
atomicOr(pAddr, value);
if((tid & 0x1f) == 0)
a_stencilBuffer[tid >> 5] = *pAddr;
}
#else
__shared__ uint active_threads[CURR_BLOCK_SIZE];
active_threads[threadIdx.x] = value;
if((threadIdx.x & 0x1) == 0)
active_threads[threadIdx.x] = value | active_threads[threadIdx.x+1];
if((threadIdx.x & 0x3) == 0)
active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+2];
if((threadIdx.x & 0x7) == 0)
active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+4];
if((threadIdx.x & 0xf) == 0)
active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+8];
if((threadIdx.x & 0x1f) == 0)
active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+16];
uint* perWarpArray = active_threads + ((threadIdx.x >> 5) << 5);
if((tid & 0x1f) == 0)
a_stencilBuffer[tid >> 5] = perWarpArray[0];
#endif
}
Таким образом мы можем делать атомики в шаред-память + доступны 2 функции голосования, __any и __all, так что мы их можем использовать. В остальных случаях остается только классическая редукция.
Тестируем Stencil
Для нужд рейтрейсинга, такой стенсил буфер подошел довольно удачно. На GTX560 моего старенького ноута получается около 4 миллиардов вызов кернелов в секунду (то есть 4 миллиарда пустых вызовов в секунду) — неплохо, правда?! При увеличении глубины трассировки производительность немного падала в соответствии с тем, насколько реально много отраженных объектов мы видим. Тесты специально производились на как можно более простой отражающей сцене:
Динамика FPS следующая: 30, 25, 23.7, 20, 19.4, 18.8
Автор: trollsupport
