Perl meets GPU

Perl Advent Calendar, 20日目担当しますpapixです. 宜しくお願い致します.
今日は, ｢Perl meets GPU｣というタイトルで, PerlとGPU, そしてGPUを利用したGPGPUについてのお話をさせて頂きたいと思います.

GPGPUとは?

GPGPUとは, General-Purpose computing on Graphics Processing Unitsの略語で, グラフィックボード(GPU)の演算資源を, 画像処理以外の汎用演算に用いる技術のことです.
GPUには, 画像処理を高速に行うために大量のコアが搭載されているので, これらのコアで処理を分散することで, 高速化を実現します.
GPGPUを上手く利用すれば, シミュレーション, 暗号解読, 音声処理などのデータ並列性のある処理を, 高速かつ低コストで実現できるので, 特に学術分野ではいろいろな試みが行われています.
具体的な研究例としては, 地震の解析の高速化, 高速なフーリエ変換, GPGPUを使ったSSLリバースプロキシの高速化などなど, 探せば結構見つかります.

かつてGPUは結構高価で, 個人では手が出しにくかった部分もありましたが, 最近はGPUの価格も安くなってきていますし, 何よりAWSではGPUが利用可能なインスタンスもあるということで, GPGPUに挑戦する敷居はかなり低くなってきています.

というわけで, 今日はそんなGPGPUを, Perlから楽しむためのモジュールを紹介したいと思います.

GPGPUプログラミング

...その前に, GPGPUプログラミングの特徴について, 軽く紹介したいと思います.

GPGPUを利用したプログラムを実装する際には, GPU側(デバイス側)とCPU側(ホスト側)を区別して考えなければなりません.
デバイス側で実行するコードとホスト側で実行するコードは明確に区別され, 特にGPUで実行するコードは｢カーネル関数｣と呼ばれる特殊な関数として記述します.

このように区別しなければならない理由は, デバイス側のメモリはホスト側のメモリと別に用意されているためです.
GPUで並列実行したい計算データは, ホスト側のメモリからデバイス側のメモリに転送する必要がありますし, 逆にGPUの実行結果はデバイス側からホスト側に転送しなければなりません.
よって, GPUプログラミングは次のような流れになります.

• 必要なメモリをデバイス側で確保
• 計算データをホスト側からデバイス側へ転送
• カーネル関数を呼び出し, GPUで並列実行
• 計算データをデバイス側からホスト側へ転送
• デバイス側で確保したメモリを開放

｢メモリの確保｣, ｢デバイス･ホスト間のデータ転送｣などの処理は, GPGPUフレームワークに任せるのが一般的です.
GPGPUフレームワークは様々な実装が開発･提供されていますが, 今日紹介するモジュールは, 全てNVIDIAが開発した｢CUDA｣というGPGPUフレームワークを利用したものです.

あまりここで長々と説明するのはどうかと思うので(反響? があれば, この辺りについてもいずれ詳しくご紹介できればと思います...), そろそろモジュール紹介へ入って行きましょう.

CUDA::Minimal

まずはCUDA::Minimal. 作者はDavid Mertensさんです.
こちらのモジュールはまだCPANにはあがっておらず, GitHubから入手する必要があります: https://github.com/run4flat/perl-CUDA-Minimal

作者の方が, blogs.perl.orgにこのモジュールについての記事を書いています: Perl's first real CUDA bindings released
こちらの記事にサンプルコードがあるので, 早速見てみましょう(元のサンプルコードにあったコメントは削除しています).
このスクリプトは, $N_values個の要素を持つ配列に対して, 各要素を3倍するコードになっています.

use strict;
use warnings;
use CUDA::Minimal;
use ExtUtils::nvcc;
use Inline C => DATA => ExtUtils::nvcc::Inline;

# 計算データを生成.
# 今回は $N_values = 10 なので, 1から10までのリストを生成してから, packでfloat型のバイナリへ.
my $N_values = 10;
my $host_data = pack('f*', 1..$N_values);

# 生成した計算データを格納できるだけのメモリをデバイス側で確保.
# MallocFromはメモリの確保と転送を行い, Mallocはメモリの確保のみを行う関数.
# これらの関数は, デバイス側で確保したメモリのポインタを返すので, スカラ変数で受け取る.
my $input_dev_ptr = MallocFrom($host_data);
my $output_dev_ptr = Malloc($host_data);

