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];