GpuJoin + GpuPreAgg combined kernel

以下のクエリは、t0とt1の2つのテーブルをJOINし、その結果をGROUP BYして出力するものである。
しかし、EXPLAIN ANALYZEの出力には奇妙な点がある。

postgres=# explain analyze
           select cat,count(*),avg(ax) from t0 natural join t1 group by cat;
                                        QUERY PLAN
--------------------------------------------------------------------------------
 GroupAggregate  (cost=955519.94..955545.74 rows=26 width=20)
                 (actual time=5964.955..5964.972 rows=26 loops=1)
   Group Key: t0.cat
   ->  Sort  (cost=955519.94..955523.12 rows=1274 width=44)
             (actual time=5964.943..5964.947 rows=26 loops=1)
         Sort Key: t0.cat
         Sort Method: quicksort  Memory: 28kB
         ->  Gather  (cost=955323.19..955454.23 rows=1274 width=44)
                     (actual time=5964.756..5964.914 rows=26 loops=1)
               Workers Planned: 7
               Workers Launched: 7
               ->  Parallel Custom Scan (GpuPreAgg)  (cost=954323.19..954326.83 rows=182 width=44)
                                                     (actual time=5596.730..5596.735 rows=3 loops=8)
                     Reduction: Local
                     GPU Projection: cat, pgstrom.nrows(), pgstrom.nrows((ax IS NOT NULL)), pgstrom.psum(ax)
                     Unified GpuJoin: enabled
                     ->  Parallel Custom Scan (GpuJoin) on t0  (cost=45544.82..840948.19 rows=100000000 width=12)
                                                               (never executed)
                           GPU Projection: t0.cat, t1.ax
                           Outer Scan: t0  (cost=0.00..976191.14 rows=14285714 width=8)
                                           (actual time=50.762..891.266 rows=100000000 loops=1)
                           Depth 1: GpuHashJoin  (plan nrows: 14285714...100000000, actual nrows: 100000000...0)
                                    HashKeys: t0.aid
                                    JoinQuals: (t0.aid = t1.aid)
                                    KDS-Hash (size plan: 10.78MB, exec: 16.00MB)
                           ->  Seq Scan on t1  (cost=0.00..1935.00 rows=100000 width=12)
                                               (actual time=0.010..26.491 rows=100000 loops=1)
 Planning time: 0.501 ms
 Execution time: 6008.393 ms
(22 rows)

GpuScanとGpuPreAggの間に挟まれたGpuJoinが(never executed)となっているのである。
これは、PostgreSQLのエグゼキュータから見た時に、GpuJoinが一度も呼び出されていないことを意味する。
ただ、その割には、OUTER側の t0 テーブルのスキャン、INNER側の t1 テーブルのスキャン(Seq Scan)はしっかり実行されている。

これは、GpuPreAggの直下にGpuJoinが存在している場合の特別な最適化で、GpuPreAggのGPUカーネルがJOINとGROUP BYの両方のタスクを一気に実行しており、GpuJoinを実行する必要がなかった事を意味している。
(但し、GpuPreAggへの入力として、t0およびt1テーブルを読み出す必要はある)


複数のテーブルを(場合によってはWHERE句付きで)スキャンし、それを何らかのキーで結合して、最終的にGROUP BYを使って集約するというのは非常によくある処理である。
PostgreSQL*1これを内部的にいくつかのステップに分解し、順を追って処理していく。

例えば、以下のような非常に単純なクエリの場合

SELECT cat, count(*), avg(x)
  FROM t0 JOIN t1 ON t0.id = t1.id
 WHERE y like ‘%abc%’
 GROUP BY cat;

先ず最初にSCANのロジックが動作する。幸い、WHERE句の条件は複数のテーブルに跨るものではないので、条件に合致しないレコードはこの時点で捨てられる。これによって、JOINすべき行数を減らすことができる。
次に、SCANの出力はJOINの入力となり、t0.id = t1.id という条件によって t0 テーブルと t1 テーブルを結合する。JOINの結果生成されたレコードは次のAggregation/GROUP BYの入力となり、ここでcat列の値ごとに集計され、行数とx値の平均が出力される。

WHERE句の評価、JOIN、GROUP BY/AggregationはそれぞれPG-StromがGPUで実行可能なワークロードであるが、PostgreSQLの実行計画におけるScan, Join, Aggを単純に置き換えるとどうなるか。

実はCPUとGPUの間でデータ転送のピンポンが発生してしまう。
CPUはまずストレージからデータを読み出し、これをDMAバッファにロードする*2
GpuScan kernel関数はWHERE句の評価を並列実行し、結果をDMAバッファに書き戻す。これは次いでGpuJoin kernel関数の入力になり、t0, t1テーブルの結合処理を並列実行し、同じく結果をDMAバッファに書き戻す。さらにこれは次のGpuPreAgg kernel関数の入力になり、GROUP BYと集計演算の結果をDMAバッファに書き戻す。最後に、CPUが最終ステージの集約を実行して、結果をユーザへ返却する。

いくら非同期処理とはいえ、これだけ何度もCPU⇔GPU間でデータ転送を行うと色々と辛い。

なので、PG-Stromには以前から OUTER SCAN Pull-up という仕組みを実装していた。

これは、比較的単純なWHERE句の評価を GpuJoin 側に取り込んで実行するもので、WHERE句の評価をくぐり抜けたレコードのみをJOIN処理の対象とする事で実現している。

スキャン処理(WHERE句評価)の仕組みは単純なので、例えば、GROUP BYの直下にスキャンがあるようなケース。
例えば以下のクエリのような場合、

SELECT cat, count(*), sum(x) FROM big_table;

同様にGpuPreAggもWHERE句の評価を取り込んで実行するという芸当が可能であった。


今回、新たに実装し、冒頭のEXPLAIN ANALYZE文で実行してみたのは、元々の(GpuScan + GpuJoin)をGpuPreAgg側に取り込み、SCAN→JOINを実行し、さらに集約演算を実行した結果を書き戻すという機能。

一般に、GROUP BYや集約演算の実行によって大規模なデータから非常に小さな集計結果を出力する事が期待されているため、元々のSCAN+JOINやSCAN+PreAggといったシンプルなワークロードの結合に比べると、転送すべきデータ量を削減する効果は非常に大きい事が期待できる。

実際、CUDAのプロファイラでタイムラインを採取してみると

MemCpy(HtoD)(CPU側からGPU側へのデータ転送)と、それを処理するGPU kernel関数の実行は非常に多く発生しているにも関わらず、MemCpy(DtoH)(GPU側からCPU側へのデータ転送)は、最後に一回だけ、しかも僅か 1.2kB が書き戻されているだけである。

このようにシンプルな集計処理であれば、SCAN→JOIN→GROUP BYを1回のGPU kernel呼び出しで済ませてしまう事で、PCIeバス上のデータ転送量を大幅に減らし、必要最小限のデータ転送で済ませる事ができるようになる。

なお、Visual Profilerで見た感じ、各タスクのスケジューリングにまだ相当改善できる余地が残っているので、この辺も追って改良を加えていきたい。

*1:別にPostgreSQLに限った事ではないと思うが

*2:SSD-to-GPUダイレクトSQL実行の事はここでは忘れてほしい

Pascal以降のUnified Memoryを使いたおす。

今でこそTESLA P40に24GBのRAMが載り、コンシューマ向けでもGTX1080Tiに11GBのRAMが搭載されてたりと、GPU側でも10GBを越えるメモリを積むことは珍しくなくなってきた*1
長らく自分の開発環境で頑張ってくれたGTX980は(当時のハイエンド製品だったにも関わらず)4GBのRAMしか積んでおらず、基本的には、希少資源であるデバイスRAMをどのようにアロケーションするかというのは、データベースのワークロードをGPUで処理する上での大問題であった。

