雑記帳

ちょっとした文章とかメモ書きとか。

HaskellでOpenCLを使ってみる その2

 Control.Parallel.OpenCLを使う

今回はその2。

Control.Parallel.OpenCLには通常のOpenCLと同じく
OpenCLのカーネル関数をビルドすることや、
ホストコードの実行などをHaskellから呼び出せるようになっています。

Control.Parallel.OpenCLのGithub:https://github.com/zhensydow/opencl にexampleがあるのでそれを見ていけば基本的には問題はないです。

しかしこのライブラリ、単にラップしているだけなので便利なユーティリティ関数は用意されていません。

という事で、少し手を加えた実行例を示します。

Haskellでは型情報が大いに情報を持っているのでそれを参考にしながら書いていきます。

OpenCLではカーネル関数を読み込んでデバイスドライバ側でビルドを走らせます。
そのカーネル関数を読み込む関数を見てみましょう。

clCreateProgramWithSource関数ですね。

clCreateProgramWithSource :: CLContext-> String -> IO CLProgram

見たところ第二引数にStringが渡されているので、
これを外部ファイルから読み込む関数を使って渡してやればOpenCLの本領(?)が発揮できるでしょう。

ファイル読み込み

ファイル読み込みはエラーも伴うのでとりあえずcatchでベタ書きしてみます。

onError関数でエラーハンドリングをしています。

import Prelude hiding ( catch )
import System.IO 
import System.Exit
import Control.Exception
--error handling
onError :: String -> IOError -> IO String
onError filename error = do
  hPutStrLn stderr $ "File not found: " ++ filename
  exitWith(ExitFailure 1)

--read from file
programSourceFromFile :: IO String
programSourceFromFile = do let clprogramName = "testdouble.cl"
                           catch (readFile clprogramName)
                                 (onError clprogramName)

と書いてみました。はい、これでIOの作用をもつStringを返す関数が書けました。

ホスト側のコード

これを用いて書いたOpenCLのHaskellによるホストコードは以下:
–testcl.hs–

import Prelude hiding ( catch )
import Control.Parallel.OpenCL
import Foreign( castPtr, nullPtr, sizeOf )
import Foreign.C.Types( CFloat,CDouble )
import Foreign.Marshal.Array( newArray, peekArray )
import System.IO
import System.Exit
import Control.Exception

--error handling
onError :: String -> IOError -> IO String
onError filename error = do
  hPutStrLn stderr $ "File not found: " ++ filename
  exitWith(ExitFailure 1)

--read from file
programSourceFromFile :: IO String
programSourceFromFile = do let clprogramName = "testdouble.cl"
                           catch (readFile clprogramName)
                                 (onError clprogramName)

main :: IO ()
main = do
  -- Initialize OpenCL
  (platform:_) <- clGetPlatformIDs
  (dev:_) <- clGetDeviceIDs platform CL_DEVICE_TYPE_GPU
  context <- clCreateContext [] [dev] print
  q <- clCreateCommandQueue context dev []

  -- Compile and get binaries
  clsource <- programSourceFromFile
  program <- clCreateProgramWithSource context clsource
  clBuildProgram program [dev] ""
  bins <- clGetProgramBinaries program

  -- Create kernel from binaries
  (program2,_) <- clCreateProgramWithBinary context [dev] bins
  let clkernelBuildOption = ""
      clkernelName = "arraydev"
  clBuildProgram program2 [dev] clkernelBuildOption
  kernel <- clCreateKernel program2 clkernelName
  -- Initialize parameters
  let original = [0 .. 20] :: [CDouble] -- or [CFloat]
      elemSize = sizeOf (0 :: CDouble) -- or CFloat
      vecSize = elemSize * length original
  putStrLn $ "Original array = " ++ show original
  input <- newArray original

  mem_in <- clCreateBuffer context [CL_MEM_READ_ONLY, CL_MEM_COPY_HOST_PTR] (vecSize, castPtr input)
  mem_out <- clCreateBuffer context [CL_MEM_WRITE_ONLY] (vecSize, nullPtr)

  clSetKernelArgSto kernel 0 mem_in
  clSetKernelArgSto kernel 1 mem_out

  -- Execute Kernel
  eventExec <- clEnqueueNDRangeKernel q kernel [length original] [1] []

  -- Get Result
  eventRead <- clEnqueueReadBuffer q mem_out True 0 vecSize (castPtr input) [eventExec]
  result <- peekArray (length original) input
  putStrLn $ "Result array   = " ++ show result
  let resultcpu = map (*2)[0.0,1.0 .. 20.0]
  putStrLn $ unlines ["compare"
                     ,"cpu : " ++ show resultcpu
                     ,"gpu : " ++ show result
                     ,"same? : " ++ show (resultcpu == result)]

  return ()

