GPUでレイトレーシングを並列処理

未分類

以前やったGPUで並列処理の続きとして今回はGPUでレイトレーシングを行いたいと思います。
レイトレーシングはリアルな画像を作り出す反面、計算に大変時間がかかることはよく知られています。
その処理をいかに高速化できるか？というのが今回のコラムの目的です。

基本的なピクセル並列

GeForce8x 系のアーキテクチャが基本的にマルチスレッドで動作することは以前に触れました。
まず考え付く方法としてピクセル単位の並列化が考えられると思います。
具体的に言うとX軸の画素をブロック内のスレッド１つ１つに割り当て、Y軸の画素をそれぞれのブロック単位で処理します。
この方法はオブジェクトが少なく、またピクセル数が並列動作の最大パフォーマンスが出せるスレッド数 (192スレッド以上) に近づくほど効率がよくなります。
ソースコードは以下のような感じとなります 

////////////////////////////////////////////////////////////////////////
// GPU処理エントリポイント
#define SCREEN_WIDTH  512
#define SCREEN_HEIGHT 512// 512*512の画像を出力するとしました。
__global__
void GPUMain(
Display* fOut,       // 画面ピクセルの位置と色を格納した構造体
Object* fIn,         // オブジェクト情報を格納した構造体
Environment *fIn_env // 環境情報を格納した構造体
)
{
const unsigned int pixelx = threadIdx.x; // ピクセルXにスレッド1個を対応
const unsigned int pixely = blockIdx.x;  // ピクセルYにブロック１個を対応
__shared__ Object object[MAXOBJ];        // オブジェクトのデータ量により
// シェアードメモリに蓄えれる物体数が変化します。
if(0<=pixelx && pixelx < MAXOBJ)
{
object[pixelx] = fIn[pixelx];        // sharedメモリにストア
}
__syncthreads();                         // ストア終了まで待つ
// レイトレーシング機能をまとめた関数とします。
// ピクセル位置、オブジェクト情報、環境情報から視線追跡を行い出力ピクセルを算出します。
// 今回この内部は並列化していませんので割愛させていただきます。
RayTrace(pixelx, pixely, object, fIn_env, &fOut[SCREEN_WIDTH * pixely + pixelx]);
}

しかし一般的にこのような限られた状況で行うレイトレーシングは稀です。
以前も触れましたがsharedメモリの容量は小さいので、このままでは完全にストアできるオブジェクトはわずか 数十個〜数百個です。 オブジェクト情報をシェアードメモリを使用せずグローバルメモリから読み込めば難しく考える必要もありませんが、高速化といった観点からははずれてしまいます。

ブロック単位のピクセル分割

よって解決策として画面分割とピクセル単位の並列化を併用します。 画面分割によってブロック内に含まれるオブジェクトに対してのみ計算処理を行えるようになります。 また、ピクセル１個に対しスレッド１個を割り当て、2*2, 4*4, 8*8,... ピクセルのブロックといった感じで マルチコアに割り当てます。 この方法を実装したプログラムは次のようになります。 

#define SCREEN_WIDTH  512
#define SCREEN_HEIGHT 512 // 512*512の画像を出力するとしました。
#define PIXLE_BLOCK_SIZE 8 // 8*8ピクセルのブロック単位で処理をすると設定
__global__
void GPUMain(
Display* fOut,        // 画面ピクセルの位置と色を格納した構造体
Object* fIn,          // オブジェクト情報を格納した構造体
Environment *fIn_env  // 環境情報を格納した構造体
)
{
const unsigned int pixelx = threadIdx.x; // ピクセルXにスレッド1個を対応
const unsigned int pixely = threadIdx.y; // ピクセルYにスレッド1個を対応
__shared__ Object object[MAXOBJ];        // シェアードメモリはブロック内で有効
if(0<=pixelx && pixelx < MAXOBJ)
{
object[pixelx] = fIn[pixelx];        // sharedメモリにストア
}
__syncthreads();
RayTrace(pixelx + (blockIdx.x % 64) * PIXEL_BLOCK_SIZE,  // 走査するピクセルのX座標を設定
pixely +  (blockIdx.x / 64) * PIXEL_BLOCK_SIZE,      // 走査するピクセルのY座標を設定
object, fIn_env, &fOut[SCREEN_WIDTH * pixely + pixelx]);
}

この方法を使用すれば以前のプログラムより多くのオブジェクトを高速に処理することが可能になります。
しかしこの方法を実装しても画面いっぱいにオブジェクトが含まれる場合やブロック内の数個のピクセル処理に多くの時間を必要とした場合、ブロック全体のパフォーマンスが落ちます。
また、CUDAには10秒ルールという制約があり、１つのジョブに10秒以上時間がかかると強制的に処理が中断されます。
そういうタスクを行う場合はCPUに一度処理を戻さなければならないのですが、それを繰り返すことによりメモリ内容を転送するディレイが発生します。
これらの問題解決策として空間分割を行い、レイ、ボクセル単位で並列化する方法があるのですが、それはまた次の機会にということで・・・。

後これは愚痴なんですが、計算した内容を直接画面にキックする関数くれよ・・・と。
CPUに戻す手間なんとかなりませんか・・・。

担当: 松浦