RubyでOpenCLを試してみた

OpenCL?

MacOSX10.6では、OpenCLというGPUを扱うための環境が標準で搭載されています。例えば私のMacbookPro(2009Midモデル)ではGeForce9400Mが付いていますが*1、OpenCLを使うことで並列計算処理にこのGPUを利用することが出来ます。GPUはCPUに比べ多くのコアを搭載しており、これを利用することで並列計算速度の大幅な向上が期待できます。同様のGPGPU技術としてNVIDIA向けのCUDAが存在しますが、OpenCLはNVIDIAGPU以外のデバイス(ATIGPUやCell/B.E.等)にも利用できるなど、異種混合な環境での動作を主眼に置いているようです。


RubyでOpenCL

本来OpenCLを利用するにはXCode開いてC言語で色々書いて…という処理が必要なのですが、今回はBarracudaというOpenCLのRubyBinding*2を利用することで、簡単にOpenCLを利用してみました。OpenCLでは、GPU等のデバイスに処理を依頼する『ホスト側』のコードと、依頼されたデバイスが行う処理である『デバイス側』のコードの記述が必要になります。Barracudaを使うことで、この『ホスト側』のコードは全てRubyで記述することができます。『デバイス側』のコードは文字列OpenCL-C言語のコードを書いて渡すような形式になります。

Install

BarracudaをMac OSX 10.6環境にインストールします。
自分の場合RVMでRubyを入れているため、インストールスクリプトに少し変更が必要でした。

git clone git://github.com/lsegal/barracuda.git
cd barracuda
vim Rakefile # sudo権限で入れようとしてしまうので、SUDO = "" に書き換える
rake install

HelloWorld

まずは並列計算とは全く関係なく、単純にGPUを利用してHello,world!を表示してみます。OpenCL-C言語には標準入出力が無いため、デバイス側で生成した文字列をホスト側で受け取って出力します。GPUを利用するための初期化処理等は全てBarracudaが行ってくれるため、実際の処理にのみ注力してコードを記述できます。

# Barracudaの読込み
require "rubygems"
require "barracuda"
include Barracuda

# デバイス側プログラムをOpenCL-C言語で定義
program = Program.new <<-"eof"
  __kernel void hello(__global char* output) {
    output[0] = 'H';
    output[1] = 'e';
    output[2] = 'l';
    output[3] = 'l';
    output[4] = 'o';
    output[5] = ',';
    output[6] = ' ';
    output[7] = 'W';
    output[8] = 'o';
    output[9] = 'r';
    output[10] = 'l';
    output[11] = 'd';
    output[12] = '!';
  }
eof

# 13個のChar型要素を持つ配列を用意
output = Buffer.new(13).to_type(:char)

# 結果を格納するための配列を渡してGPUに処理させる
program.hello(output)

# ASCIIコードで返ってくるので、Stringクラスに変換後、連結してから出力する
# [72,  101,  108,  108,  111,  44,  119,  111,  114,  108,  100,  33]
# -> ['H', 'e', 'l', 'l', 'o', ',', 'w', 'o', 'r', 'l', 'd', '!']
# -> "Hello, World!"
puts output.map{|num| num.chr }.join

並列処理

次は並列処理をCPUとGPUそれぞれに任せて、比較してみます。
モンテカルロ法を用いてn回の試行でdim次元球の体積を求めるというプログラムを作りました。

require "rubygems"
require "barracuda"
include Barracuda

### これは何?
# モンテカルロ法を用いたN次元球の体積計算をベンチマークに用いた、
# CPUとGPUの実行時間の比較を出力するプログラムです

# 乱数配列生成、半径内要素数計算、体積計算
def cpu(n, dim, rad)
  points = (1..n).map{|i| Array.new(dim){ rand } }
  count = points.select{|point|
    point.map{|x| x**2 }.inject{|sum,x| sum+=x } ** (0.5) < rad
  }.count
  volume(n, dim, rad, count)
end

# 実行時間を計測して出力する
def calc_time(opt={}, &block)
  opt[:times] ||= 1
  opt[:format] ||= "%f"
  start = Time.now
  opt[:times].times{|i| yield }
  finish = Time.now

  opt[:format] % [(finish - start) / opt[:times]]