例えば、OUTER側から20万行を読み出して*2、INNER側の1万行とJOINする処理を考えた場合、最悪ケースでは20万行×1万行で20億行が生成されることになる(CROSS JOIN)。

もちろん、PostgreSQLの統計情報からある程度の推計は可能であるし、JOINの結果生成される行数というのはコスト推計値にダイレクトに響いてくる要素なので、オプティマイザは効率の悪いJOINを避けるように実行計画を立てる。
ただ、テーブルから抽出するレコード件数の推計なんて割と適当だし、それが何段も重なった結果、データの分布や行数の推計値なんてのは間違う時には派手に間違うものである。

なので、GpuJoinのロジックの中でGPUでの処理結果を格納する結果バッファの大きさを推定するロジックというのは、ある種黒魔術のような状態になっており、極めて保守性が悪かった。また、それだけ頑張って効果があるかというと、外れる時には外れるし、データの分布に偏りがある場合には統計情報など何の役にも立たなかった事もある。
結果がバッファに入りきらない場合、例えば前述の20万行×1万行の場合だと、20万行のうち4万行だけを使ってJOINを行い、その結果を書き戻してから次の4万行を処理するという動作を行う。もちろん、こういったフォールバックとやり直しはコストが高いのでできれば避けたいが、結果バッファのマージンを増やすと、貴重なデバイスRAMを無駄に消費するというのがジレンマであった。

一方で、2012年からずっとPG-Stromを開発している中でGPUも徐々に進化していき、例えば、Kepler世代のTESLAモデルで導入されたDynamic ParallelismはGpuJoinを自然な形で実装する事を可能にした。
同様に、Pascal世代で強化されたUnified Memory*3を使う事で『実行してみるまで結果サイズが分からない(しかも予想以上に増える事がある)』問題に対しても、効率的にデバイスメモリを割り当てる事が可能になる。

以下の図は、PG-Stromのデバイスメモリアロケーションと、物理GPUバイスメモリの消費量を示す模式図である。

PG-Stromでは細かなcuMemAlloc()/cuMemFree()の呼び出しに伴うオーバーヘッド*4を避けるため、1GB単位でデバイスメモリを確保し、内部的にはbuddy allocatorを使って管理している。
Buddy allocatorは単純だがメモリ利用効率はそれほど高くなく、例えば20MBを確保するにも32MB分の領域を必要とする。つまり、この場合12MB分は完全に無駄である。

Kepler/Maxwell世代はデマンドページングに対応していないため、1GBのセグメントを確保した時点で物理メモリを割り当て、それが実際にGPUプログラムで利用されたかどうかに関わらず、物理メモリを確保する。
Pascal/Volta世代はデマンドページングに対応しているため、1GBのセグメントを割り当てたとしても、この時点では物理メモリは割当てられていない。Buddy allocatorで割り当てた領域を、GPUプログラムが実際に使用した時点で初めて物理メモリを割り当てる。
そのため、Buddy allocator側で未割当ての領域まで物理メモリを消費する事はないし、また、32MBや64MBといった"キリのいい"サイズを割り当てる事はあまり多くないため、平均するとBuddy allocatorの割当てたメモリ領域の70~80%程度しか物理メモリは消費していないようだ。

GPU上で仮想メモリ空間が利用でき、実際に使用した分しか物理メモリが消費されないという事は、結果バッファのマージンぎりぎりの所で調整していたロジック自体を不要にできるという事でもある。

以下の図は、PG-Stromが結果バッファをどのように使用しているかを示す模式図である。

PG-Stromの結果バッファ*5は前と後ろから同時に消費されていく。
先頭からはvalues/isnullペアがレコード数の増加に伴って、末尾からは文字列など可変長データの格納用バッファとして。これが衝突すると結果バッファの不足となり、再度バッファを確保し直して再実行という事になる。

これもKepler/Maxwellの世代だとバッファを広く取りすぎると同時に物理メモリを消費してしまうため、結果の行数をある程度精緻に予測して*6バッファを割り当てなければならないが、Pascal/Volta世代では、GPU上の仮想アドレス空間マッピングされるだけなので、気楽にドカンと巨大なバッファを確保し、使った分だけ物理メモリがデマンドアロケーションされるという方針にできる。

この辺の制御をCUDAのインフラに任せてしまえると、実装がかなり楽になる。
実際、黒魔術チックな結果行数の推定に関わるロジックをざっくりと消してしまえたので、PG-Strom v2.0develのGpuJoinの実装は、v1.0に比べると1500行程度小さくなっている。

また、バッファ溢れ時に再度アロケーションをやり直してGPU kernelを再実行したり、部分的に完了した結果だけを書き戻すというロジックはバグの多い箇所だったので、この辺を無くしてしまえるというのはソフトウェアの品質的にもメリットが大きい。

という訳で、今後はPascal世代のGPUを前提とすることにしたい。

*1:まぁ、こういったGPUを搭載するサーバには100GB以上のホストRAMが平気で積んであったりして、相対的にはGPU上のデバイスRAMが貴重である事に変わりはないが

*2:仮に数億行のテーブルを読み出す場合であっても、一回のGPU Kernel呼び出しで処理するのは20万行だけと仮定する。20万行の処理を500回繰り返せば1億行を処理できる。

*3:仮想アドレス空間/ページフォルト/デマンドページングに対応した

*4:細かな領域が多数存在し、マルチスレッドで同時にバシバシalloc/freeすると結構待たされる!

*5:正確にはKDS_FORMAT_SLOT形式の場合

*6:それでも外れる時は外れる

NECを退職し、新会社を立ち上げました。

ご報告が遅れましたが、6月30日付で新卒の2003年から14年あまり勤務したNECを退職しました。
また、本日、東京法務局品川出張所においてヘテロDB株式会社の登記申請を行い、また、併せて新会社のチーフアーキテクト兼代表取締役社長に就任しました。

今後は、前職では実現できなかった、GPUSSDなどヘテロジニアスな計算機資源を活用する事で、高性能、低価格、使いやすさを両立するデータベース製品の事業化を目指していく事になります。

どうぞよろしくお願いいたします。

web: http://heterodb.com/


弊社が入居する西大井創業支援センター(品川区)


10年以上も勤務した会社を辞めてスタートアップを立ち上げるというのは、おそらく人生の中でも上位に食い込むビッグイベントの一つだと思うので、今の決意や創業に至る一連の流れについて記録を残しておこうと思います。
(書き下してみたら意外と長かった!ごめんなさい)

オープンソースとの関わり

大学院を修了しNECに入社したのは2003年4月。
新人研修を経て配属されたのがHPC事業部配下で、今は亡きIA-64向けLinux kernelの開発とサポートを行う部隊。
大学の卒論で少しLinux kernelのデバイスドライバを書いた事があるとはいえ、本格的にLinux kernelのソースコードと格闘するのはこれが初めてで、HPC風味のパフォーマンスを重視するコーディングを徹底的に叩き込まれた。

一つキャリアの転機となったのは、当時リリース間近に控えていたLinux v2.6系の新技術調査という事で、全く偶然にSELinuxの担当となった事。
SELinuxの挙動を調べる中で、AVC(access vector cache)という極めて高い頻度で参照されるにも関わらず排他ロックで保護されている箇所が、HPC視点からはとても許容できないロック競合を引き起こしているように見えた。
IA-64 Linuxの製品では2004年時点でも既に32CPUのマシンがあり、『こんなロック競合を見過ごしたまま客先でSELinuxが使われたら大変なことになる!』と上司に相談し、当時、Linux kernelで使われ始めたばかりのRCU(Read Copy Update)によるSELinux AVCのロックレス化とコミュニティへの提案を行う事となった。
最終的にこの提案はLinux v2.6.9でマージされ、32CPU環境では修正の有無で1000倍以上の性能差が出るほどの顕著な改善となった。自分の中でコミュニティ指向が明確になったのは、この頃、NSARedHatやHPのエンジニアと議論を重ね、世界中で使われることになるソフトウェアの標準を自分が主導して作り上げたのだという経験が大きいと思う。

