OpenCL Mandelbrot Demo on Alteraを読む&試す

RocketBoards.orgの[OpenCL Mandelbrot Demo on Altera](https://rocketboards.org/foswiki/view/Projects/OpenCLMandelbrotDemoOnAtlasSoC)を読んで試してみた，という話．

## ソースコード
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](https://rocketboards.org/foswiki/view/Projects/OpenCLMandelbrotDemoOnAtlasSoC)の手順でコンパイルすればいい...けどカーネルのコンパイルは面倒という場合はイメージファイル(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](https://qiita-image-store.s3.amazonaws.com/0/25342/42a0740c-d038-041d-fbda-37ccb8c1145d.png)

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


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

![frame0.png](https://qiita-image-store.s3.amazonaws.com/0/25342/f2e8e72a-e842-7f9b-45e4-2e8ef35d9935.png)

![frame4.png](https://qiita-image-store.s3.amazonaws.com/0/25342/87167b8d-432a-3123-38d2-310d0473e315.png)

![frame9.png](https://qiita-image-store.s3.amazonaws.com/0/25342/1d7be804-8b87-11eb-9859-beb6a400e9b2.png)

## OpenCLカーネルを変更してみる...rbfファイル変更編
プログラム実行時にハードウェアを変更することができなかった(**要確認**)ので，コンフィギュレーションファイルを変更する方法を試す．この環境ではSDカードの先頭パーティション(FAT32領域)のsoc_system.rbfがFPGAのコンフィギュレーションファイルに相当．
というわけで，aocを実行してOpenCLコードを合成すると生成されるtop.rbfで上書きしてみた．
結果はこんな感じ．
![frame0_mod.png](https://qiita-image-store.s3.amazonaws.com/0/25342/524645bb-93d1-6e0e-6dcf-dbe651dde437.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];
```