end

def volume(n, dim, rad, count)
  count.to_f / n * (rad ** dim) * 8
end

# 乱数配列生成、GPU処理、半径内要素数計算、体積計算
def gpu(n, dim, rad, program)
  randoms = Array.new(n * dim){ rand }
  output  = Buffer.new(n)
  program.inner(dim, rad, randoms, output, :times => n)
  count = output.select{|x| x == 1 }.count
  volume(n, dim, rad, count)
end

# 試行回数x次元個の乱数配列から、試行回数個の配列Aを返す
# Aの中身は、ランダムに生成された1点が半径内にあれば1である要素
program = Program.new <<-"eof"
  __kernel void inner(
    __global int dim,
    __global float rad,
    __global float* randoms,
    __global int *out) {

    int id = get_global_id(0);
    int base = id * dim;
    float sum = 0.0;

    for (int i=0; i<dim; i++) {
      sum += randoms[base+i] * randoms[base+i];
    }
    out[id] = isless(rootn(sum, 2), rad);
  }
eof

n     = 1000
times = 10
rad   = 1.0
dims  = (0..10).map{|x| 2**x }

puts "n:#{n}, times:#{times}"
puts "%6s, %8s, %8s" % ["dim", "GPU", "CPU"]
dims.each do |dim|
  puts "%6d, %8f, %8f" % [
    dim,
    calc_time(:times => 10){ gpu(n, dim, rad, program) },
    calc_time(:times => 10){ cpu(n, dim, rad) },
  ]
end

=begin
n:1000, times:10
   dim,      GPU,      CPU
     1, 0.041787, 0.002888
     2, 0.040524, 0.004046
     4, 0.044142, 0.007514
     8, 0.045412, 0.009012
    16, 0.055316, 0.015790
    32, 0.057869, 0.029709
    64, 0.074391, 0.067222
   128, 0.099503, 0.121594
   256, 0.143888, 0.212903
   512, 0.272613, 0.420411
  1024, 0.439119, 0.838492

n:10000, times:10
   dim,      GPU,      CPU
     1, 0.044383, 0.021451
     2, 0.045894, 0.032384
     4, 0.060762, 0.058683
     8, 0.079788, 0.089311
    16, 0.109284, 0.165430
    32, 0.170022, 0.322847
    64, 0.277251, 0.563807
   128, 0.559522, 1.125359
   256, 1.050201, 2.252186
   512, 1.969891, 4.539807
  1024, 3.854168, 8.958749
=end


次元数dimが少ない状態ではCPUの方が高速ですが、dim=64辺りからGPUの方が高速になっています。試行回数をn=1000(立方体の中に1000個の点を置いて半径内外で判断)で試して更に10回の平均を取っていますが、nの数が増えるに従ってその差は顕著に現れると思われます(但しGPUのコア数や転送量の関係でどこかで頭打ちになる)。

注意点

なお、GPU内でOpenCL-C言語を用いて乱数を生成する方法が分からなかったので、乱数配列を渡しています。書籍『OpenCL入門』にMersenneTwisterで乱数を生成する実装が書いてあったのですが、上手く利用できませんでした。また本家MTのC言語実装であるmt19937ar.cも、そのままでは動かせず*3、乱数を生成する種となる『状態』の並列化を考えないといけない等の問題があったため、利用できませんでした。

まとめ

以上、Barracudaを利用してOpenCLをRubyで扱う方法を簡単に説明しました。現状BarracudaはMacOSX10.6にしか対応していないのですが、今後FedoraLinux系OSで動作するようになれば、AmazonEC2 GPU InstanceでRubyを利用して簡単にジョブオフロードを行うようなことも出来るのではないでしょうか。


OpenCLの情報は日本語圏ではまだ乏しく、ましてRubyBindingの情報など恐らく皆無ですが、何かの参考になれば幸いです。
幸い最近のMacではOpenCL環境をOS標準機能としてサポートしているようなので、興味がある方は是非利用してみてブログ等でアウトプットして頂ければと思います。

*1:システムプロファイラ等で確認できます

*2:OpenCLをRubyで記述できるようにしたもの

*3:OpenCL-C言語で対応していない計算箇所がどこかにあるのかも。具体的にはTempering直前辺り