# MallocFrom/Mallocで確保したデバイス側のポインタと, $N_valuesを引数として渡して, カーネル関数を起動.
# カーネル関数は後述.
invoke_the_kernel($input_dev_ptr, $output_dev_ptr, $N_values);

# SetSizeで, GPUで計算したデータを格納するためのスカラ変数 $result_array を生成.
SetSize(my $results_array, length($host_data));

# Transferで, デバイス側のメモリに確保されているGPUで計算したデータを, ホスト側のメモリへ転送する.
Transfer($output_dev_ptr => $results_array);

# 計算結果を表示. Transferで $result_array にfloat型のバイナリで計算結果が格納されているので, unpackで変換する.
print "$_\n" foreach (unpack 'f*', $results_array);

# デバイス側のメモリを開放する.
Free($input_dev_ptr, $output_dev_ptr);

__END__

__C__

// CUDAで記述したカーネル関数. 先頭の __global__ でカーネル関数であることを示している.
__global__ void triple(float * in_g, float * out_g) {
    out_g[threadIdx.x] = in_g[threadIdx.x] * 3;
}

// カーネル関数のラッパー.
void invoke_the_kernel(SV * in_SV, SV * out_SV, int N_values) {
    float * d_in = INT2PTR(float *, SvIV(in_SV));
    float * d_out = INT2PTR(float *, SvIV(out_SV));

    triple <<<1, N_values>>>(d_in, d_out);
}

CUDA::Minimalは, CUDA C言語で記述したカーネル関数を, Inline::CのようにPerlのスクリプト中に記述することができます.
ただ, CUDA::Minimalはメモリの確保, データ転送, カーネル関数の起動など, GPGPUプログラミングを実装する上で必要最低限
の機能しか提供していません.
そのため, Perlのデータ構造とバイナリを変換する処理(上記コードのpack/unpack)などは, 全て手動で記述する必要があります.

更にCUDA::Minimalを利用したコードは, 言うまでもありませんがGPUが搭載されている環境でしか実行できません.
これらの問題を解決すべく, PerCUDAと呼ばれるPerl向けのGPGPUフレームワークを実装しました.

PerCUDA

PerCUDAは, この記事の筆者(papix)が大学院の研究として開発したGPGPUフレームワークです.
CUDA::Minimalを発展させ, PerlからGPGPUを｢お手軽に楽しめる, 試せる｣ことを目指しています.

PerCUDAは, バックエンドとしてCUDA::DriverAPIと呼ばれるモジュールを利用しています.
このモジュールは, CUDAが提供するAPIのうちの1つ, ｢ドライバAPI｣をPerlから利用するためのモジュールです(ちなみに, CUDAが提供するAPIのもう1つは｢ラインタイムAPI｣と呼ばれていて, CUDA::MiminalはこのAPIを利用しています).
CUDA::DriveAPIについては, 以前｢Hachioji.pm日めくりテックトーク｣に記事を書いたので, そちらも参考にして頂けると幸いです: CUDA::DeviceAPIでPerlからお手軽GPGPUライフ!

...ちなみに余談ですが, 記事では｢CUDA::DeviceAPI｣になっていますが, その後で｢利用してるAPI, デバイスAPIじゃなくてドライバAPIじゃないか! 名前間違ってた!!!｣という事に気がついて, ｢CUDA::DriverAPI｣に改名しています.
こちらのモジュールはまだCPANにアップしていないので, GitHubからご覧下さい: https://github.com/papix/CUDA-DriverAPI

さて, PerCUDAを利用することで, GPGPUをどのように実装できるか, 見てみましょう.
次のコードは, 3 x 3 のサイズの配列に対して, 配列の乗算を行うスクリプトです.

use strict;
use warnings;
use PerCUDA;

my $array = [
    [ 1 .. 3 ],
    [ 4 .. 6 ],
    [ 7 .. 9 ],
];

