Help us understand the problem. What is going on with this article?

ZEAM開発ログv0.1.6 Elixir から Rustler で GPU を駆動しよう〜ElixirでAI/MLを高速化

More than 1 year has passed since last update.

(この記事は、「fukuoka.ex(その2) Elixir Advent Calendar 2017」の12日目,「機械学習と数学 Advent Calendar 2017の6日目です)

昨日は @twinbee さんの「Elixirから簡単にRustを呼び出せるRustler #4 SHIFT-JIS変換を行う」でしたね。

「ZEAM開発ログ 目次」はこちら

おしらせ

お礼:各種ランキングに69回のランクインを達成しました

4/27から、30日間に渡り、毎日お届けしている「季節外れのfukuoka.ex Elixir Advent Calendar」「季節外れのfukuoka.ex(その2) Elixir Advent Calender」ですが、Qiitaトップページトレンドランキングに8回入賞、Elixirウィークリーランキングでは5週連続で1/2/3フィニッシュを飾り、各種ランキング通算で、トータル69回ものランクイン(前週比+60.1%)を達成しています

みなさまの暖かい応援に励まされ、合計452件ものQiita「いいね」(前週差+103件)もいただき、fukuoka.exアドバイザーズとfukuoka.exキャストの一同、ますます季節外れのfukuoka.ex Advent Calendar、頑張っていきます

image.png

さて本題〜はじめに

ElixirでAI/MLを高速化すべく,今回はElixirからGPUを駆動しました。

本連載の前回記事はこちら
|> ZEAM開発ログv0.1.0 Flow / GenStage による並列プログラミング入門
|> ZEAM開発ログv0.1.1 AI/MLを爆速にしたい! Flow / GenStage でGPUを駆動できないの?
|> ZEAM開発ログv0.1.2 AI/MLを爆速にしたい! Flow のコードを OpenCL で書いてみる〜CPU編
|> ZEAM開発ログv0.1.3 AI/MLを爆速にしたい! Flow のコードを OpenCL で書いてみる〜GPU編
|> ZEAM開発ログv0.1.4 Python/NumPyとElixir/Flow一本勝負!ElixirはAI/ML業界に革命をもたらすか!?
|> ZEAM開発ログv0.1.5 Elixir から Rustler でネイティブコードベンチマークを呼び出してみよう〜ElixirでAI/MLを高速化

今までのまとめとしては次の通りです。

  1. OpenCLでGPUを利用した時にはC言語の1並列と比べて3.95倍の速度向上,Elixirの同等プログラムと比べて10.8〜11.9倍の速度向上になりました。
  2. OpenCLを使わずにマルチコアかつSIMD命令やAVX命令を使った場合は,GPUの場合より高速になる可能性があります。見積もりではElixirの同等プログラムと比べて13倍前後の速度向上を期待できそうです。
  3. GPUの場合は,データの転送に時間がかかっているので,データの転送量に比べて演算負荷が大きくなればなるほど,CPUよりGPUの方が有利になると思われます。
  4. CPUとGPUで適性を見極めて適切に負荷分散をすること,さらにCPUとGPUを並列実行することで,さらなるパフォーマンスを引き出せる可能性があります。
  5. Elixir / Flow は Python / NumPy より1.5倍前後速いです。
  6. OpenCL(GPU) は Python / NumPy より17倍前後速いです。
  7. 今後,研究・開発が進んだ暁には,PythonからElixirに移行することで,AI/ML処理の実行速度が大幅に向上する可能性があります。
  8. Rustler を使っても pure Elixir から高速化できませんでした。
  9. Flow は便利ですが,オーバーヘッドがかなりあります。
  10. Rustler とC言語やOpenCLと比べるとまだまだ速度差は歴然としています。さらなる高速化に向けてはデータ表現の相互変換をどのように工夫するかが鍵になりそうです。LLVMで直にアセンブリコードを出力したらどうなるか興味があります!

今回はいよいよElixirからRustlerとOpenCL経由でGPUを駆動してみたいと思います。GPU駆動について書いたこれまでの連載の集大成です!

Rust の最適化

Rust でのコード最適化のやりかたを @tatsuya6502 さんに教わりました! ありがとうございます。

