rust
Elixir
GPU

ZEAM開発ログ v.0.3.0 Elixir から GPU を駆動するベンチマークをいろいろアップグレードしてみた

(この記事は「fukuoka.ex x ザキ研 Advent Calendar 2017」の25日目です)

昨日は @zumin さんの「RubyMineにElixir Pluginを導入してみた」でした。

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

はじめに

今週からまたGPU駆動の話に戻ります。

さて,今回は @twinbee さんの 「Elixir-NIF-Rustボイラープレート Ruster0.17.1の注意点」 にもありますように,Elixir / Rust 関係で下記のようなアップデートがありましたので,追従してみたいと思います。

  • Elixir 1.6.6 (compiled with OTP 21)
  • Rust 1.27.0
  • Flow 0.14.0
  • Rustler 0.17.1

なお,ocl はバージョン変わらず 0.18 です。

Rustler 0.17.1 へのアップグレード

@twinbee さんの 「Elixir-NIF-Rustボイラープレート Ruster0.17.1の注意点」 にもありますが,Nif*というプレフィックスがほとんどなくなりました。

  • Rename NifEncoder/NifDecoder to Encoder/Decoder
  • Rename NifTerm to Term
  • Rename NifEnv to Env
  • Rename NifError to Error
  • Rename NifAtom to Atom
  • Rename NifBinary to Binary
  • Rename NifPid to Pid
  • Rename NifListIterator to ListIterator
  • Rename NifMapIterator to MapIterator

一部は Nif プレフィックスは残っていますので,注意しましょう。私が使った中では NifResult だけは残っています。
過去に紹介した下記の記事は,現状ではすべて Nif プレフィックスがついていますのでご注意ください。そのうち更新しようとは思っています。
|> ZEAM開発ログv0.1.5 Elixir から Rustler でネイティブコードベンチマークを呼び出してみよう〜ElixirでAI/MLを高速化
|> ZEAM開発ログv0.1.6 Elixir から Rustler で GPU を駆動しよう〜ElixirでAI/MLを高速化
|> ZEAM開発ログv0.1.6.1 Elixir / Rustler 小ネタ集〜 Rustler でリストからバイナリに変換

この更新はそれほど難しくありません。バージョンアップしたあと,エラーメッセージに従って淡々と修正するのみです。ただし,バージョンアップする際に,mix.exs と native/logistic_map/Cargo.toml の2箇所を修正する必要がある点にご注意ください。私は最初 mix.exs だけ修正していたので,コンパイル結果が変わらずに不審に思っていました。

修正したコードは下記の通りです(主要部分を抜き出しました)。

native/logistic_map/src/lib.rs

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

extern crate ocl;

use rustler::{Env, Term, NifResult, Encoder, Error};
use rustler::types::list::ListIterator;
use rustler::types::binary::Binary;

use ocl::{ProQue, Buffer, MemFlags};

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

fn loop_calc(num: i64, init: i64, p: i64, mu: i64) -> i64 {
    let mut x: i64 = init;
    for _i in 0..num {
        x = mu * x * (x + 1) % p;
    }
    x
}

fn map_calc_list<'a>(env: Env<'a>, args: &[Term<'a>]) -> NifResult<Term<'a>> {
    let iter: ListIterator = try!(args[0].decode());
    let num: i64 = try!(args[1].decode());
    let p: i64 = try!(args[2].decode());
    let mu: i64 = try!(args[3].decode());

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

    match res {
        Ok(result) => Ok(result.iter().map(|&x| loop_calc(num, x, p, mu)).collect::<Vec<i64>>().encode(env)),
        Err(err) => Err(err),
    }
}

fn trivial(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: Env<'a>, args: &[Term<'a>]) -> NifResult<Term<'a>> {
    let iter: ListIterator = try!(args[0].decode());
    let p: i64 = try!(args[1].decode());
    let mu: i64 = try!(args[2].decode());

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

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

あ,trivial っていう関数名,変えたいな。。。あとで変えよう。

Flow 0.14.0 へのアップグレード

次は Flow 0.14.0 による変更点です。Flow の GitHub 上の CHANGELOG くらいしか情報がなく,とても難儀しました。

具体的には,Window で指定する :resetmap_state が廃止されて,代わりに on_trigger で調整するという感じです。

今回影響を受けるのは,benchmarks8で使用している mapCalc8 です。

lib/logistic_map.ex

Flow 0.13.0 までのコード

  def mapCalc8(list, num, p, mu, stages) when stages > 1 do
    window = Flow.Window.global
    |> Flow.Window.trigger_every(LogisticMapNif.get_env("LogisticMapNif_map_calc_list"), :reset)

    list
    |> Flow.from_enumerable
    |> Flow.partition(window: window, stages: stages)
    |> Flow.reduce(fn -> [] end, fn e, acc -> [e | acc] end)
    |> Flow.map_state(& &1 |> LogisticMapNif.map_calc_list(num, p, mu))
    |> Flow.emit(:state)
    |> Enum.to_list
  end

Flow 0.14.0 のコード

  def mapCalc8(list, num, p, mu, stages) when stages > 1 do
    window = Flow.Window.global
    |> Flow.Window.trigger_every(LogisticMapNif.get_env("LogisticMapNif_map_calc_list"))

    list
    |> Flow.from_enumerable
    |> Flow.partition(window: window, stages: stages)
    |> Flow.reduce(fn -> [] end, fn e, acc -> [e | acc] end)
    |> Flow.on_trigger(fn acc -> {(acc |> LogisticMapNif.map_calc_list(num, p, mu)), []} end)
    |> Enum.to_list
  end

違いは3箇所。

