ステンシルバッファは、実際にはそうでない場所の反射をマスクするために使用されます。 ステンシル技術は、OpenGLとDirectXで使用されます。 アルゴリズムを適用する前に、ステンシルテストが実行され、画像がない場所では、ピクセルシェーダーはレンダリングされません。 したがって、不必要な作業を抑制します。
ステンシルは深さのあるバッファに保存されます。 たとえば、D3DFMT_D24S8形式では、24ビットは深度ビットで、8ビットはステンシルです。 簡単にするために、ステンシルは最後のビットに保存されるとさらに仮定します。 このビットが1の場合、ピクセルはアクティブです。 したがって、簡略化されたアルゴリズムは次の手順を示します。
- ゼロ（ゼロで埋める）ステンシルバッファー。
- 記録を開始し、ステンシルバッファーに平面を描画します。これに対して、反射を考慮します。 ミラーがある場合はユニットが格納され、ミラーがない場合はゼロが格納されます。
- 別のマトリックスを使用して、すべてのジオメトリを平面に対して反射し、反射を描画しながら、同時にステンシルテストを実行します。
したがって、鏡が画像のどこにあったか、反射が表示されます。 そして何もなければ、何も変わりません。
CUDAソフトウェアの実装
残念ながら、CUDAにはステンシルテストメカニズムがありません。 これは非常に便利なトリックであり、次の記事でこの制限を回避する方法を説明しますが、実装の詳細を確認します。
そのため、ステンシルバッファは、正確に（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など）では、加速がより重要になる場合があることに注意してください。
そのため、ステンシルバッファへの書き込みのみを実装することに変わりはありません。 まず、ワープ内のすべての値をステルスバッファからactiveWarp変数まで読み取ります。 次に、各スレッドは論理＆を使用してこの変数からビットを受け取り、アクティブ変数に格納します。 カーネルの最後に、特定のワープのすべてのアクティブな変数から値を1つの32ビットuintに戻し、ゼロワープスレッドが結果をメモリに書き込みます。
// (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関数コードに到達することが保証されています。 何らかの理由でコード内でリターンを行うことにした場合、すべてが壊れます（少し後で説明します）。 代わりに、リターンを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です。 つまり、そこで必要なことを正確に行い、ワープ内のさまざまなスレッドからフラグをステッチに戻します。
__ballot（）関数は、いわゆる「ワープ投票関数」に属し、非常に高速に動作します。 残念ながら、コンピューティング機能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ステートメント内に入らないため、そのようなコードの__ballot（）関数結果の1..31ビットは常にゼロになります。 ここから、真実は興味深い結論に従います。 Fermiアーキテクチャのビデオカード用に書き込みが保証されている場合、ステンシルバッファに書き込むカーネルでも、次のようにストリームを強制終了できます。
if(want to kill thread) return;
したがって、返されたストリームはマスクされ、__ ballot（）は結果の対応するビットにゼロを返します。 本当に微妙な点が1つあります。 少なくとも、ワープ内のゼロフローの場合、これを行うことはできません。そうしないと、結果が単純に書き戻されません。 したがって、実際には、これしかできません
if(want to kill thread && (tid&0x1f!=0)) return;
または、上記のフォームを使用します。
if(want to kill thread) { active = 0; goto WRITE_BACK_STENCIL_DATA; }
古いハードウェアの実装機能（G80-GT200）
Stencilが古い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を使用できるので、それらを使用できます。 その他の場合、古典的な削減のみが残ります。
ステンシルのテスト
レイトレーシングのニーズに合わせて、このようなステンシルバッファーが非常にうまく機能しました。 私の古いラップトップのGTX560では、1秒あたり約40億のカーネルコール（つまり、1秒あたり40億の空のコール）を受け取ります-悪くないでしょう？！ トレースの深さが増加すると、反映される実際のオブジェクトの数に応じてパフォーマンスがわずかに低下しました。 テストは、可能な限り単純な反射シーンで特別に実行されました。
FPSのダイナミクスは次のとおりです。30、25、23.7、20、19.4、18.8