Cargo.toml に次のような記述を加えます。

# 最高の性能を得るためにreleaseプロファイルを調整する 
[profile.release]
# 最適化レベル。デフォルトは2
opt-level = 3

# LLVMのcodegenの多重度。デフォルトは16。
# 数字を増やすとコンパイル時間が短くなるが、最適化の機会が減ってしまう。
# 1にするとコンパイル時間が長くなるが、コンパイル後のバイナリでは最高の性能が得られる。
codegen-units = 1

また,C言語のプログラムを int64 を使うようにしないとオーバーフローするということだったので,修正しました。

以上を踏まえて全面的にベンチマークを取り直すことにしました。

Rust から OpenCL を利用するには

ocl というライブラリがありました。GitHubはこちら

使い方は README や ocl/examples に書かれているコードを読んでください。

Elixir / Rustler から OpenCL を利用する

GitHub レポジトリはこちら

現状では,OpenCLの有無の判定のしかたがわからなかったので,OpenCL ブランチに push しています。下記のようにしてください。

$ git clone git@github.com:zeam-vm/logistic_map.git
$ git checkout OpenCL

ソースコードはこんな感じです。

native/logistic_map/src/lib.rs の OpenCL 周り

#[macro_use] extern crate rustler;
#[macro_use] extern crate lazy_static;

extern crate ocl;

use rustler::{NifEnv, NifTerm, NifResult, NifEncoder, NifError};
use rustler::types::list::NifListIterator;
use ocl::{ProQue, Buffer, MemFlags};

rustler_export_nifs! {
    "Elixir.LogisticMapNif",
    [("call_ocl", 3, call_ocl)],
    None
}

fn logistic_map_ocl(x: Vec<i64>, p: i64, mu: i64) -> ocl::Result<(Vec<i64>)> {
    let src = r#"
        __kernel void calc(__global long* input, __global long* output, long p, long mu) {
            size_t i = get_global_id(0);
            long x = input[i];
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            x = mu * x * (x + 1) % p;
            output[i] = x;
        }
    "#;

    let pro_que = ProQue::builder()
        .src(src)
        .dims(x.len())
        .build().expect("Build ProQue");

    let source_buffer = Buffer::builder()
        .queue(pro_que.queue().clone())
        .flags(MemFlags::new().read_write())
        .len(x.len())
        .copy_host_slice(&x)
        .build()?;

    let result_buffer: Buffer<i64> = Buffer::builder()
        .queue(pro_que.queue().clone())
        .flags(MemFlags::new().read_write())
        .len(x.len())
        .build()?;

    let kernel = pro_que.kernel_builder("calc")
        .arg(&source_buffer)
        .arg(&result_buffer)
        .arg(p)
        .arg(mu)
        .build()?;

    unsafe { kernel.enq()?; }

    let mut vec_result = vec![0; result_buffer.len()];
    result_buffer.read(&mut vec_result).enq()?;
    Ok(vec_result)
}

