LoginSignup
1
2

More than 5 years have passed since last update.

@miyox(たけふみ みよし)inインテル® FPGA ユーザーコミュニティ

OpenCL Mandelbrot Demo on Alteraを読む&試す

Last updated at Posted at 2016-04-13

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

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

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

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

frame0.png

frame4.png

frame9.png

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

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

これは，マンデルブロー描画ルーチンで，黒にしていた部分を色テーブルの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];
1
2
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
1
2