Пример реализации Stencil буфера с помощью CUDA

в 15:38, , рубрики: CUDA, stencil, Анимация и 3D графика, метки: ,

Пример реализации Stencil буфера с помощью CUDA
Stencil buffer используется для маскировки отражений в тех местах где их на самом деле нет. Техника Stencil используется в OpenGL и DirectX. Перед применением алгоритма производится Stencil тест и в тех местах где нет изображения пиксельный шейдер не будет рендерится. Таким образом мы пресекаем лишнюю работу.

Stencil хранится в буфере с глубиной. Например в формате D3DFMT_D24S8 24 бита — биты глубины и 8 бит — Stencil. Для упрощения, далее мы будем считать, что Stencil хранится в последнем бите. Если этот бит = 1, то пиксель активен. Таким образом упрощенный алгоритм представляет следующие шаги:

  1. Обнуляем (заполняем нулями) Stencil buffer.
  2. Начинаем запись и рисуем в Stencil buffer плоскость, относительно которой будем считать отражение. Там где есть зеркало будут храниться единицы, а там где зеркала нет — нули.
  3. Отражаем всю геометрию относительно плоскости при помощи отдельной матрицы, и затем рисуем отражение, попутно запуская 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 миллиарда пустых вызовов в секунду) — неплохо, правда?! При увеличении глубины трассировки производительность немного падала в соответствии с тем, насколько реально много отраженных объектов мы видим. Тесты специально производились на как можно более простой отражающей сцене:
Пример реализации Stencil буфера с помощью CUDA
Динамика FPS следующая: 30, 25, 23.7, 20, 19.4, 18.8

Автор: trollsupport

* - обязательные к заполнению поля


https://ajax.googleapis.com/ajax/libs/jquery/3.4.1/jquery.min.js