skymatix Developers Blog

株式会社スカイマティクスの開発チームによるDevelopers Blogです。

CUDA プログラムの性能分析と高速化

Researcher の尾川です。今回は、弊社の MVS 機能の高速化についてとりあげます。

MVS は、画像と SfM の結果(画像の撮影位置、特徴点における画素座標と3次元座標の対応関係)をもとに3次元点群を作成するプログラムです。この手法は大別すると4種類ありますが、弊社では2枚の画像の比較によって奥行きを推定する PatchMatch を採用しています。とても単調な作業を大量に繰り返すので、GPGPU に適した処理です。

高速化と言われたら最初にすること、皆はわかるかな? そうだね、測定だね。 ということで Nsight Systems を使って nsys profile ./a.out で調べてみると、実は CUDA を使っていない時間がかなり多いことが判明しました。GPU のあるマシンで十分に使えていないのはもったいないですね。

高速化前(全体、10秒)

NVTX によるマークアップなどで細かく調べていった結果、1回の CUDA 実行(カーネル計算)の前後に関連するすべてのファイルを読み書きしているのが一番の原因であることがわかりました。いわゆる IO バウンドです。CUDA 実行中は CPU は暇を持て余しているので、k 回目の実行中に k+1 回目の準備を行い、k-1 回目の片付けをすることで、ほぼ連続した CUDA 実行が可能になることが期待できます。

コードで書くと、

for (int i = 0; i < n; ++i) {
    before_execute(i);
    execute_kernel(i);
    after_execute(i);
}

を次のように組み替えるイメージです。

auto after_task = std::async(std::launch::deferred, [] {});
auto before_task = std::async(std::launch::deferred, [] { before_execute(0); });
for (int i = 0; i < n; ++i) {
    after_task.wait();
    before_task.wait();
    if (i + 1 < n) {
        before_task = std::async(std::launch::async, [i] { before_execute(i + 1); });
    }
    execute_kernel(i);
    after_task = std::async(std::launch::async, [i] { after_execute(i); });
}
after_task.wait();

これでかなり改善されました。さらに、数カ所に #pragma omp parallel for を入れたり、二重ループの添字の順を変更したり、処理の順序を入れ替えてキャッシュする(すべてキャッシュすると RAM が不足する)ことで IO 回数を減らしたりといった追加の工夫により、全体の処理時間は約150秒 → 約60秒、GPU 使用率は概算で 40% → 80% に改善されました。

高速化後(全体、10秒)

以上、CUDA プログラムと言っても .cu ファイル以外にもみるべき場所はありますよ、というお話でした。

なお、Nsight compute の印象からはカーネル計算そのものの高速化は最大 30% あるかどうかだと見込んでいます。

www.slideshare.net