fn call_ocl<'a>(env: NifEnv<'a>, args: &[NifTerm<'a>]) -> NifResult<NifTerm<'a>> {
    let iter: NifListIterator = try!(args[0].decode());
    let p: i64 = try!(args[1].decode());
    let mu: i64 = try!(args[2].decode());

    let res: Result<Vec<i64>, NifError> = iter
        .map(|x| x.decode::<i64>())
        .collect();

    match res {
        Ok(result) => {
            let r1: ocl::Result<(Vec<i64>)> = logistic_map_ocl(result, p, mu);
            match r1 {
               Ok(r2) => Ok(r2.encode(env)),
               Err(_) => Err(NifError::BadArg),
            }
        },
        Err(err) => Err(err),
    }
}

lib/logistic_map_Nif.ex の該当部分

defmodule LogisticMapNif do
  use Rustler, otp_app: :logistic_map, crate: :logistic_map

  # When your NIF is loaded, it will override this function.
  def call_ocl(_x, _p, _mu), do: :erlang.nif_error(:nif_not_loaded)

lib/logistic_map.ex の該当部分

defmodule LogisticMap do

  @logistic_map_size      0x2000000
  @default_prime 6_700_417
  @default_mu 22
  @default_loop 10

  @doc """

  ## Examples

      iex> 1..3 |> LogisticMap.mapCalc9(61, 22, 1)
      [28, 25, 37]
  """
  def mapCalc9(x, p, mu, _stages) do
    x
    |> Enum.to_list
    |> LogisticMapNif.call_ocl(p, mu)
  end

このコードだとリスト長がGPUで確保できるメモリを超えてしまった時にエラーが発生します。その問題を修正するならこちらです。

defmodule LogisticMap do

  @logistic_map_size      0x2000000
  @default_prime 6_700_417
  @default_mu 22
  @default_loop 10

  @doc """

  ## Examples

      iex> 1..3 |> LogisticMap.mapCalc9(61, 22, 1)
      [28, 25, 37]
  """
  def mapCalc9(x, p, mu, _stages) do
    x
    |> Enum.to_list
    |> Stream.chunk_every(100000000000)
    |> Enum.map(& &1 |> LogisticMapNif.call_ocl(p, mu))
  end

ただし実験してみると, Stream.chunk_every でけっこうオーバーヘッドが発生することがわかったので,今回は前者で実験してみました。この辺りは課題ですね。

コンパイルしてみると,OpenCLがインストールされている場合には,完全に次の1コマンドで実行できます。ビルドと実行が容易なのは,とても嬉しいですね! mix と cargo のおかげです!

$ mix run -e "LogisticMap.allbenchmarks"

@twinbee さんの実験によると,GCE の GPU 付きでも,OpenCL をインストールすることで動作したそうです! そのうち Qiita の記事にしていただけるんじゃないかと期待しています!

Rust単体実行バージョン

比較のためにRust単体で実行するベンチマークも作成しました。 @tatsuya6502 さん,ありがとうございます!

GitHub レポジトリはこちら

master ブランチがCPU1並列バージョン,OpenCLブランチが ocl を利用しての OpenCL バージョンです。

CPU1並列バージョンの実行方法

$ git clone git@github.com:zeam-vm/logistic_map_rust.git
$ cd logistic_map_rust
$ cargo run --release

OpenCLバージョンの実行方法

$ git clone git@github.com:zeam-vm/logistic_map_rust.git
$ cd logistic_map_rust
$ git checkout OpenCL
$ cargo run --release

実行結果

実行環境は下記の通りです。

Mac Pro (Mid 2010)
Processor 2.8GHz Quad-Core Intel Xeon
Memory 16GB
ATI Radeon HD 5770 1024MB

主要な実行結果は次の通りでした。(条件を揃えるため,すべてベンチマークを取り直しています)

なお,benchmarks_emptyは,Rustler の呼び出しコストを測定するため,一切の計算を無しに引数を読み込むだけ読み込んで結果を Elixir にそのまま返すという処理を記述してみました。

stages benchmarks1 benchmarks3 benchmarks5 benchmarks8 benchmarks9 benchmarks_empty C Rust CPU Rust OpenCL Python
pure Elixir pure Elixir Elixir/Rustler Elixir/Rustler Elixir/Rustler Elixir/Rustler clang Rust Rust Python
loop inlining inside of Flow.map loop, passing by list passing by list, with Window OpenCL(GPU), inlining Ruslter empty CPU, loop CPU, loop OpenCL(GPU), inlining NumPy, CPU
1 53.223249 44.075415 12.042998 9.702550 6.571943 6.027101 2.727346 2.720996 1.857096 17.749182
2 26.474494 20.632323 36.622580 17.419324
4 15.186692 13.497807 36.566198 16.233323
8 12.796444 11.570678 36.226341 15.953845
16 12.962736 11.607420 36.405342 16.846778
32 12.988720 11.488179 39.785561 19.186600
64 13.288242 11.561640 39.646524 24.126142
128 13.214251 12.116703 40.206917 35.373501
  1. benchmarks9(Elixir/Rustler, OpenCL(GPU), inlining)は,pure Elixir(benchmarks1, benchmarks3)のベストタイム(stages=8,32)よりも,1.75〜1.95倍高速です。
  2. benchmarks9(Elixir/Rustler, OpenCL(GPU), inlining)は,Elixir/Rustler/CPU(benchmarks5, benchmarks8)のベストタイム(stages=1)よりも1.5〜1.83倍高速です。
  3. benchmarks9(Elixir/Rustler, OpenCL(GPU), inlining)とbenchmarks_empty(Elixir/Rustler)の差は0.55秒で,Rust OpenCL が1.86秒,この間くらいが正味の Rust による OpenCL の実行時間だと思われます。オーバーヘッド分は Erlang VM の実行コストと,リスト構造から配列に変換するコストです。Erlang VM で実行するのではなく,リスト構造を最初から配列にする最適化も含めたElixirソースコードからの静的コンパイル/最適化をかけてやると,Rust OpenCL くらいの実行時間(pure Elixir との比較で6.18〜6.88倍高速)になる潜在的可能性があります。
  4. C言語の実行時間とRust CPUの実行時間はほぼ等しいです。最適化をかけた場合には,Rustそのものによる実行時間のオーバーヘッドはないものと考えて良さそうです。
  5. Python から pure Elixir は1.39〜1.54倍,Python から benchmarks9(Elixir/Rustler, OpenCL(GPU), inlining) は2.7倍,Python から Rust OpenCL は9.54倍の速度向上です。Python から Elixir に置き換えることで,このくらいの速度向上を期待できそうです!

おわりに

ついに Elixir から Rustler と OpenCL 経由で GPU を駆動することに成功しました!

  1. Elixir/Rustler/OpenCL(GPU) は pure Elixir よりも1.75〜1.95倍高速です。
  2. Elixir/Rustler/OpenCL(GPU) は Elixir/Rustler/CPU よりも1.5〜1.83倍高速です。
  3. Erlang VM の実行コストと,リスト構造から配列に変換するコストが結構かかっています。リスト構造を最初から配列にする最適化も含めたElixirソースコードからの静的コンパイル/最適化をかけてやると,6.18〜6.88倍高速になる潜在的可能性があります。
  4. C言語の実行時間とRust CPUの実行時間はほぼ等しいです。最適化をかけた場合には,Rustそのものによる実行時間のオーバーヘッドはないものと考えて良さそうです。
  5. Python から pure Elixir は1.39〜1.54倍,Python から benchmarks9(Elixir/Rustler, OpenCL(GPU), inlining) は2.7倍,Python から Rust OpenCL は9.54倍の速度向上です。Python から Elixir に置き換えることで,このくらいの速度向上を期待できそうです

とにかく,ビルド設定に関しては,ほとんど何も考える必要なく OpenCL を駆動できたのは素晴らしいです! これは,Rust の ocl ライブラリと,Rustler の功績です。NumPy 互換の CuPy で CUDA 経由で GPU を使うのにあれこれ煩雑な設定がたくさん要ることを考えると,この利点だけでも相当なアドバンテージです。

ただし,Rust プログラミングはかなり熟練を必要とし, @twinbee さんがおっしゃるように,コード1行書くのに1時間かかるというのも決して誇張ではない状況です。今後の研究で Elixir のコードからじかに コンパイルして GPU 駆動できるようにライブラリを整備していきたいと思います。

今後にもどうぞご期待くださいませ。

今回で連載は一区切りですが,次回は Elixir / Rustler の小ネタを披露します。お楽しみに!

明日は @koga1020 さんの「Phoenix + Vue.js 入門」です。こちらもお楽しみに!

p.s.「いいね」よろしくお願いします

よろしければ,ページ左上の image.pngimage.png のクリックをお願いしますー:bow:
ここの数字が増えると,書き手としては「ウケている」という感覚が得られ,連載を更に進化させていくモチベーションになりますので,もっとElixirネタを見たいというあなた,私たちと一緒に盛り上げてください!:tada:

Why not register and get more from Qiita?
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
Comments
No comments
Sign up for free and join this conversation.
If you already have a Qiita account
Why do not you register as a user and use Qiita more conveniently?
You need to log in to use this function. Qiita can be used more conveniently after logging in.
You seem to be reading articles frequently this month. Qiita can be used more conveniently after logging in.
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
ユーザーは見つかりませんでした