(1) Flow.Window.trigger_every から :reset を取り除く

    |> Flow.Window.trigger_every(LogisticMapNif.get_env("LogisticMapNif_map_calc_list"), :reset)
    |> Flow.Window.trigger_every(LogisticMapNif.get_env("LogisticMapNif_map_calc_list"))

(2) Flow.map_stateFlow.on_trigger に置き換えて,:reset の場合の書き方にする

    |> Flow.map_state(& &1 |> LogisticMapNif.map_calc_list(num, p, mu))
    |> Flow.on_trigger(fn acc -> {(acc |> LogisticMapNif.map_calc_list(num, p, mu)), []} end)

(3) Flow.emit を取り除く

    |> Flow.emit(:state)

とくに困ったのが,(3)がわからなかったことです。Flow.emitを書いたままにすると謎の実行時エラーに見舞われるのですが,この対処方法が探せど探せど見つからない。困った挙句,先頭から1行1行実行を確かめながら対応しました。

Elixir 1.6.6 (compiled with OTP 21) へのアップグレード

この機会に Elixir と Erlang の両方をバージョン切り替えできるようにセットアップしなおしました。

Elixir 1.6.6 (compiled with OTP 21) をインストールする手順は次の通りです。

Erlang を erlenv でインストールする

すでに Erlang を brew 等でインストールしていた場合には一度アンインストールしましょう。(erlenv でインストールしていた場合を除く)

その後,次のコマンドを順に実行します。

$ cd
$ git clone https://github.com/talentdeficit/erlenv.git .erlenv

ここで ~/.bash_profile に次の記述を加えます。

export PATH=$HOME/.erlenv/bin:/usr/local/opt/openssl/bin:$PATH
eval "$(erlenv init -)"

それから次のように Erlang をインストールします。(Intel Mac の場合)

$ brew install openssl
$ curl -O http://erlang.org/download/otp_src_21.0.tar.gz
$ tar zxf otp_src_21.0.tar.gz 
$ cd otp_src_21.0
$ ./configure --prefix=$HOME/.erlenv/releases/21.0 --enable-smp-support --enable-threads --enable-darwin-64bit --enable-kernel-poll --enable-hipe --without-javac --with-ssl=/usr/local/opt/openssl 
$ touch lib/wx/SKIP lib/odbc/SKIP
$ make
$ sudo make install
$ erlenv global 21.0
$ erlenv rehash

これでインストールできたはずです。erlを実行して下記のように表示されれば成功です。

$ erl 
Erlang/OTP 21 [erts-10.0] [source] [64-bit] [smp:4:4] [ds:4:4:10] [async-threads:1] [hipe]

Eshell V10.0  (abort with ^G)
1> 

ここで CTRL+C を押し,a キーを押して抜け出ます。

チェックポイントは,次の通りです。

  • OTP 21 と表示されているか
  • 64-bit になっているか
  • SMP環境だった時に [smp:*:*] となっているか
  • [hipe] がついているか

もしこれらが異なっていたら,configure で正しく指定されていないか,ハードウェア的にサポートされていないかだと思います。

Elixir を exenv でインストールする

次に Elixir をインストールします。これも brew 等ですでにインストールされていたらアンインストールします。(exenv でインストールしていた場合を除く)

$ brew install elixir-build exenv
$ exenv install -l
$ exenv install 1.6.6
$ exenv global 1.6.6

次のコマンドを実行して次のように表示されれば成功です。

$ elixir -v
Elixir 1.6.6 (compiled with OTP 21)

ベンチマークのソースコードの変更点はありません。

Rust 1.27.0 へのアップグレード

rustup を使っている場合には次のコマンドだけで最新版がインストールされます。超簡単ですね。

$ rustup update

ベンチマーク結果

さてバージョンアップで性能は向上するんでしょうか? 気になるベンチマーク結果は...

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→47.472903 44.075415→37.091920 12.042998→10.768412 9.702550→9.487438 6.571943→6.879992 6.027101→4.520966 2.727346 2.720996→2.931215 1.857096→1.561099 17.749182
2 26.474494→22.943836 20.632323→19.870662 36.622580→32.880617 17.419324→24.325771
4 15.186692→13.509106 13.497807→13.612359 36.566198→35.468315 16.233323→18.880198
8 12.796444→12.146374 11.570678→14.655723 36.226341→35.201705 15.953845→17.437002
16 12.962736→12.180955 11.607420→10.867108 36.405342→35.324410 16.846778→19.406703
32 12.988720→12.366302 11.488179→11.144992 39.785561→38.734335 19.186600→22.267963
64 13.288242→12.297703 11.561640→13.347286 39.646524→39.217209 24.126142→28.936603
128 13.214251→12.196510 12.116703→11.158467 40.206917→39.528575 35.373501→38.189437

Rustler のベンチマークを中心に若干スピードが向上しているようにも思います。が,誤差の範囲かもしれません。バージョンアップによる速度向上は気持ち速くなったかなくらいのようです。

まとめ

バージョンアップにより次のような変更点があります。

  • Rustler 0.17.1 へのアップグレードにより,Nifプレフィックスがほとんど不要になりました。NifResult だけは依然として必要です。
  • Flow 0.14.0 へのアップグレードにより,Window で指定する :resetmap_state が廃止されて,代わりに on_trigger で調整するようになりました。
  • バージョンアップによる速度向上はほとんどないか,あったとして気持ち速くなったかぐらいです。

次回は新しい Rustler の機能を使って,ベンチマーク高速化の試みをやってみましょう。お楽しみに!

明日は @takasehideki さんの「ElixirでIoT#1.3.3:logistic_mapをIoTボードで性能評価してみた」です。こちらもお楽しみに!

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

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