my $runner = PerCUDA->new(
    dim      => 2,        # 配列の次元
    size     => 'array1', # 配列のサイズ
    engine   => 'gpu',    # 使用する処理系
    argument => [         # 引数とデータ型
        array1 => 'f',
        array2 => 'f',
        array3 => 'f',
        N => 'i',
    ],
    retval => ['array3'], # 返り値
    source => <<'KERNEL', # カーネル関数
    
my $tmp = 0;
for (my $i = 0; $i < $N; $i++) {
$tmp = $tmp
    + $array1->[$i]->[$y]
    * $array2->[$x]->[$i];
}
$array3->[$x]->[$y] = $tmp;
    
KERNEL
);

my $retval = $runner->run($array, $array, []);

PerCUDAの特徴は, ｢カーネル関数をPerlで記述できる｣という点です.
現状, Perlで利用できる全ての構文が利用できるわけではありませんが, 使い慣れたPerlでカーネル関数を記述できるというのは大きなメリットになるのではないでしょうか.

また, カーネル関数をPerlで記述しているので, PerCUDAでは, カーネル関数の処理系をPerl処理系とGPUから選択することができます.
そのため, GPUが搭載されていない環境であってもスクリプトを実行することができるのです.
上記のコードでは, engineという部分でGPUを利用してカーネル関数を実行するよう指定しています.
ちなみに, 指定がない場合はPerCUDAが自動的に判定を行い, GPUが搭載されている環境であればGPUを, そうでなければPerl処理系を利用して, カーネル関数を実行します.

PerCUDAの難点は, GPU向けのコードをPerlで記述したカーネル関数から自動生成するので, カーネル関数のチューニングが難しいという点です.
チューニングのしやすさと, それによる高速化という点では, CUDA C言語でカーネル関数を記述できるCUDA::MinimalやCUDA::DriverAPIに多少劣ってしまいます.

PerCUDAの性能

PerCUDAを利用して, 約10^6の画素(1024*1024画素)に対するマンデルブロ集合の計算を行った結果を紹介します.
｢GPU｣はカーネル関数の処理系としてGPUを使った場合, ｢Perl｣はカーネル関数の処理系としてPerl処理系を利用した場合の実行時間になります.

n	GPU	Perl
2^0	0.14419s	2.63360s
2^1	0.14523s	3.27114s
2^2	0.14588s	4.24830s
2^3	0.14409s	5.49464s
2^4	0.14605s	7.26325s
2^5	0.15051s	10.08421s
2^6	0.15296s	15.55460s
2^7	0.15346s	26.60217s
2^8	0.15998s	49.52719s
2^9	0.17723s	91.99616s
2^10	0.20646s	175.36165s
2^11	0.26290s	348.06578s
2^12	0.34298s	687.46580s

...というわけで, 圧倒的にPerCUDA(というかGPU)が早いことがわかります.
実行する処理にもよりますが, PerCUDA(そしてGPGPU)を利用すれば, これだけの高速化が実現できるのです.

2013/12/23追記:
｢PerCUDAと素のCUDAの性能評価｣について, ｢Hachioji.pm 日めくりテックトーク｣に寄稿させて頂きました: PerCUDAの性能評価

PerCUDAのこれから

PerCUDAのプロトタイプ(?)は完成していますが, いろいろと出来が酷いので, 来年3月の公開を目指して研究室の後輩でもある@bool_foolと一緒にリファクタリング作業を続けています.

最初に書いたように, PerCUDAは, ｢PerlからGPGPUをお手軽に楽しむ, 試す｣ためのモジュールです.
最初から最後までPerCUDAを利用して高速化... というのではなく, CUDA::MinimalやCUDA::DriverAPIの利用, あるいはPerlからCUDAへの移植を試みる前の足がかりとして使ってもらえるようになればいいな, と思っています.

Perlで記述したコードがあって, ｢これ, GPUで処理したら早くなるかな...?｣という時に, まずPerCUDAで試してみる.
PerCUDAで, ある程度の成果(高速化)が実現できたのであれば, CUDA::MinimalやCUDA::DriverAPIを利用したGPGPUや, PerlからCUDAへの移植によって, さらなる高速化を試みる... そういう風に使ってもらうのが理想的と思っています.

CUDA::MinimalやCUDA::DriverAPIなど, PerlからGPGPUを利用する環境はかなり整ってきている... という印象を最近持っています.
この記事を読んで興味を持った方は, 是非｢PerlからのGPGPU｣を試してみて頂けると嬉しいです.

...PerCUDAのこれからに, どうぞご期待下さい!