sonots:blog

2017年07月08日23:54
カテゴリ

GPUを支える技術読了

GPUを支える技術を読んだのでメモ。会社のお金で買ってもらいました。ありがとうございます。


1章 [入門]プロセッサとGPU

  • SMが独立の命令列を実行できる単位
  • CPUがアクセスするメモリとGPUがアクセスするメモリは別になっているのが現状は一般的
    • GPUメモリは高バンド幅メモリが必要。CPUメモリはバンド幅はそこまで必要ないが大容量であって欲しい
    • 特質が違って統一させるともったいない
    • Skylakeぐらいになってくると、同じぐらいのメモリでも有効に使える、とかなんとか
  • グラフィックスでは1ピクセルちょっと違っても気にならないので、エラー検知をしない
    • GPGPU向け Fermi では特別に誤り訂正ユニットを載せた
  • SIMDとSIMT
    • SIMD: 複数データにたいして１つのインストラクションを１度に実行する。全体で１つのインストラクション。命令ユニットが少なくてすむ。
    • SIMT: NVIDIA独自(というか提案?)のアーキテクチャ。SIMDの改善。32本のスレッド(１ Warp)に対して１つのインストラクションを１度に実行する。全体としては同時に複数のインストラクションを発行可能(別Streamを指定して発行)。SIMT アーキテクチャでないと実現不可能なのかがいまいちわからないが、Predicate という機構が入っていて、条件分岐処理を書くことができるようになっている。ただし、MIMDのように if の TRUE 側も FALSE 側も並列実行できるわけではなくて、どちらか一方の path のみが同時実行になる。
    • ref. このサイトが詳しかった http://yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html    


4章 [詳説] GPUの超並列処理

  • kepler 以降で、gpu カーネルから gpu カーネルを起動するダイナミックパラレリズムという機能が入っている
  • float16を使うと32に比べて2倍になると思いきや、メモリバンド幅のほうがネックになっていて、4倍ぐらいになる。とにかくバンド幅の進化が、GPUの進化においついていない


5章 GPUプログラミングの基本

  • PTX アセンブラっぽいやつ

    • nvidia は機械語を公開していない。互換性をばんばん壊してる
    • が、ドライバでPTXを機械語に実行時に変換していて、PTXレイヤで互換性を保っている

  • 同期

    • __syncthreads(): 同一ブロック内の全スレッド待ち合わせ
    • __threadfence_block(): メモリアクセス(グローバルメモリとシェアードメモリ)全てが デバイス内(シェアードメモリの場合はブロック内)の全てのスレッドで確認できるようになるまで待機
    • cudaStreamSynchronize(): 指定したストリームの処理完了を待つ
    • cudaDeviceSynchronize(): 指定したデバイスの処理完了を待つ(全ストリームの処理完了を待つ)

  • 一つのワープに含まれる32のアクセスのアドレスが128バイト境界にアラインされた128バイトの範囲内(補足: 128の倍数でメモリ領域が分割されている)にまとまっていれば１回のメモリアクセスでまとめて処理(コアレス)できる。

    • １回で処理できず、もう一度アクセスしにいくことをリプレイという
    • リプレイを避けることが高速化の鍵

  • シェアドメモリは on chip なので、グローバルメモリより速い

    • 列方向で飛び飛びアクセスになる場合は、シェアドメモリに転送しておくと改善できるかも(転送するには一回アクセスはするんだよな？何度もアクセスする場合は一回ですむ、ということかな)
    • シェアドメモリも同じバンクからアクセスすると遅くなるので、ズラしたほうが良いらしい。そこまでしなくても十分速いけど

  • ダブルバッファ

    • ストリームを使って、非同期にメモリー転送すると、転送中に計算させられるから良い。
    • 計算中のメモリ領域は上書きできないので、半分しか使えない。ダブルバッファというテクニック

  • DMA転送する場合に、スワップアウトが起きると困る。起きないようにするのが、pinned memory


感想

  • 単著の割に、内容が重複していることが多いと感じた。
    • 例えば、SIMDとSIMTの話は、何度も出てくる。
  • 内容は網羅的だが、詳細までは説明されていないと感じた。
    • 重要なキーワードは漏れなく載っている。
    • しかし、わかりにくいものは、わかりにくいままで、結局ウェブ検索した。
    • 例えば、SIMT がわかりにくかった。
    このエントリーをはてなブックマークに追加
2017年06月08日18:51
カテゴリ

gdb で ruby プロセスのC レベルとRubyレベルのバックトレースを表示する方法

ちょっと思いついて、sigdump gem を入れていなくても、gdb で ruby プロセスにアタッチすれば、Ruby レベルのバックトレースを簡単に取れるんじゃなかろうかと思って試してみた。C レベルのバックトレースも取りたい。

(1) rb_print_backtrace() と rb_backtrace()

CRubyの内部に、Cレベルバックトレースを出力するrb_print_backtrace関数と、Rubyレベルバックトレースを出力する rb_backtrace関数があるので、それを利用する。