"好きな事"を仕事にするには?

ただ一方で、自分がSELinuxに関わり続けたいと思っても、事業会社である以上、例えば『Linux v2.6の新機能を調査する』という活動が終わってしまえば、少なくとも表向きは別の業務をこなしつつ、余った時間でSELinuxの活動をするしかなくなる。少なくとも肩身は狭い。
そういった経験から、自分の好きな事に時間を使うにはビジネスとしての裏付けを作らないとダメだという思いを固めつつあった。

幸い、当時のLinux推進センターでは、コミュニティに機能提案やパッチを投稿する事が推奨されており、(主業務は別にあるので100%コミットという訳にはいかないものの)手の空いた時間に開発するテーマを選択する事は比較的自由にできた。
2005年頃、jffs2ファイルシステムのxattr対応やbusyboxへのSELinux関連コマンドなど、組込み領域へのSELinux対応強化を行っていた。これは、製品出荷後長期間メンテナンスフリーで稼働せねばならない組込み機器の特性とSELinuxがマッチして付加価値に繋がるのではないかというアイデアが背景にあった。

また、2006年にはIPA未踏ソフトウェア創造事業に『SELinuxによるPostgreSQLのアクセス制御強化』のテーマで採択され、SE-PostgreSQLの原型を開発した。これも当時、セキュリティソリューションとしてのSELinuxの可能性を考えた時に『OSだけの制御では低レイヤすぎてお金を払うに値する利用シーンを考えられない』という声を多々頂いていたので、『じゃあ、上物のアプリケーションへ範囲を広げる。ビジネスの可能性は分からないので先ずは個人的に作る』という事で始めたものに、IPAがお金を出してくれたというもの。

Security Enhanced PostgreSQL - System-wide consistency in access cont…

その後、2009~2010年頃は『LAPP/SELinux』という事を言っていたと思う。これは、LAMPならぬLAPPスタックのセキュリティポリシーSELinuxで一元管理しようという試みで、モチベーションは全く同じ。より上位のアプリケーションまで対応範囲を広げれば、SELinuxのビジネスとしての価値を見出せる領域を作り出せるのではないかというものである。

LAPP/SELinux - A secure web application stack using SE-PostgreSQL

ちょうど30才前後のこの頃、未踏でスーパークリエータの認定を頂いた事もあり、社内では割と自由に働かせてもらっていた。ただその一方でフラストレーションも溜まっており、『開発したものを早くビジネスにしなさい』と度々お小言を頂いていた一方で、どうすればビジネスにできるのか、上司も周囲も、誰も教えてくれなかったのである。

ただ、これは今だから分かるのは、SE-PostgreSQLのように新しいコンセプトで、まだ市場が存在するのかどうかも分からないような事業企画を、大組織の中で承認プロセスを通せる程度の隙の無さでビジネスプランを作成する方法なんて、上司にも分からなかったのではなかろうか。*1

私がやりたかった事とは?

端的に言えば技術は目的、ビジネスは手段とでもなるだろうか。

自分が楽しいと思う技術、手と頭を動かしたいと思えるソフトウェアを開発するために、自分とチームメンバーの人件費や開発費その他をペイできる程度のビジネスを作りたいという思いは一貫している。
ただ、それでユニコーン企業を目指すとか、世界を変革するために寝食を惜しんで事業を拡大するというのは少し違和感がある。やはり起点はテクノロジであり、面白い事、わくわくする事を(ニッチでもいいので)世の中で役立てたい。その結果として、ホームランかシングルヒットか、もしかしたら空振り三振かもしれないが。

一方で、まだ20代のこの頃は知る由もなかった事だが、NECでは何か新ビジネスを立ち上げる際、事業開始から3年後に数十億円レベルの売上が見込める事というのが一つの判断基準になっており、スモールビジネスというのはそもそも折り合いが悪かった。
この辺は他の大企業でも同様だと思うが、新しい技術を使ったソリューションを事業化するとして、また市場に受け入れられるのかどうかも分からない、ビジネスモデルは仮説段階、市場規模の推定も難しいような状況で、これだけの規模のビジネスを社員にコミットさせるというのは流石に無理筋であるとは思う。

新事業社内公募

2011年2月~2013年11月までNECヨーロッパに出向し、SAP社とのアライアンスを担当する事になった。
当時、サーバ製品の拡販のためSAP HANAの認証取得が重要な課題となっており、Linuxやデータベースに詳しく、外国人との折衝もこなせる人という事で白羽の矢が当たったと伝え聞いている。

これまでオープンソースに関わる仕事だけをやってきた自分にとって、ドイツでの3年間は、NECの顔としての立場で製品認証や保守スキームの折衝から契約交渉、プロプラ技術の検証、果ては機材の購買に至るまで、芸風を広げるのに大いに良い経験であった。

この時期、SAP社との協業の傍ら自分にとっては2つの大きなトピックがあり、これらが今後の方向性を決定づけた事になる。

一つは、PGcon 2011でのTim Child氏のPG/OpenCLの発表に刺激を受けて勉強を始めたGPUによる並列処理。
2012年の初頭にFDWベースで最初の実装を作ってみたところGPUの値段の割に性能改善効果が著しく、まさに"面白い!"技術であった。これがGPUによるデータベース高速化をメインの仕事にする出発点となった。また、この結果をブログにまとめたところ、同時期に中央研究所でGPUによるインメモリDBの研究を行っていた柏木の目にとまる事になる。後に彼はヘテロDB社の共同創業者となる。

もう一つは社内新事業公募。ドイツからの帰任後、また本業の仕事を抱えつつ"余った時間で"PostgreSQL関連のコミュニティ活動を行うという状況に嫌気が差していたので、実は帰国と同時にPostgreSQLの会社にでも転職しようと考えていた*2
ただ、自分の手で開発した技術を軸にビジネスを作りたいという思いをまだ社内で真剣にぶつけてみた訳ではなかったので、一度くらい、社内新事業公募という形で会社のトップにぶつけてみて、それでダメなら辞めれば良いじゃないかという気持ちで、公募に誘ってくれた同期と、柏木、自分の3名で提案を書いてみた。
まぁ、S常務(当時)も社内ブログで『ダメなら会社辞めても事業を興すという気概のある社員の応募を求む』と吠えていた訳なので。

で、結果は応募110件中の2件で採択。辞めさせてもらえなかった。

PG-Stromプロジェクト

帰国後、PG-Stromを軸とする事業立ち上げを目指して、マーケティングとソフトウェア開発を並行して進める事となった。
この段階でのマーケティングとはつまり、この技術を事業化した時に本当に3年後〇〇億円の売上げを生み出せる可能性があるのかどうかを調べ、芽がないのであれば早々に投資を打ち切って無用な出血を抑えるという判断のために行うものである。

もちろん、技術の価値は顧客の課題解決手段としての価値とは異なるし、価値提案(VP: value proposition)の設定次第で、顧客にとってお金を払ってでも使いたい技術なのか、単なるwhat's newに過ぎないのかは変わってくる。当時、社内ではリーン・スタートアップが流行っており、仮説検証と称して、ユーザにVPをぶつけてその反応を伺いながら、提案シナリオや製品機能を修正するというアプローチを採っていた。

ただ、例えばWebサービスであればモックアップのWebサイトでサービスのコンセプトを議論するという事も可能であろうが、基盤ソフトウェアであるデータベース製品が『時々クラッシュします』では顧客にとっての価値を実証するどころでは無いわけで、まだソフトウェアがロクに動きもしない段階で、真っ先にマーケティング活動を始めて(そして、その分開発時間が割かれる事になる!)果たしてそれが正しいアプローチだったのかは今でも疑問が残る。

