このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。
フィックスターズは今年も ACM ICPC つくば大会 にスポンサーとして協賛しております。今年はコンテスト後の懇親会でクイズを出題しているのですが、その問題準備の際に作ったはいいものの結局使われなかったコードが大量に出てきてしまったため、それらをまとめて簡単な解説とともに公開いたします。
全点対間最短路問題はN個の頂点を持つ有向グラフにおける任意の2点間の最短路を求める問題です。密なグラフの全点対間最短路を求めるアルゴリズムとしてワーシャル-フロイド法が広く知られており、以下のような非常にシンプルなコードで求めることができます。
// 入力: A[i][j] = 頂点iと頂点jを結ぶ辺の重み
// 出力: A[i][j] = 頂点iから頂点jまでの最短路の重みの総和
for(int k = 0; k < N; ++k)
for(int i = 0; i < N; ++i)
for(int j = 0; j < N; ++j)
A[i][j] = min(A[i][j], A[i][k] + A[k][j]);
上に示したワーシャルフロイド法のアルゴリズムは i, j ループを並列化した以下のようなコードでも正常に動作します。
for(int k = 0; k < N; ++k)
#pragma omp parallel
for(int i = 0; i < N; ++i)
for(int j = 0; j < N; ++j)
A[i][j] = min(A[i][j], A[i][k] + A[k][j]);
……正常に動作して正しい結果は返してくれるのですが、演算資源より先にメモリ帯域を使い切ってしまうため少ないコア数でもあっという間に性能が頭打ちしてしまいます。
ICPCの際に出題したクイズには、上記のコードと同様の方針で並列化したCUDAコードを並べてどれくらいの性能差か予測するという問題を出したのですが、どちらのコードもまだかなりの性能向上の余地があります。
愚直な i, j ループの並列化がスケールしないのはメモリ律速になってしまっているのが原因です。メモリアクセスのコストを減らすためには、キャッシュやレジスタに収まるサイズに問題を分割し、その範囲の演算をまとめて行うことでメインメモリへのアクセスを減らすブロッキングがよく用いられます。ワーシャルフロイド法の高速化においてもブロッキングは有効で、分割の方法もいくつか提案されているのですが今回は再帰的に分割していく方法 [1] を使用しました。
ある程度以上のパフォーマンスチューニングは非常に労力がかかるのですが、問題の一部を既存の研究されつくしたカーネルに任せることで比較的容易に非常に効率の良いアルゴリズムを実装できることがあります。再帰的に分割するワーシャル-フロイド法でもそのようなアプローチをとることができ、処理の大部分をトロピカル半環(加算と乗算の代わりにminと加算を用いる)の行列積に帰着させることができます [2]。
行列積と聞くと理論ピークに近い性能が出せそうな気がしてきますね!
行列積といえばプロセッサの浮動小数点数演算回路をほぼフルに使う処理のひとつとして有名ですが、今回扱う行列積は add, mul の代わりに min, add を使用しているため、FMA が利用できないなど通常のそれとまったく同じとはいかない箇所があります。
ここでは、今回使用したプロセッサでトロピカル半環の行列積を求める際の理論ピーク性能を求めておきます。
今回使用したプロセッサ (Core i7-4930K) では 3.4 [GHz] × 6 [コア] × 8 [ベクトル長] で 163.2 [GFLOPS] が限界となります(TurboBoostを無視した場合)。通常の行列積ではさらにスーパースカラが効いて2倍の 326.4 [GFLOPS] となるのですが、IvyBridge では add と min が両方とも同じポートを使用するため、スーパースカラの効果が得られず限界値が低くなってしまいます。
今回使用した GPU (GeForce GTX 980) では 1.126 [GHz] × 16 [SMM] × 64 [SMMあたりのminのスループット] × 2 [スーパースカラ] で 2306 [GFLOPS] が限界となります(Boostを無視した場合)。CPU の時とは対照的に、こちらはスーパースカラが効いて min と add が並列に処理できているようでした。一方、通常の行列積では FMA を用いることでさらに倍の演算を行うことができるため、通常の行列積との比較だと GPU でも半分がピークとなります。
cuBLASなどのライブラリではあまりトロピカル半環の行列積はサポートされていないため、今回はキャッシュやレジスタでブロッキングする程度のチューニングを施したカーネルを自前で用意しました。
# 行列積のチューニングはまじめに書くとそれだけで記事1本以上のボリュームになってしまうので割愛します……
これらのコードとビルドスクリプトなどはBitBucketで公開しています。
アルゴリズム | 処理時間 [ms] | GFLOPS | 愚直な実装に対する性能向上 |
---|---|---|---|
CPU: 愚直な並列化 | 257656 | 4.26 | |
CPU: 再帰的分割 | 28405 | 38.70 | ×9.07 |
CPU: 再帰的分割+行列積 | 10213 | 107.65 | ×25.22 |
GPU: 愚直な並列化 | 25460 | 43.18 | |
GPU: 再帰的分割+行列積 | 604 | 1817.70 | ×42.09 |
アルゴリズム | 処理時間 [ms] | GFLOPS |
---|---|---|
CPU: 再帰的分割+行列積 | 84806 | 103.71 |
GPU: 再帰的分割+行列積 | 4685 | 1877.21 |
それぞれのプロセッサ用の愚直な実装と比べてみると、CPUでは約25倍、GPUでは約42倍の性能向上となりました。
また、CPUでは限界の65%程度、GPUでは限界の80%程度の性能を出すことができています。
なんとなくではありますがCPU版はもう少し詰められそう、GPU版はCUDA Cでここから先はきつくなりそうといった気がします。
ワーシャル-フロイド法を通して、同じハードウェア・同じ計算量であっても性能に大きな差が出る例を示しました。
ICPCのようなコンテストではあまり定数倍のチューニングは重要視されませんが、このような方法でのアルゴリズムの高速化にも興味を持っていただければ幸いです。
keisuke.kimura in Livox Mid-360をROS1/ROS2で動かしてみた
Sorry for the delay in replying. I have done SLAM (FAST_LIO) with Livox MID360, but for various reasons I have not be...
Miya in ウエハースケールエンジン向けSimulated Annealingを複数タイルによる並列化で実装しました
作成されたプロファイラがとても良さそうです :) ぜひ詳細を書いていただきたいです!...
Deivaprakash in Livox Mid-360をROS1/ROS2で動かしてみた
Hey guys myself deiva from India currently i am working in this Livox MID360 and eager to knwo whether you have done the...
岩崎システム設計 岩崎 満 in Alveo U50で10G Ethernetを試してみる
仕事の都合で、検索を行い、御社サイトにたどりつきました。 内容は大変参考になりま...
Prabuddhi Wariyapperuma in Livox Mid-360をROS1/ROS2で動かしてみた
This issue was sorted....