gdb で ruby プロセスにアタッチして、

gdb $HOME/.rbenv/versions/2.4.1/bin/ruby [PID]

以下のようにコマンドを打つ。

call write(2, "== c backtrace ==\n", 18)
call rb_print_backtrace()
call write(2, "== ruby backtrace ==\n", 21)
call rb_backtrace()

しかし、これらではカレントスレッドのバックトレースしか出せない。
また、gdbプロセスの標準出力ではなく、対象プロセスのSTDERRに出力がでるので、どこに結果が出たのかわからなくなりそう、という懸念があった。

(2) backtrace と rb_eval_string

Cレベルバックトレースは gdb の機能である backtrace コマンドで表示できるので、スレッド一覧を info threads で取って、全スレッドに対して backtrace コマンドを打ち込めば取れる。これは何も問題ない。

Rubyレベルバックトレースを、call rb_eval_string("コード") に以下のような Ruby コードを打ち込んで取得するというアイデアを思いついた。

File.open('/tmp/dump', 'a') {|f|
  Thread.list.each {|th|
    f.write %Q[  Thread #{th} status=#{th.status} priority=#{th.priority}\n]
    th.backtrace.each {|bt|
      f.write %Q[      #{bt}\n]
    }
  }
}

実際にやってみると、gdb を detach してもプロセスが復帰しなくなってしまった。。。 gdb から ruby コードを実行するとなんやかんやあって壊れる。

メモ: perl なら gdb から perl コードを実行できるのか ... ? 羨ましい => 追記: この任意の perl コードを gdb から実行するやつは、安全にperlコードを実行できるように perl のイベントループの特定の箇所に breakpoint を貼ってから perl コードを実行する工夫をしているらしい。ただそれでも centos 5.8 の system perl だと segv 起きたりしたとのこと by hirose31

(3) rubyレポジトリにある .gdbinit で定義されている rb_ps を使う

もうすでにアイデア倒れだった感があるのだが、rubyレポジトリに .gdbinitというデバッグ用 gdb スクリプトがあり、そこに C レベルおよび Ruby レベルのスタックトレースを表示する rb_ps という関数が実は定義されているので、それを使うことができる。

gdb -x .gdbinit $HOME/.rbenv/versions/2.4.1/bin/ruby [PID]

(gdb) rb_ps
* #<Thread:0x7f46bb0a5ee8 rb_thread_t:0x7f46bb0725d0 native_thread:0x7f46ba514740>
0x7f46ba16d700 <thread_join_m at thread.c:980>:in `join'
loop.rb:17:in `<main>'
* #<Thread:0x7f46bb202750 rb_thread_t:0x7f46bb3e03d0 native_thread:0x7f46b89c0700>
0x7f46ba0e4f30 <rb_f_sleep at process.c:4388>:in `sleep'
loop.rb:6:in `block (2 levels) in <main>'
0x7f46ba1a72b0 <rb_f_loop at vm_eval.c:1137>:in `loop'
loop.rb:4:in `block in <main>'
* #<Thread:0x7f46bb202660 rb_thread_t:0x7f46bb3e47e0 native_thread:0x7f46b87be700>
0x7f46ba0e4f30 <rb_f_sleep at process.c:4388>:in `sleep'
loop.rb:13:in `block (2 levels) in <main>'
0x7f46ba1a72b0 <rb_f_loop at vm_eval.c:1137>:in `loop'
loop.rb:11:in `block in <main>'

表示に 2.6 sec ぐらいかかってちょっと遅い感はあったが、しっかり動く。
rb_ps は関数呼び出しを使ってないので、core ファイルに対しても使える。便利。

結論

CRuby の C level interaface の変更に追随するのは大変なので、ruby core team にメンテナンスされている .gdbinit を持ってきて使うのが一番楽、という結論に至ってしまった。

gdb を live process に対してアタッチすると、その間プロセスが止まってしまうので注意。固まったプロセスの調査に使うとか、rb_ps は core ファイルに対しても使えるので gcore コマンドで core を吐かせてからゆっくり調査する、とかやると良いだろう。

追記: gem にしました > https://github.com/sonots/gdbdump-ruby

    このエントリーをはてなブックマークに追加
2017年06月07日22:35
カテゴリ

CアプリケーションからRubyスクリプトを呼び出す方法

Cアプリケーションから libruby を使ってRubyコードに定義されている関数を呼び出す方法について記載する。
対象の Ruby バージョンは 2.4.1。

libruby の生成

libruby は rbenv で ruby をインストールすると、実はすでにできているのでご利用いただける。

~/.rbenv/versions/2.4.1/lib
~/.rbenv/versions/2.4.1/include/ruby-2.4.0/

ここに libruby-static.a とヘッダファイルがある。

利用中の ruby から lib と include パスを動的に取り出す

以下のコマンドでビルドに必要なパスを動的に取り出せる

$ ruby -e 'puts RbConfig::CONFIG["libdir"]'
/Users/seo.naotoshi/.rbenv/versions/2.4.1/lib

$ ruby -e 'puts RbConfig::CONFIG["LIBS"] + " " +  RbConfig::CONFIG["LIBRUBYARG_STATIC"]'
-lpthread -ldl -lobjc -lruby-static -framework CoreFoundation

$ ruby -e 'puts RbConfig::CONFIG["rubyhdrdir"] + File::SEPARATOR + RbConfig::CONFIG["arch"]'
/Users/seo.naotoshi/.rbenv/versions/2.4.1/include/ruby-2.4.0/x86_64-darwin16

$ ruby -e 'puts RbConfig::CONFIG["rubyhdrdir"]'
/Users/seo.naotoshi/.rbenv/versions/2.4.1/include/ruby-2.4.0

これを利用して Makefile を書くとこんなかんじになる。

TARGET = sample
LIBS = -L $(shell ruby -e 'puts RbConfig::CONFIG["libdir"]') $(shell ruby -e 'puts RbConfig::CONFIG["LIBS"] + " " +  RbConfig::CONFIG["LIBRUBYARG_STATIC"]')
INCLUDE = -I $(shell ruby -e 'puts RbConfig::CONFIG["rubyhdrdir"] + File::SEPARATOR + RbConfig::CONFIG["arch"]') -I $(shell ruby -e 'puts RbConfig::CONFIG["rubyhdrdir"]')

all : $(TARGET)

$(TARGET) : sample.c
        gcc $(INCLUDE) $(LIBS) -o $(TARGET) sample.c

clean :
        rm -f $(TARGET)

Cアプリケーションからrubyコードを呼び出す例

呼び出し対象のRubyコードが以下のようなものだとして、Test::Callee.new#foo を呼び出したいとする。

callee.rb
module Test
  class Callee
    def foo(a)
      puts a
    end
  end
end

Cアプリケーションは以下のように書けば良い。

sample.c
#include "ruby.h"
#include "ruby/encoding.h"


VALUE $kernel;

void init()
{
    // Ruby初期化のおまじない
    ruby_init();
    ruby_init_loadpath();
    rb_enc_find_index("encdb"); // encodingライブラリの初期化
    rb_require("rubygems");
    rb_require("./callee");
}

void run()
{
    VALUE module = rb_const_get(rb_cObject, rb_intern("Test"));
    VALUE klass = rb_const_get(module, rb_intern("Callee"));
    VALUE obj = rb_class_new_instance(0, NULL, klass); // Test::Callee.new
    VALUE str = rb_str_new2("こんにちは");
    rb_funcall(obj, rb_intern("foo"), 1, str); // obj.foo(str)
}

int main()
{
    init();
    run();
}

コメントに書いているのだが、一応解説しておくと、

    ruby_init();
    ruby_init_loadpath();
    rb_enc_find_index("encdb");

が Ruby 初期化のおなじないである。Ruby のスタートアップ処理については usa さんが詳細な記事を書いていたので読むと良い（かもしれない）。

    rb_require("rubygems");
    rb_require("./callee");

次にここで rubygems を require しつつ、今回読み込みたい ruby スクリプトを require している。require_relative 相当のものは使えなかったので、./ をつけているが、カレントディレクトリが変わると動かなくなるので実は微妙な気はしている。

rb_require("callee") として、ビルド後

RUBYLIB=. ./sample

のようにして RUBYLIB を指定しながら実行する方が良いかもしれない。いや、面倒だな。宿題。

    VALUE module = rb_const_get(rb_cObject, rb_intern("Test"));
    VALUE klass = rb_const_get(module, rb_intern("Callee"));
    VALUE obj = rb_class_new_instance(0, NULL, klass);

これは ruby で書くと obj = Test::Callee.new に相当する。

    rb_funcall(obj, rb_intern("foo"), 1, str);

最後にこれで、obj.foo(str) を呼び出している。完。

今回のコード

https://github.com/sonots/libruby-sample においてあります。

    このエントリーをはてなブックマークに追加
そのっつ (SEO Naotoshi)@sonots
A Ruby and Fluentd committer working at DeNA. 記事本文および記事中のコード片は引用および特記あるものを除いてすべて修正BSDライセンスとします。 #ruby #fluentd #growthforecast #haikanko #yohoushi #specinfra #serverspec #focuslight
はてぶ人気エントリー
人気記事
最新記事
ブログ移行
ISUCON8 の予選問題出題を担当した
objdump でインライン展開を追う
perf でアプリケーションのプロファイリング
Rubyアソシエーション開発助成成果報告会で Cumo の報告をしてきた
Optimized C++ を読んだ
RubyKaigi 2018 で Cumo の話をしてきた
pybind11 に contribution した
valgrind --tool=callgrind で CPU profiling
「Linux の仕組み」を読んだ

Google

Template by decoweb