libSGMがOpenCVにマージされました

はじめに

初めまして、アルバイトの大塚です。adaskitという社内の自動運転関連のオープンソースプロジェクトに微力ながら携わっており、弊社が公開している視差計算ライブラリlibSGMの開発を行っていました。

libSGMが機能追加などで成熟してきたことを受け、adaskitチーム内でOpenCVに取り込んでもらうことができないか検討するようになり、調査と移植作業の末にPull Requestを送りました。

その結果Pull Requestが承認され、version 4.5.1にて利用可能になりました。

動作確認

OpenCVでlibSGMが利用できるようになったことを確認してみましょう。opencv/opencv、opencv/opencv_contribから4系の最新版を入手し、ビルド・インストールします。libSGMはcudastereo moduleに追加されているため、BUILD_opencv_cudastereoをONにします。

OpenCVはPython bindingsを備えており、開発者は関数などにマクロを追加するだけでPythonから利用できるようになります。今回は簡単な動作確認としてPythonスクリプトを用意しました。

OpenCVのperfテストで利用されているaloeの画像を入力に与えてみます。

期待した結果が得られました。libSGMの開発に携わった一個人としては、OpenCVの機能にlibSGM（の移植）が入り、動かせることに感動します。

速度は移植前のlibSGMと比べてどうでしょうか。libSGMのbenchmark sampleを用いて計測しました。計測条件と結果は以下の通りです。

移植前と遜色ないパフォーマンスを発揮できています。若干の差は移植するにあたり、既存のクラス・構造体を用いたことに依るものではないかと考えています。

Pull Requestを送るにあたって

ここでは今回苦労した・体験した点を書いていこうと思います。

OpenCVは開発者向けに以下のドキュメントを用意しています。

• https://github.com/opencv/opencv/wiki/How_to_contribute
• https://github.com/opencv/opencv/wiki/Coding_Style_Guide

これらを読む他、既存のPull Requestのレビューでどのような点が指摘されているか確認しました。

以上を元にlibSGMをOpenCVスタイルに移植し、動作確認しました。

どのmoduleに移植するか

CUDAでステレオマッチングを行う、という機能に合ったmoduleがopencv_contribのcudastereoに当たるため、こちらで作業しました。

CUDA 8.0, Compute Capability 2.0のサポート

移植作業当時、OpenCVのCMakeビルド設定CUDA_ARCH_BINに2.0が含まれ、Warp Shuffle関数がないとエラーが発生しました。libSGMではCompute Capability 3.5以上をサポートしていますが、OpenCVに取り入れてもらうため共有メモリを使った実装をラッパーに追加し、テストしました。

// https://github.com/opencv/opencv_contrib/blob/4.5.1/modules/cudastereo/src/cuda/stereosgm.cu#L100-L119
template <unsigned int WARPS_PER_BLOCK, typename T>
__device__ __forceinline__ static T shfl_xor(T var, int laneMask, int width = cudev::WARP_SIZE, uint32_t mask = 0xFFFFFFFFU)
{
#if __CUDA_ARCH__ >= 300
#if CUDA_VERSION >= 9000
    return __shfl_xor_sync(mask, var, laneMask, width);
#else
    return __shfl_xor(var, laneMask, width);
#endif // CUDA_VERSION
#else
    static __shared__ T smem[WARPS_PER_BLOCK][cudev::WARP_SIZE];
    smem[cudev::Warp::warpId()][cudev::Warp::laneId()] = var;
    T ret = var;
    if (((cudev::Warp::laneId() % width) ^ laneMask) < width)
    {
        ret = smem[cudev::Warp::warpId()][cudev::Warp::laneId() ^ laneMask];
    }
    return ret;
#endif // __CUDA_ARCH__
}

この他、同様の理由で__ldg関数のラッパーを作成しました。

テストの作成

OpenCVに新しい機能を追加する場合、テストが必要です。libSGMには既存のテストがありますが、これに加えてOpenCVのStereoSGBMで用いられているテストを追加しました。

このテストの仕組みは簡単で、まずパラメータファイルを用意してテストを実行するとGround Truthとの比較結果が保存されます。次回実行時にはその比較結果と相違ないか確認が行わる、というものです。この比較結果はopencv/opencv_extraをForkしたリポジトリにcommitしておき、opencv_contribと同時にPull Requestを送ります。

終盤近くになってCIを回すようにしました。CI用のイメージはopencv-infrastructure /opencv-worker-configリポジトリにあるDockerfileを参考にしました。

Pull Request作成後の流れ

文面を整え、緊張の中いざPull Requestを作成した後は以下のように事が進みました。

• Pull RequestがOpenされてから3分後、labelが貼られる
• OpenCV buildbot用の記述が書き込まれる
  • このため、こちらで記述する必要はない

force_builders=Custom
buildworker:Custom=linux-4
build_image:Custom=ubuntu-cuda:16.04

• 1週間後、レビューが開始される
• 約2時間後、approvedがもらえる
• 翌日opencv-pushbotによってmergeされる

問題なくmergeされ、とても安心しました。

おわりに

libSGMがOpenCVに移植されたことで、より多くの方に触れていただける機会が増えることを願っています。

アルバイトでは日々貴重な体験をさせていただいていますが、「OpenCVに大きなサイズのPull Requestを送る」という経験ができてよかったです。

事前のレビューや相談に対応して下さった高木章洋さん、OpenCVに関する知識を提供して下さった吉村康弘さん、Pull Requestに対応して下さったalalekさんにこの場を借りて感謝申し上げます。

なお、今回のPull Requestでは移植作業全般を担当したことを理由に、commitのauthorを大塚とさせていただきました。
libSGMの実装が含まれるcudastereo/src/cuda/stereosgm.cuにはauthorをadaskit Teamと記しております。

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。