もちろん教科書的には、顧客にとっての価値シナリオとその時点での製品の完成度は別の概念だが、価値シナリオを納得するためには"実際に動くモノ"を見せるのが一番、百聞は一見に如かずで、ソフトウェアの完成度が未成熟な段階では確実に仮説検証の成功率は下がってしまうハズ。
結局、会社の方針もあり最初に米国市場でのユーザ開拓を行ってはみたものの『it's too early』という事でユーザとの仮説検証までは進む事ができず、自分たちの足で訪問でき、話を聞ける国内ユーザを対象にするよう方針転換する事となった。

退職、そしてヘテロDB社の創業

幸い、国内のユーザ様には何件か興味を持ってもらい、利用シナリオをあれこれ考えたり、ディスカッションを重ねる中で新機能の構想を具体化し、ようやく、技術と顧客にとっての価値が一本の線で繋がる実感は徐々に出てきた。
今になって思い返せば、8年前に『開発したものを早くビジネスにしなさい』とお小言を言われた時に、こういった活動をすれば(それなりに)根拠を持って話ができたんじゃなかろうか。
また、コンセプトや価値提案は様々な伝え方にはなるが、そのベースになるPG-Stromの機能や安定性が時間をかけて徐々に強化されてきたという事も、ユーザ様の納得感を得られる一助になっているだろう。

ただ一方で、少なくとも我々に見えている範囲では、NECの新事業のボリューム感として一桁以上小さい事は否めず、これはまだ2016年の段階でもまだ未解決の課題として残っていた。
しかも売上が年々縮小していく中、新領域に対する投資は非常に厳しくなっているため、2017年度からはマネジメント系の別業務をメインにした上で、PG-Stromはオープンソース活動として"空いた時間で"行わざるを得ないという事になった。

要は、新事業立ち上げを目的とした活動ではなく、空き時間のサイドジョブに戻ったわけで、いわば『10年前の振り出しに戻る』という事である。

自分としては、大組織における壁を打ち破れなかったという事で忸怩たる思いがある一方、ユーザ様とのディスカッションや共同検証を経て、テクノロジーは確実に顧客課題を解決するソリューションに進化しつつあるという確信は持っている。

で、あれば、5年前に自分が作り出したソフトウェアが飼い殺しのような状況のまま、後発の競合が次々に先を行くのを指を咥えてただ眺めているという選択肢はあり得ない。そう考え、NECを出て自らPG-Stromを事業化するための新会社を立ち上げる事を決意した。

転職という選択肢はあったかもしれない。実際『PG-Stromの開発をそのまま継続していいのでウチに来てほしい』という誘いもあったが、投資予算の打ち切りをチラつかせながら、まだソフトウェアの機能・品質が未成熟な状態でマーケティング活動を行わざるを得なかった*3事を思い出すと、プロジェクトのガバナンスを現場から離れた人が握る事への不安はどうしても拭えなかった。

新会社でもやる事は基本的に同じ。PG-Stromをエンタープライズ領域で使える製品として世に出し、データベースという最も広範に利用されているソフトウェアの技術革新を通じて、GPUSSDなどヘテロジニアスな計算機資源をユーザや社会の課題解決に役立て、ビジネスとして成り立たせていくために必要なあらゆる事をやっていく。

がんばるぞー

*1:『いや馬鹿にすんな、そんなの知ってるわ!』という反論は受け付けます。

*2:ただし帰国前に転職すると帰りの飛行機代が出ない

*3:もちろん、幹部も悪意でやらせている訳ではない。彼らは株主から預かった資産を最も効率的に投資する立場なので、パフォーマンスの悪い投資先には早めに見切りをつける必要がある。

スキャン速度 10GB/s への挑戦~その①~

PCIe直結のNVMe-SSDは、コントローラの性能にもよるものの、PCIe x4接続のコンシューマ製品であれば一枚あたり1.8GB/s~3.5GB/s、PCIe x8接続のエンタープライズ製品であれば一枚あたり5.0GB/s~6.0GB/sものスループットを出すことができる。

ただ、実際にはサーバのPCIeスロットに空きがあっても、PCIe x8スロットなので、コンシューマ製品を使うと帯域を余らせてしまったり、エンタープライズ製品では少々お高くなるといった問題がある。

例えば、私が持っているSuperMicro 5018GR-Tというサーバは、TESLAなどパッシブファン型のGPUを搭載できるモデルで、PCIe x16スロットが2つ(GPU用)と、x8スロットが1つ(HHHL)用意されている。
Supermicro | Products | SuperServers | 1U | 5018GR-T
この場合、x16スロットにGPUを搭載すると、物理的に残りスロットは2つ。
エンタープライズ向け製品はそこそこ値段が張るので、なかなか気軽に買うわけにはいかないので、ソフトウェアの開発・デバッグをコンシューマ向けNVMe-SSD製品で行うことになる。そうすると、2枚合わせても帯域は3.6GB/s~7.0GB/s程度となり、GPU側のPCIe x16スロットの帯域を埋め尽くすところまでは到達できない。

そんな折、面白いガジェットを見つけた。
カナダのAmfeltecという会社が開発、販売しているデバイスで、PCIe x16スロットに接続のキャリアボードの上にPCIeスイッチが搭載されており、その先にM.2 SSDを最大4枚搭載することができる。
PCI Express Gen 3 Carrier Board for 4 M.2 SSD modules - Amfeltec

形状は1Slot幅のフルハイト・ハーフレングス(FHHL)なので、少なくともGPUを搭載するようなスロットであれば余裕で搭載できる*1。今まで最大でもSSD x3枚構成でしかスループットを計測できていなかったので、果たして、理論合計帯域がPCIe x16に達した状態で、SSD-to-GPUダイレクト転送が想定通りに機能するのかというのを試すため、このデバイスをカナダから取り寄せてみた。

お値段 540.14USD也 (本体 465.00USD + 送料/手数料 75.14USD)

で、届いたブツがこれ。

これに、M.2接続のコンシューマ製品では最高のスループットを誇るSamsung社のSSD PRO 960を4枚挿すことにする。


※写真では箱が2つですが、512GB版を4枚買いました。

ヒートシンクアイネックスHM-21という高さの低いものを購入。結果的に、このH5.1mmという背の低さが1Uサーバの筐体にうまくフィットすることになった。


こういった感じで、表面/裏面それぞれ2枚ずつM.2 SSDを搭載することができる。


実際にはヒートシンクを搭載するので、SSDのロゴは隠れます。


表面/裏面にそれぞれ2枚ずつ、計4枚のSSD 960PROが搭載されているのがお分かりだろうか。


5018GR-Tサーバに搭載した様子。
奥側のPCIe x16スロットにはTesla P40が、手前のPCIe x16スロットにはキャリアボード経由でSSDが4枚搭載されている。
Samsungのカタログスペックによると、SSD 960PROのSeqReadは3.5GB/sという事なので、もし4枚のSSDからフルスピードでデータを読み出せれば、合計の転送速度は14.0GB/sとなり、目標とする10GB/sを上回ることになる。

さて、どうなるか。

まず、4枚のSSDがOSで認識されていることを確認。

[kaigai@saba ~]$ lspci -v | grep Samsung
05:00.0 Non-Volatile memory controller: Samsung Electronics Co Ltd Device a804 (prog-if 02 [NVM Express])
        Subsystem: Samsung Electronics Co Ltd Device a801
06:00.0 Non-Volatile memory controller: Samsung Electronics Co Ltd Device a804 (prog-if 02 [NVM Express])
        Subsystem: Samsung Electronics Co Ltd Device a801
07:00.0 Non-Volatile memory controller: Samsung Electronics Co Ltd Device a804 (prog-if 02 [NVM Express])
        Subsystem: Samsung Electronics Co Ltd Device a801
08:00.0 Non-Volatile memory controller: Samsung Electronics Co Ltd Device a804 (prog-if 02 [NVM Express])
        Subsystem: Samsung Electronics Co Ltd Device a801

