OpenCL Mandelbrot Demo on Alteraを読む&試す

12396 ワード
FPGA OpenCL Altera OpenCL テキストリンク

RocketBoards.orgのOpenCL Mandelbrot Demo on Alteraを読んで試してみた，という話．

ソースコード

https://github.com/altcrauer/mandelbrot_demo で公開されている．
プログラムはARM上で実行されるものとFPGAで実行されるものから構成される．

ARM上で実行されるもの

  • main.cpp
    • colorTableInit 色テーブル[0～999]の初期化
    • main 初期化と実行
  • MandelbrotWindow.cpp マンデルブロー描画の処理全体の統括
    • mandelbrotWindowInitialize hardwareInitializeとsoftwareInitializeを呼び出す．また，SDL_CreateRGBSurfaceFromを呼び出してマンデルブロー画像保持用のバッファを2面用意(バッファへのポインタはtheFrames[0], theFrames[1]で保持される)
    • mandelbrotWindowRelease
    • mandelbrotWindowResetView
    • mandelbrotWindowUpdate
    • mandelbrotWindowMainLoop マンデルブローの計算と描画処理のループ
    • mandelbrotCalculateFrameを呼び出してtheFrames[]に画像を書かせる
    • 結果をディスプレイに出力する，ファイルに保存する，など
    • mandelbrotWindowRepaint
    • mandelbrotDumpFrame
  • Mandelbrot.cpp
    • mandelbrotSetColorTable
    • mandelbrotSwitchCalculationMethod マンデルブロー計算をHWで行うかSWで行うか切り替える．デフォルトはHW．実行中にSDLK_h(SDLのhキー押下イベント)で切り替え．
    • mandelbrotCalculateFrame マンデルブロー計算をHWまたはSWで実行
    • mandelbrotRelease
  • HardwareMandelbrot.cpp
    • hardwareInitialize FPGAにマンデルブロー計算カーネルをロード
    • hardwareSetColorTable
    • hardwareCalculateFrame FPGA上のカーネルでマンデルブロー計算
    • hardwareRelease
    • cleanup opencl.cppで利用されるクリーンアップ関数．実際はhardwareReleaseを呼んでいる
  • SoftwareMandelbrot.cpp
    • softwareInitialize HardwareInitializeと違って何もしてない
    • softwareSetColorTable
    • softwareCalculateFrame ソフトウェアでのマンデルブロー計算
    • softwareRelease
  • opencl.cpp OpenCLコアを利用するためのユーティリティ
    • setCwdToExeDir
    • findPlatform
    • getPlatformName
    • getDeviceName
    • getDevices
    • createProgramFromBinary
    • createProgramFromSource
    • loadBinaryFile
    • fileExists
    • getBoardBinaryFile
    • getCurrentTimestamp
    • getStartEndTime
    • getStartEndTime
    • waitMilliseconds
    • oclContextCallback
    • isAlteraPlatform
    • findAnyPlatform
  • Keyboard.cpp
    • keyboardPressEvent
  • Mouse.cpp
    • mousePressEvent
    • mouseReleaseEvent
  • StopWatch.cpp
    • startTime
    • getElapsedTime
  • options.cpp

FPGA上で実行されるもの

  • device/mandelbrot_kernel.cl
    • hw_mandelbrot_frame グローバルIDにしたがって並列計算

ソフトウェア処理とOpenCL処理の実装

マンデルブロー画像の描画ルーチンはCPUで実行するソフトウェアとOpenCLの両方で実装されている．それぞれの実装をみてみる

ソフトウェアバージョン

あるピクセルの値を決定するmandel_pixelが定義されていて

inline unsigned int mandel_pixel(MANDELBROT_SW_PRECISION x0, MANDELBROT_SW_PRECISION y0, unsigned int maxIterations)
{
// 略
  while (xSqr + ySqr < 4.0 && iterations < maxIterations)
  {
    // perform the current iteration                                                                                                 
    xSqr = x*x;
    ySqr = y*y;

    y = 2*x*y + y0;
    x = xSqr - ySqr + x0;

    // increment iteration count                                                                                                     
    iterations++;
  }

  // return the iteration count                                                                                                      
  return iterations;
}

softwareCalculateFrameの中で，全ピクセルを求めている

  // for each pixel in the y dimension window                                                                                        
  for (j = 0, cur_y = y; j < theHeight; j++, cur_y -= cur_step_size)
  {
    // for each pixel in the x dimension of the window                                                                               
    for (cur_x = x, k = 0; k < theWidth; k++, cur_x += cur_step_size)
    {
      // set the value of the pixel in the window                                                                                    
      pixel = mandel_pixel(cur_x, cur_y, theSoftColorTableSize);
      if (pixel == theSoftColorTableSize)
        *fb_ptr++ = 0x0;
      else
        *fb_ptr++ = theSoftColorTable[pixel];
    }
  }

OpenCLバージョン

OpenCLカーネルとしては，↓の処理がオフロードされる

        const size_t windowPosX = get_global_id(0);
        const size_t windowPosY = get_global_id(1);
        const MANDELBROT_HW_PRECISION stepPosX = x0 + (windowPosX * stepSize);
        const MANDELBROT_HW_PRECISION stepPosY = y0 - (windowPosY * stepSize);
