////////////////////////////////////////////////////////////////////////
// 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]);
}

#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]);
}