きちんとデバイスファイルも作成されている。

[kaigai@saba ~]$ ls -l /dev/nvme?n1
brw-rw----. 1 root disk 259, 1 May 28 14:08 /dev/nvme0n1
brw-rw----. 1 root disk 259, 3 May 28 14:08 /dev/nvme1n1
brw-rw----. 1 root disk 259, 2 May 28 14:09 /dev/nvme2n1
brw-rw----. 1 root disk 259, 0 May 28 14:09 /dev/nvme3n1

とりあえず、RAID0構成の違いによるスループットの差を見るため、1個のSSDをそれぞれ3つの区画に分け、2枚構成/3枚構成/4枚構成でのデータ転送速度を計測してみる。

# mdadm -C /dev/md2 -c 128 -l 0 -n 2 /dev/nvme0n1p1 /dev/nvme1n1p1
# mdadm -C /dev/md3 -c 128 -l 0 -n 3 /dev/nvme0n1p2 /dev/nvme1n1p2 \
                                     /dev/nvme2n1p2
# mdadm -C /dev/md4 -c 128 -l 0 -n 4 /dev/nvme0n1p3 /dev/nvme1n1p3 \
                                     /dev/nvme2n1p3 /dev/nvme3n1p3

まず、SSD 2枚によるSSD-to-GPUダイレクト転送の実験。

[kaigai@saba utils]$ ./nvme_test /opt/nvme2/100GB
GPU[0] Tesla P40 - file: /opt/nvme2/100GB, i/o size: 100.00GB, buffer 32MB x 6
read: 100.00GB, time: 15.36sec, throughput: 6.51GB/s
sem_wait: 10177ms, nr_ram2gpu: 0, nr_ssd2gpu: 13107200, average DMA blocks: 256.00

スループットは 6.51GB/s で、理論限界3.5GB/sのSSDを2枚束ねたという点では上出来。

avg-cpu:  %user   %nice %system %iowait  %steal   %idle
           0.06    0.00    1.36    0.00    0.00   98.58

Device:            tps    MB_read/s    MB_wrtn/s    MB_read    MB_wrtn
nvme0n1       26398.00      3299.75         0.00       6599          0
nvme3n1           0.00         0.00         0.00          0          0
nvme1n1       26261.50      3282.69         0.00       6565          0
nvme2n1           0.00         0.00         0.00          0          0
md2           52659.00      6582.38         0.00      13164          0
md3               0.00         0.00         0.00          0          0
md4               0.00         0.00         0.00          0          0

実行中、iostatも概ねその辺の値を記録している。

次に、SSD 3枚によるSSD-to-GPUダイレクト転送の実験。

[kaigai@saba utils]$ ./nvme_test /opt/nvme3/100GB
GPU[0] Tesla P40 - file: /opt/nvme3/100GB, i/o size: 100.00GB, buffer 32MB x 6
read: 100.00GB, time: 10.73sec, throughput: 9.32GB/s
sem_wait: 6193ms, nr_ram2gpu: 0, nr_ssd2gpu: 13107200, average DMA blocks: 256.00

スループットは 9.32GB/s まで上がり、理論限界10.5GB/sに対してなかなかの値。

avg-cpu:  %user   %nice %system %iowait  %steal   %idle
           0.08    0.00    1.80    0.00    0.00   98.12

Device:            tps    MB_read/s    MB_wrtn/s    MB_read    MB_wrtn
nvme0n1       25458.00      3182.25         0.00       6364          0
nvme3n1           0.00         0.00         0.00          0          0
nvme1n1       25464.50      3183.12         0.00       6366          0
nvme2n1       25458.50      3182.31         0.00       6364          0
md2               0.00         0.00         0.00          0          0
md3           76382.00      9547.75         0.00      19095          0
md4               0.00         0.00         0.00          0          0

実行中のiostatによれば、各デバイスは概ね3.1GB/sのスループットを記録していた模様。

最後に、SSD 4枚によるSSD-to-GPUダイレクト転送の実験。

[kaigai@saba utils]$ ./nvme_test /opt/nvme4/100GB
GPU[0] Tesla P40 - file: /opt/nvme4/100GB, i/o size: 100.00GB, buffer 32MB x 6
read: 100.00GB, time: 10.56sec, throughput: 9.47GB/s
sem_wait: 5818ms, nr_ram2gpu: 0, nr_ssd2gpu: 13107200, average DMA blocks: 256.00

おや、流石に10GB/s近くなり、他のところで苦しくなってきたのか?
9.5GB/s近辺でデータ転送のスループットが頭打ちとなっている。

avg-cpu:  %user   %nice %system %iowait  %steal   %idle
           0.06    0.00    1.88    0.00    0.00   98.05

Device:            tps    MB_read/s    MB_wrtn/s    MB_read    MB_wrtn
nvme0n1       19392.00      2423.94         0.00       4847          0
nvme3n1       19407.00      2425.88         0.00       4851          0
nvme1n1       19415.50      2426.94         0.00       4853          0
nvme2n1       19407.50      2425.94         0.00       4851          0
md2               0.00         0.00         0.00          0          0
md3               0.00         0.00         0.00          0          0
md4           77621.00      9702.62         0.00      19405          0

iostatを眺めてみると、一枚あたりの転送性能は2.4GB/s程度に落ち込んでいる。
GPU側でデータを受ける方の限界や、あるいはライザカード上のPCIeスイッチが被疑という事も考えられるが、この辺は別な環境を用意して切り分けができるようになるまではペンディングといったところか。

この手のカードは、製品構成を考える時に候補に入れるかどうかは(主に保守サポートの問題で)結構悩ましいが、開発・デバッグ用途でなら単にデータ転送レート10GB/s近くを出せるデバイスという位置づけで十分使える。

というわけで、ストレージの帯域にお悩みの週末プログラマの方は、Amfeltec社 SKU-086-34(PCIe Carrier board for 4 M.2 SSD modules)ボードでも試してみてはいかがでしょうか?

※ なお『その②』以降の予定は今のところ未定です。

*1:さらにパッシブファン用GPU向けの風量が供給されているので、ヒートシンクさえ付ければThermal Throttlingが発動するレベルまで温度は上がらないハズ

GTC2017 - 個人的トピック

個人的トピックまとめ。

ポスター発表

今回、急遽渡米することを決めた理由がこのポスター発表。

セッションやトレーニングと並行して、GTCではGPUを用いた研究開発ネタのポスター展示が行われており、今年は約140件のポスター展示が採択されている。
このうち、プログラム委員の評価が高い20件が、S7480 - Fast Forward Poster Program for the Top 20 Posters のセッションで各々4minづつの発表を行う。さらに Top-20 の中からプログラム委員が事前に選定した5件が "Top-5 Poster Finalist" という扱いで、コンベンションセンターの正面、最も目立つ場所に掲載され、参加者がモバイルアプリ経由でどのポスターが気に入ったかを投票する。
最も票を集めたポスターには GTC2017 Poster Winner Award と副賞$5,000が贈られることになる。

今年は旅費の関係もあって元々渡米する予定はなかったので、セッション発表には申し込まず、ポスターを2枚投稿しただけだった。

このうち、SSD-to-GPUダイレクトSQL実行機能を紹介した P7130 - An intelligent storage for PostgreSQL database がTop-5に選定されたという連絡が来たのが4月29日(土)。Top-20の4分発表だけならちょっと考えたが、Poster Winner Awardの可能性があるならという事で急遽渡米準備。フライトとホテルを抑えてサンノゼに向かったワケである。

ただ、残念ながら結果は及ばず。
NVIDIA-JPの皆様からも他のゲストにご紹介頂いたり、偶然その場にいたNECのメンバに応援してもらったものの、7000人のGTC参加者の投票の結果*1、台湾の学生さんのMACHINE LEARNING領域の研究がPoster Winner Awardに選ばれた。おめでとうございます。
最終選考は一般からの投票だった訳だが、やはり応援してもらったのに結果につながらなかったというのは申し訳ない。次に機会があれば、今度はきちんとTop-1を取れるようにしたいものである。