// 略
        while ( xSqr + ySqr < 4.0 &&
                        iterations < maxIterations)
        {
                // Perform the current iteration                                                                                     
                xSqr = x*x;
                ySqr = y*y;

                y = 2*x*y + stepPosY;
                x = xSqr - ySqr + stepPosX;

                // Increment iteration count                                                                                         
                iterations++;
        }

        // Output black if we never finished, and a color from the look up table otherwise                                           
        framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations)? BLACK : colorLUT[iterations];

ソフトウェアからは次のように呼び出される(エラーチェック部分を省略している)．

int hardwareCalculateFrame(MANDELBROT_HW_PRECISION aStartX, MANDELBROT_HW_PRECISION aStartY, MANDELBROT_HW_PRECISION aScale, unsigned short int* aFrameBuffer)
{
  // Make sure width and height match up                                                                                             
  hardwareSetFrameBufferSize();

  unsigned rowOffset = 0;
  for(unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++])
  {
    // Create ND range size                                                                                                          
    size_t globalSize[2] = {thePixelDataWidth, rowsPerDevice[i]};

    // Set the arguments                                                                                                             
    unsigned argi = 0;
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(MANDELBROT_CL_PRECISION), (void*)&aStartX);
    const MANDELBROT_HW_PRECISION offsetedStartY = aStartY - rowOffset * aScale;
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(MANDELBROT_CL_PRECISION), (void*)&offsetedStartY);
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(MANDELBROT_CL_PRECISION), (void*)&aScale);
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theHardColorTableSize);
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&thePixelData[i]);
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&theHardColorTable);
    theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theWidth);
    // Launch kernel                                                                                                                 
    theStatus = clEnqueueNDRangeKernel(theQueues[i], theKernels[i], 2, NULL, globalSize, NULL, 0, NULL, NULL);
  }

  rowOffset = 0;
  for(unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++])
  {
    // Read the output                                                                                                               
    theStatus = clEnqueueReadBuffer(theQueues[i], thePixelData[i], CL_TRUE, 0, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned short int), &aFrameBuffer[rowOffset * theWidth], 0, NULL, NULL);
   }

簡単にためすには

OpenCL Mandelbrot Demo on Alteraの手順でコンパイルすればいい...けどカーネルのコンパイルは面倒という場合はイメージファイル(mandelbrot_demo.img)をダウンロードして試すことができる．
自分でコンパイルしたアプリケーションを追加したい場合には，イメージファイルをループバックでマウントして書けばいい．mandelbrot_demo.imgには複数のパーティションがあるので，たとえば，

fdisk -l -u mandelbrot_demo.img

で，Linuxシステムのあるパーティションの開始オフセットを確認して

sudo mount -o loop,offset=$((49152*512)) mandelbrot_demo.img /mnt

としてマウントすればいい．
OpenCLソースをLinuxな環境でコンパイルするには，たとえば

export LM_LICENSE_FILE=ライセンスファイル
export AOCL_BOARD_PACKAGE_ROOT=どこか/mandelblot_de0nanosoc/opencl_soc_bsp-de0_nano_with_display/c5soc
export ALTERAOCLSDKROOT=/opt/Altera/15.1/hld
source /opt/Altera/15.1/hld/init_opencl.sh
aoc -v device/mandelbrot_kernel.cl --board de0_nano_sharedonly_with_spi_tft

とか

ディスプレイがない場合の実行

ディスプレイがない場合にふつうに起動するとSDLのディスプレイチェックで終了する．なのでオプションが必要．
サンプルスクリプトの起動を参考に

source /de0_ocl_150_tft/init_opencl.sh
source /de0_ocl_150_tft/disable_auto_reprogram.sh
./mandelbrot --display=0 --test=1 --test-frames=10 --test-dump=10

とか．↓のように指定したらFPS計測の結果が出力された．ちょっとオプションのつけ方がわかってない

./mandelbrot --display=0 --test=true

実行している様子はこんな感じ．

標準出力に示されている生成速度をみると，18FPS-30FPSくらい，ってところのよう．

最初のオプションの場合には生成した画像がframe0.ppmからframe9.ppmに保存される．確認してみるとこんな感じ(convertでppm→png変換した結果)

OpenCLカーネルを変更してみる...rbfファイル変更編

プログラム実行時にハードウェアを変更することができなかった(要確認)ので，コンフィギュレーションファイルを変更する方法を試す．この環境ではSDカードの先頭パーティション(FAT32領域)のsoc_system.rbfがFPGAのコンフィギュレーションファイルに相当．
というわけで，aocを実行してOpenCLコードを合成すると生成されるtop.rbfで上書きしてみた．
結果はこんな感じ．

これは，マンデルブロー描画ルーチンで，黒にしていた部分を色テーブルの999に設定するよう変更したハードウェア．たしかに，オリジナルでは黒かったところが色テーブル999(黄色: main.cppで色テーブルを作っている)になっていることが確認できた．

        // Output black if we never finished, and a color from the look up table otherwise                                           
        //framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations)? BLACK : colorLUT[iterations];          
        framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations)? colorLUT[999] : colorLUT[iterations];