Stencil buffer або буфер трафарету зазвичай використовується в графіку (OpenGL, DirectX) для того, щоб замаскувати деякі пікселі на ізображеніівизови пиксельного шейдера для деяких областей зображення. У тексті спеціально підкреслено, що stencil test проводиться ще до виклику пиксельного шейдера і, таким чином, в тих місцях де зображення відсутній, піксельний шейдер взагалі не буде викликатися і ніяка зайва робота виконуватися не буде.
Фізично, СТЕНС зберігається на GPU в тому ж буфері, де зберігається глибина і буває різного формату. Наприклад, найбільш широко використовуваний формат D3DFMT_D24S8 означає, що 24 біта відводяться в бек-буфері на глибину і 8 біт на СТЕНС. У даній статті, ми буде використовувати спрощення і вважати, що СТЕНС-буфер зберігає на кожен піксель (або на потік) всього один біт. Якщо біт дорівнює 1, то піксель (потік) активний. Якщо 0, то неактивний. Це дозволить заощадити трохи пам'яті і спростить виклад.
Stencil Test часто використовують для побудови відображень таким методом:
Малюнок 1. Stencil buffer потрібен для маскування відображень в тих місцях де їх насправді немає (як на рис. Праворуч).
1. Очищаємо СТЕНС-буфер нулями.
2. Включаємо запис в СТЕНС буфер і малюємо в нього площину, щодо якої будемо вважати відображення. Записуємо завжди одиничку. Виходить, що в буфері маски зберігається бінарне зображення нашого дзеркала (тобто там де є дзеркало будуть зберігатися одиниці, а там де дзеркала немає - нулі).
3. Відображаємо всю геометрію щодо площині за допомогою спеціальної матриці, і малюємо її, включаючи СТЕНС-тест. Таким чином, там де на зображенні знаходилося дзеркало, буде виведено відображення. А там де його немає, нічого не зміниться.
Програмна реалізація на CUDA
На жаль, в куде, як і у всіх інших 'compute' технологіях (DX11 CS, OpenCL) механізм СТЕНС-тесту просто відсутня. У той же час, це річ дуже корисна, особливо якщо ваші обчислення реалізовані у вигляді довгого конвеєра з декількох (часто досить невеликих) ядер (kernels). Припустимо у вас є N потоків.
Наприклад, така ситуація зустрічається при реалізації на куде трасування променів. При глибині перевідбиттів близько 5, на деяких сценах, менше 10% потоків буде активні на останньому рівні.
Для того, щоб не виконувати роботу для неактивних потоків ви, швидше за все, заведете прапор в якомусь буфері і будете перевіряти, якщо це прапор дорівнює 0, то нічого не робити.
uintactiveFlag = a_flags [tid];
У даній статті пропонується зберігати в СТЕНС-буфері на 1 потік всього один біт і уникнути масових трансфертів даних по шині (або принаймні значно їх скоротити, ефективно використовуючи кеш).
Отже, заводимо СТЕНС буфер розміром рівно на (N / 32) * sizeof (int) байт. І прив'язуємо до нього текстуру.
cudaBindTexture (0, stencil_tex, m_stencilBuffer, N * sizeof (int) / 32);
Сама текстура оголошена в якомусь хедері (.h файл) наступним чином:
Texture
Далі, в тому ж файлі оголосимо такий допоміжний масив:
0x00000001, 0x00000002, 0x00000004, 0x00000008, 0x00000010, 0x00000020, 0x00000040, 0x00000080,
У цьому масиві зберігаються маски, з якими ми будемо робити логічний для того, щоб швидко отримати потрібний потоку біт. Тобто отримати рівно той біт, номер якого дорівнює номеру потоку всередині warp-а. Ось як буде виглядати stencil test:
if (! (tex1Dfetch (stencil_tex, (tid) >> 5) g_stencilMask [(tid) 0x1f])) \
Для тих Кернел, які тільки читають СТЕНС буфер, застосовувати макрос слід на початку Кернел наступним чином:
__global__void my_kernel (...)
uinttid = blockDim.x * blockIdx.x + threadIdx.x;
На практиці (GTX560) такий СТЕНС тест приблизно на 20-25% швидше, ніж проста перевірка перевірка виду:
uintactiveFlag = a_flags [tid];
Отже, залишилося реалізувати лише запис в СТЕНС-буфер. Спочатку читаємо значення для всього в warp-а з стелс-буфера в змінну activeWarp; Потім кожен потік отримує з цієї змінної свій біт за допомогою логічного і зберігає його у змінній active. В кінці Кернел ми зберемо з усіх змінних active для даного warp-а значення назад в один 32 розрядний uint, і нульовий потік warp-а запише результат назад в пам'ять.
// (tid 0x1f) same as (tid% 32)
__global__void my_kernel2 (..., uint * a_stencilBuffer)
uinttid = blockDim.x * blockIdx.x + threadIdx.x;
uint activeWarp = a_stencilBuffer [tid >> 5];
if (activeWarp == 0) // all threads in warp inactive
// each threads will store it's particular bit from group of 32 threads
uint active = activeWarp g_stencilMask [tid0x1f];
Працює це таким чином. Коли warp завершив свою роботу (наприклад він знаходиться в кінці циклу while або був відкинутий СТЕНС-тестом), він краде наступну порцію роботи для себе, збільшуючи глобальний лічильник відразу на 32. Лічильник вказує на те, скільки ще вільної роботи залишилося.
На G80 саме так persistent threads реалізувати не вийде, внаслідок відсутності атомарних операцій. Але можна просто зробити цикл виду "for (int i = 0; i<8;i++) doMyWork(i);” для того, чтобы увеличить количество работы, выполняемое одним warp-ом. На GT200 в некоторых случаях, использование persistent threads давало прирост производительности до 2 раз.
тестуємо Stencil
Власне для потреб рейтрейсінга, такий СТЕНС буфер підійшов досить вдало. Якщо уткнуться в порожнечу, на GTX560 можливо виходить близько 4 мільярдів виклик Кернел в секунду (тобто 4 мільярди порожніх викликів в секунду). При збільшенні глибини трасування продуктивність практично не падала (вірніше падала відповідно до того, наскільки реально багато відображених об'єктів ми бачимо). Тести спеціально проводилися на якомога простішою відбиває сцені і на повністю дифузійної, де відображень немає взагалі. На глибині трасування> = 2 всі потоки неактивні. На жаль не всі Кернел в моєму рейтрейсере можна було відкинути стенсілом, тому з ростом глибини відображень навіть для дифузійної сцени 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 тис. Примітивів.