ホストコードはこのようになりました。

以下の部分でOpenCLのカーネル関数のソースを読み込んでいます。

-- Compile and get binaries
clsource <- programSourceFromFile
program <- clCreateProgramWithSource context clsource

この関数により、OpenCLのソースを書き換えただけではホスト側のコンパイルは不要になりました。

ここで、HaskellからFFIを用いてCの関数を呼び出しているので、Foreignモジュールをインポートしています。–Get Result から後は得られた値を表示したり、比較したりしています。

[CDouble]のListを2倍にした値を求めるのにOpenCLのカーネル関数に値を渡したり、値を戻してきたりするのにかなりステップが要るのがめんどくさい。

CPU側でその値を得るには

let resultcpu = map (*2)[0.0,1.0 .. 20.0]

だけで良かった。かんたん!

Control.Parallel.OpenCLのGithubの例では文字列にして渡していますが、このようにしたほうがOpenCL側のコードを書きやすくなります。

OpenCLのカーネル関数側のコード

次にOpenCL側のカーネル関数のコードです。

–testdouble.cl–

#ifdef cl_khr_fp64
  #pragma OPENCL EXTENSION cl_khr_fp64 : enable
#elif defined(cl_amd_fp64)
  #pragma OPENCL EXTENSION cl_amd_fp64 : enable
#else
  #error Double precision floating point not supported by OpenCL implementation.
#endif
__kernel void arraydev(__global const double *in, __global double *out ){
  int id = get_global_id(0);
  out[id] = 2.0*in[id];
}

何やら悲しみに満ち溢れたpragmaがありますね。これはOpenCL1.1のデフォルトのプロファイルではdouble型の変数がカーネル関数で許可されていないのでそれを有効にするpragmaです。
これをさっきのtestcl.hsと同じフォルダに入れておきます。

もしもエラーになってしまうなら使っているOpenCLデバイスがdouble型に対応してないので

  • CDouble ->CFloat
  • double -> float

の書き換えと#error行の消去を行なってみてください。

実行

コンパイルするには
Linux:

$ ghc --make tecstcl.hs

OSX:

$ ghc -framework OpenCL -o testcl testcl.hs

とします。

これを実行すると以下の結果が得られます。

./testcl
Original array = [0.0,1.0,2.0,3.0,4.0,5.0,6.0,7.0,8.0,9.0,10.0,11.0,12.0,13.0,14.0,15.0,16.0,17.0,18.0,19.0,20.0]
Result array   = [0.0,2.0,4.0,6.0,8.0,10.0,12.0,14.0,16.0,18.0,20.0,22.0,24.0,26.0,28.0,30.0,32.0,34.0,36.0,38.0,40.0]
compare
cpu : [0.0,2.0,4.0,6.0,8.0,10.0,12.0,14.0,16.0,18.0,20.0,22.0,24.0,26.0,28.0,30.0,32.0,34.0,36.0,38.0,40.0]
gpu : [0.0,2.0,4.0,6.0,8.0,10.0,12.0,14.0,16.0,18.0,20.0,22.0,24.0,26.0,28.0,30.0,32.0,34.0,36.0,38.0,40.0]
same? : True

OpenCLデバイス側で実行した値は渡す前に比べて2倍になっていますね。

HaskellからOpenCLを使ってみた一例でした。

広告

コメントを残す

以下に詳細を記入するか、アイコンをクリックしてログインしてください。

WordPress.com ロゴ

WordPress.com アカウントを使ってコメントしています。 ログアウト / 変更 )

Twitter 画像

Twitter アカウントを使ってコメントしています。 ログアウト / 変更 )

Facebook の写真

Facebook アカウントを使ってコメントしています。 ログアウト / 変更 )

Google+ フォト

Google+ アカウントを使ってコメントしています。 ログアウト / 変更 )

%s と連携中

%d人のブロガーが「いいね」をつけました。