ただ、今回ポスター発表という事で気付いた事も。GTC16, GTC15と過去2回の参加はいずれもセッション発表のみの参加で、これはある種一発勝負の要素がある。つまり、同じ時間帯に注目度の高いセッションが入っていたり、講演タイトルの"引き"がいま一つだと、わざわざ米国まで出かけて聴衆が10人くらいだとか、そういう悲しい事になる可能性も多いにある。

一方で、空き時間に人がぶらぶら眺めに来るポスターであれば、ちょっと熱心に読んでいる人がいたら『Hello, I'm author of the poster. Let me introduce our research briefly.』なんて声をかければ*2何回でもディスカッションのチャンスはあるし、場合によっては、セッションにただ参加して聞いているだけの人よりも高いレベルの認知/関心を得られるかもしれない。

来年以降、研究開発の発表を行う時には、単にセッションにプロポーザルを出すだけじゃなく、これまでリーチできなかった人の認知/関心を得るためにどういった発表の仕方がベストかもう少し考えてみた方がよさげ。例えば、ピュアな技術の部分はポスターとして、製品や事例に関してはセッションでというやり方があり得るかもしれない。

ちなみに、ポスターに申し込む時には一つ注意があり、プロポーザルを出す段階で既にポスターを作成していなければならないという事。
確か1月下旬が〆切だったが、担当のKaraさんから『ポスターのPDFが添付されてないですよ』と突っ込みがあり、ポスター2枚を大慌てで仕上げる羽目になった*3。ただ、これもちょっと考え物で、1月以降の進捗が発表に反映されないというデメリットはあるとは思う。

[idea] Visual Profiler

その他、参加したセッションの中でPG-Stromに応用が利きそうなものが二つ。
一つは Visual Profiler で、不勉強にして今まで知らなかったのだが、あるマシンでCUDAプログラムを実行し、そこでプロファイラが生成したログファイルを、別のマシンにインポートして可視化するという機能があるらしい。

確かに、CUDA Driver APIを見てみると、以下のようにプロファイラを初期化する関数の中に、出力ファイルを指定するオプションがある。

CUresult
cuProfilerInitialize(const char* configFile,
                     const char* outputFile,
                     CUoutput_mode outputMode);

今まで、PostgreSQL v9.5ベースの実装(つまりCPUパラレルなし)では、例えばデータ転送時間にしてもDMAの開始時刻/終了時刻を記録して、総転送データサイズを時間で割ってやれば、プロファイリングとして比較的直観に一致する結果が得られていた。
しかし、CPU+GPUハイブリッド並列に対応するPostgreSQL v9.6以降では、他のバックエンドとDMA転送がかち合った結果、PCIeバス使用率は高い(= 実際には10GB/sとか出る)一方で、処理時間が間延びした分転送レートが非常に低く見えてしまうという問題があった。
EXPLAIN ANALYZEで手軽に性能情報を確認できるというメリットがある一方で、CPU並列と相性が悪いというのは未解決問題であった。
であれば、細かい事は全部CUDAの実行系に任せて、PG-Stromの側でプロファイリングを頑張らないのも一つの選択肢ではないのかとも思えてきた。

実際、GPU kernel実行時間やDMA Send/Recvといった定量指標だけでなく、GPUタスクがうまい具合にキューに突っ込まれてSSDGPUを遊ばせていないかというのも非常に重要なチューニングポイントではあるが、現状のPG-Stromのプロファイラはこれをきちんと取る事ができない。

[idea] GDF(GPU Data Flame) プロジェクト

MAP-DやH2Oといった、データ解析分野でGPUを使っている人たちが始めたプロジェクトで、要は、GPU上のグローバルメモリに配置したデータをGDFに対応したアプリケーション間で交換するという規格らしい。
確かに、cuMemAlloc()で獲得したメモリ領域は、cuIpcGetMemHandle()を使って識別キーをエクスポートする事が可能で、識別キーを受け取った他のプロセスはcuIpcOpenMemHandle()を呼び出す事でその領域を自分のプロセスでも利用する事ができるようになる。
彼らのアイデアとしては、GPU上の統計解析・機械学習アプリケーションを連携させるために、毎度毎度GPU RAM⇔CPU RAM間のデータ交換は避けたいというもの。GDFの仕様に従ってデータ領域をエクスポートすれば、連携対象のアプリケーションも同じようにその領域を参照できるはずだというもの。

海外としては二つ狙いがある。
一つは、やはり開発リソース・得意分野の関係で自分たちが統計解析・機械学習のエンジンを全て自作するか??というもので、サードパーティと連携できる芽があり、さらにNVIDIAもこれを推しているなら悪くはないだろうという判断。
もう一つは、PL/CUDA関数で1GBを越えるデータサイズを扱う時に問題となっている、PostgreSQL可変長データサイズの問題を実質的にクリアできる可能性があること。

例えば、10GB程度の行列をGPUにロードすることを考える。Tesla P40であれば24GBのデバイスメモリを搭載しており、メモリの割り当て自体はそれほど大変な話ではない。しかし、PostgreSQL可変長データのサイズ制限により、これをPL/CUDAの引数として渡そうとすると、実質的には1個の行列であるにも関わらず複数の2次元配列に分割しなければならないなどの制約が生じる。
では考え方を変えて、GPUバイスメモリ上に確保したメモリ領域を、PostgreSQLからForeign Tableとして読み書きできるようにしてはどうか?
この方法であれば、一度に受け渡すデータサイズが1GBを越える必然性はないし、FDWドライバが内部的に管理するデータ構造であれば(varlena型ではないので)1GBの壁は生じない。

GPU RAM上にデータを保持するForeign Tableを使用する場合、おそらく、PL/CUDA関数への引数はregclass型を用いてForeign TableのOIDが渡るようにすべきだろう。元々データはGPU RAM上に存在していることが前提なので、これらのデータは毎回RAM⇔GPU間を転送する必要がなくなる。
残念ながらTop-20に残らなかった方のポスター(P7129)では、PL/CUDAを用いた類似化合物検索を行っているが、1GB制約に合わせるため化合物データ1000万件(1.5GB)を4分割してPL/CUDA関数に与えている。しかし、予め化合物データがGPU側にロードされていれば、呼び出しのたびに2次元配列をセットアップする必要はなくなり、大幅に統計解析処理の効率が上がるだろう。

ここ最近の一連の開発で学んだのは、統計解析・機械学習のワークロードにおいては何がしかの前処理は必ず必要になるため、GPU RAMと関連付けられたForeign Tableに対して SELECT ... INTO fdw_table FROM transaction_table するというのは、どうやらユーザの負担を大きく増やすという性質のものではなさそうだという点。
生データからサマリを作ったり、前処理を行うというのはPG-Stromの得意とするところなので、これをPL/CUDAと組み合わせるための良い閃きが得られたと思っている。

[idea] 機械学習でのSSD-to-GPUダイレクトの利用

ポスターの発表に関連して、『これは機械学習のデータロードに使えませんかね?』と、以前にSSD-to-GPUダイレクトをブログで書いた時に頂いたのと全く同じ反応を、しかも複数の方から頂いた。
時間が取れないこともありペンディングになっているが、それほど難しい話ではないので、もう一度真面目に検討してみないと…。

[idea] CUDA MPSの利用

Tesla V100とCUDA9.0の発表に関連して、MPSの機能拡張によって、SMに空きがあるときはどんどんGPU kernelを放り込んでいけるようになるという説明があった。
現状、PostgreSQL v9.6に対応したバージョンはMPSと同じような機能を実現するバックグラウンドワーカを使っているが、こういった機能強化を取り込むためにも、できるだけ標準的なソフトウェアスタックの上に構築したいという思いはある。
現状、MPSだとストリームにコールバックを突っ込む事ができない。ここを何とかカバーする方法があれば、、、、という事で要検討。

*1:投票率は知らないがw

*2:なお、話しかけた相手が日本人だと若干気まずい。

*3:自業自得である

GTC2017 - Tesla V100 と CUDA9.0

今年のGTCはジェンスン・ファンCEOの基調講演が3日目に設定されており、そこでVolta世代の新製品 Tesla V100 と、それを搭載するDGX-1などサーバ製品とGPUクラウドが発表された。
4日間の日程のうち、情報の解禁が3日目の正午なので、その後は慌ただしくVoltaのアーキテクチャや、関連するCUDA9新機能のセッションに参加。
聞くべきセッションが被っちゃったりもしたので、できれば基調講演は初日の方が日程に余裕はあったかなぁ…。

Tesla V100について

5120CUDAコア、16GB(HBM2; 900GB/s)のメモリを搭載するまさに化け物プロセッサ。

バイス全体で80個のSMを搭載する。SMあたりの構成は64CUDA CoreなのでPascalと同等だが、目新しい基軸としてTensor Coreという新しい計算ユニットが追加されている(SMあたり8個 = 全体で640個)。また、L1 Cacheが128KB x 80 = 10MB、L2 Cacheが6MB搭載されている。

これは4x4行列同士のFMAD演算(←と言っていいの?)を実行する専用のH/Wユニットで、従来であれば、16個のCUDAコアが協調して掛け算/足し算を行うべき箇所を命令一発で実行するというもの。

深層学習のアルゴリズムに関して詳しくないんだけども、ハマる状況であれば綺麗に高速化できるんであろうし、実際、TensorCore機能を使った機械学習ワークロードの実行速度が、P100比で9.3倍に向上したという事はアナウンスされている。

また、6本のリンクを全てフル帯域で使った場合の理論帯域だけども、NVLinkで300GB/sのスループットがあり、NVLink接続のリモートGPUに対するcache coherencyやatomic operationが可能だとのこと。

この辺の発表を聞いていての印象は、15年-20年ほど前からx86 CPUが辿った道のりに似てるなぁ~と。

Intel CPUの場合、マルチメディア処理命令としてMMX/AVXという拡張命令が加わったが、それを自分のプログラムから直接使うという人はあまり多くなく、実態としてはglibcなどのライブラリが内部的にこれらの機能を使うようになったために、それを通じて、ハマるケース限定でハードウェア専用命令を使うようになってきている。
TensorCoreの場合も、cuDNNやcuBLASなどのライブラリ群が内部的にこの機能を使うとのことなので、ユーザプログラムが行列計算や深層学習にこれらのライブラリ群を使っていれば、知らず知らずのうちに恩恵を受けることになる。というか、ラップしてくれないとちょっと使える状況が特殊すぎて辛い。

もう一つ。NVLinkの帯域が向上し、しかもUnified Memoryでアドレス空間の問題が解決し、cache coherencyやatomic operationに対応するとなれば、他のGPUバイス上のメモリをローカルメモリに近い感覚で使うことができるようになるという事だろう。
そうすると、かつてLinuxがNUMA環境への対応を行ったように、メモリの物理的な距離を意識しなければならない状況が生まれつつあるように思える。例えば Tesla V100 は 16GB(HBM2) のメモリを積んでいるが、アプリケーションが20GBのデータをロードして2台のGPUでぶん回したい場合、データ量は明らかに1台のGPUに収まらないので2台に分散ロードする必要がある。
この時、物理GPUページの割り当てが『前半10GBはGPU0、後半10GBはGPU1』みたいなざっくりで上手く動作するのだろうか?もしかしたら、メモリアクセスの局所性が期待できるタイプのワークロードなのか? それとも、全体を舐め回すのでインタリーブ型の物理データ配置を行うべきなのか、それほど簡単な問いではないように思える。


SMの構成は図の通りであるが、L1 Cacheの扱いが少し変わっている。
PascalではL1 Cacheが24KBだったのが、128KBあるL1/Shared Memory共通の領域のうち、共有メモリが使用していない領域をL1 Cacheとして利用できるわけである。
Maxwell以前にL1/Shared Memoryの大きさを16KB/48KB、32KB/32KB、48KB/16KBの中から選んでいたのに通ずるが、もう少しフレキシブルに"早いメモリ"を使ってくれているようである。
ただ、GPUのL1キャッシュの場合(少なくともPascalまでは)、キャッシュの下に転がっているGlobal Memoryの内容がGPU kernel実行中に少なくとも変更されない/Read-Onlyであるという制約があったが、これはおそらくVoltaでも変わらないものと思われる(何も言及無かったし)

また、Pascal以前のGPUと大きく変わってくる点として、プログラムカウンタ(PC)がwarp単位ではなくthread単位になるという違いがある。
従来は、if (...) { ... } else { ... } ブロックを実行するときに、warp中の一部スレッドだけが条件が真、他のスレッドが偽である場合には両方のブロックを逐次的に実行することになるため、他方のスレッド群が遊んでしまい、それ故にGPUは条件分岐に弱いと言われてきたが、各スレッドがPCを持つという事はこうした弱点に手を打ってきたという事でもある。
(とは言え、TrueブロックとFalseブロックの分量があまりにアンバランスだとアカンでしょうが)

5月13日追記:
FBでNVIDIA-JPの成瀬さんから教えていただきました。
『プログラムカウンタ(PC)がwarp単位ではなくthread単位になる』というのは必ずしも間違ってはいないのですが、CPUスレッドのように各CUDAコアが全く独立な命令を実行できるという事ではないようです。

より正確に述べると、PC(ここでは次に実行する命令のアドレスとしましょう)は、Voltaではスレッド毎に持ちます。ただし、同じwarp内のスレッドは、PCが同じであるスレッドだけが、同じクロックで実行されます。warp内で、別PCの命令が実行されることは、これまで通り、ありません。

つまり、if (...) {} else {} ブロックで分岐が発生する場合、従来、trueブロックを実行中のCUDAコアがある時はfalseブロックに分岐したCUDAコアは寝ているしかなかった。Voltaにおいては、例えばDRAMアクセスなどでメモリストールが発生しておりtrueブロックを実行中のCUDAコアが寝ている状況では、その間を利用してfalseブロックに分岐したCUDAコアも先に実行を進めることができる、という事のようです。

他にニッチなところとしては、SMが80個にも増えてアプリケーションによっては全てのSMを使いきれないという事を想定しているのか、MPS(Multi Process Service)を使用した時に複数のアプリケーションからのGPU kernel起動要求を調停して、GPU資源を切り分けるという使い方ができますよーとの事。ただ、MPSはcuStreamAddCallback()のようにコールバック関数を呼び出す事ができなかったハズなので、非同期処理が中心のPG-Stromとは今一つ相性がなぁ…。

CUDA 9について

Voltaアーキテクチャ向けのTensorCoreを使うための機能や、cuBLASなどライブラリのTensorCore対応が行われるらしい。

ただ、最も革命的と言ってよい変化は、今までスレッド間の同期は __syncthreads() を使ったSM内の同期のみであったのが、SMを越えて複数のグリッド間で/複数のGPUバイスで並行動作するカーネル間で同期を取ることが可能になったというもの。

GPU kernelを起動するときは、必要なスレッド数がN個としてN=(ブロックサイズ)×(グリッドサイズ)
という関係になるようスレッド数を与えて起動する。ブロックサイズの上限は 1024 であり、GPU kernelの実行時にはブロックサイズに依らず、特定のブロックは特定のSMで実行される。
(グリッドサイズ > デバイスのSM数 な場合は、ブロックを順番に実行する)

で、大雑把に言えば、従来はH/Wに由来するブロックを単位とする同期の他には、GPU kernel全体の実行完了を待つという同期方法しか取れなかったが、CUDA 9で導入される co-operative group という概念を使えば、H/W的な構造と論理的な処理の同期の単位を切り分けることができる。

これがなぜ嬉しいかと言うと、例えばBitonic-Sortingを実行するときに、要素数が1個のSMで処理するには辛い数(1万個とか)になってくると、ステップ事に一度GPU kernelを完了し、そこで同期した上で改めて次のステップを実行するという処理が必要になる。
これは結構しんどい処理で、PG-StromでGPUでのORDER BY処理の実装を断念した要因の一つでもある。
が、SMを越えた単位で処理を同期できるとなれば、ステップ事に同期ポイントを設定してやればよいだけである。

例えば、以下のような感じで処理を記述する。

__global__ void
particleSim(Particle *p, int N)
{
  grid_group g = this_grid();

  // 1. Jobs prior to sync
  for (i = g.thread_rank(); i < N; i += g.size())
    integrate(p[i]);

  g.sync() // 2. Sync whole grid

  // 3. Jobs post sync
  for (i = g.thread_rank(); i < N; i += g.size())
    collide(p[i], p, N);
}

grid_groupというのはco-operative groupを表現するためのクラスで、this_grid()はこのGPU kernelを起動した時の全てのスレッドを含むオブジェクト。
つまり、g.sync()は複数のSMに跨った全てのスレッドで同期を取ることを意味している。

他にも、従来同様に現在のブロックに含まれるスレッドを表現する thread_group や、逆に複数のデバイスで同時に実行中の各GPU kernelに含まれるスレッドを表現する multi_grid_group が定義されている。
なお、SM間/デバイス間での同期機構を使う場合には、GPU kernelを起動する時に用いるAPIが従来のものとは少し異なっており、Runtime APIでは以下のAPIを用いてGPU kernelをキックしなければならない。

  • grid_groupを使う場合
    • cudaLaunchCooperativeKernel()
  • multi_grid_groupを使う場合
    • cudaLaunchCooperativeKernelMultiDevice()

で、更に面白いことに、このco-operative groupというのは更に細かい単位にパーティション化する事ができる*1

こんな調子で、thread_group gに含まれるスレッドを tiled_partition() を介して4つのパーティションに分割する事ができるらしい。
なので、真ん中の tile.sync() ではパーティション毎の同期という事になる。

  thread_group g = this_thread_label()
  thread_gruup tile = tiled_partition(g, 4);
        :
  tile.sync();
        :

現状、CUDA 9では等分割のパーティショニングのみだが、将来的には labeled partition という形にコンセプトが拡張されるとの事。
"labeled" というのは、以下のように partition() 関数に与えた整数値によってそのスレッドが属するパーティションを決めるというもの。

  int label = foo() % 4;
  thread_group g = partition(this_thread_block(), label);

何が嬉しいかというと、GROUP BY/集約演算をGPUで実装するときに、GROUP BYで指定したカラムの値が同一のスレッド間でのみ同期を取り、縮約演算を行うというロジックが可能になること。これは、現状atomic演算を用いて実装せざるを得ないことを考えると非常に大きなアドバンテージになる。

CUDA 9.0の登場が待ち遠しいところである。

余談

SM間で同期を取るとき、GPU kernelのブロック数がデバイスのSM数よりも多い場合、ハードウェア的にはSMを明け渡さないといけないはずなので、そのタイミングで共有メモリやレジスタをグローバルメモリに退避する(= そんなに軽い処理じゃないのでは?)と思い、セッション後に発表者を捕まえて質問してみた。
曰く、Pascal世代で(H/W的には)対応したプリエンプションの仕組みを使って実装されており、レジスタや共有メモリの使用量がH/Wリソースを超える場合には確かにそうなる。なので、GPU kernelを起動するときのパラメータチューニングが必要で、SMあたり {(ブロック数)/(デバイスのSM数)} のブロックを保持できる程度のレジスタ/共有メモリの消費に抑えておく必要がある。
一応、その辺は Occupancy を計算するAPIを使って効率的なブロックサイズ/グリッドサイズを求めてほしいとの事。

お願い

S7285 - Unified Memory on the Volta architecture』のセッションが、CUDA9のセッションと被ってて聞けなかった!誰か内容教えて!!

NVIDIA TechBlogにも既に発表内容を踏まえた記事が掲載されています。

(*) 本文中の画像は全て上記ブログからの引用です。

*1:が、セッションの中ではthread_groupの例しか示されてなかったので、grid_groupやmulti_grid_groupでもできるかは要確認

Posters flying on GTC2017

I'm on GTC(GPU Technology Conference) this year again.

I didn't plan to trip because of travel budget, so only posters are submitted and approved.
(Posters are shown regardless of the attendance of the presenter.)

But I got an e-mail below...

Hi Kohei,

After extensive review from multiple committees and nearly 140 posters, your poster has been chosen as a Top 5 Finalist. Congratulations!!

Eventually, I arranged a trip in hurry, then now in San Jose.

The top-5 posters are shown at most conspicuous place of the convention center.

Participants can vote one of their favor poster, then the best poster this year is chosen through the democratic process. :-)

I and other top-20 presenters had a short talk on evening of the 1st day of the conference.
4-minutes are given to each presenter, and we introduce their research in this short time-slot.
(It is shorter than usual lightning talks!)

Our challenge tries to de-define the role of GPU. People usually considers GPU is a powerful tool to tackle computing intensive workloads. It is right, and widespread of workloads utilizes GPUs well like deep learning, simulations and so on.
On the other hands, I/O is one major workloads for database; including data analytics which try to handle massive amount of data.


We designed and implemented a Linux kernel module that intermediates peer-to-peer data transfer between NVMe-SSD and GPU, on top of GPUDirect RDMA feature. It allows to load PostgreSQL's data blocks on NVMe-SSDs to GPU's RAM directly.
Once data blocks get loaded to GPU's RAM, we already have SQL execution engine on GPU devices.
It means we can process earlier stage of SQL execution prior to data loading onto the system main memory. The earlier stage includes WHERE-clause, JOIN and GROUP BY. They likely work to reduce amount of data.
For example, when an WHERE-clause filters out 90% of rows (it is not a strange assumption in reporting queries) on scan of 100GB tables, we can say 90GB of I/O bandwidth and dist cache to be wipe out are consumed by junk data.
Once data blocks are loaded to GPU, it can filter out unnecessary rows (and processes JOIN and GROUP BY), then write back the results of this pre-process to host memory. The data blocks written back contains all valid data because unnecessary data are already dropped off.


In the results, it performs nearly wired speed of NVMe-SSD.
In case of single SSD benchmark, query execution throughput by vanilla PostgreSQL is around 700MB/s. On the other hands, SSD-to-GPU Direct SQL-execution runs same query with 2.2GB/s; which is almost identical with catalog spec of Intel SSD 750 (SeqRead: 2200MB/s).
In case of multi SSDs cases, it also increases the performance, however, deeper investigation *1 and improvement for md0-raid support are needed.


We intend to position this technology to make summary and pre-process of raw data/logs generated by sensor devices, mobile devices and so on. Once logs get summarized, we can apply statistical-analytics / machine-learning algorithms.
This type of workloads are advantaged with PL/CUDA which is one other our feature to utilize GPUs in-database analytics.
Eventually, GPU-accelerated PostgreSQL will be capable to handle entire life-cycle of data utilization (generation - gathering - summarizing/pre-process - analytics/machine-learning - visualization).
It makes system configuration much simple, reduce cost for production and operation, and deliver the power of analytics to wider and wider people.

My poster is below (enlarge by click):

If you are favored with our development, don't forget to vote.

Enjoy your GTC2017!

(Comment to the above photo by my wife: it looks like a zombie.)

*1:At that time, we measured with Teska K80; that is a bit old.