特徴記述アルゴリズムHashSIFTのCUDA高速化

はじめに

初めまして、アルバイトの藤本です。特徴記述は、画像における特徴点を「記述」、すなわち特徴ベクトルという数値に変換する処理です。今回のプロジェクトにおいてはSfMという３次元復元アルゴリズムで使用することを想定しています。

数ある特徴記述のアルゴリズムの中でもHashSIFTは精度と計算コストのバランスに優れており、私はインターン生の時から、GPU環境におけるさらなる高速化に取り組んでいました。

ちなみにその成果はfixstars/cuda-efficient-featuresの一部として公開しています。cuda-efficient-featuresについてはCUDAによる局所特徴量計算の高速化とソースコード公開にて紹介がされているのでご覧いただけると幸いです。

今回はHashSIFTの概要とその高速化手法について紹介していきます。

HashSIFTの概要

HashSIFT^[1]は特徴記述において長く使われているSIFTの改良版です。以下の４つの手順からなります。

1. 特徴点のオリエンテーションに対する周辺領域、32×32画素を4×4ブロックに分割
2. ブロックごとで周辺領域の輝度勾配を８方向に分類し、勾配方向ヒストグラムを作成
3. 4×4×8＝128次元の特徴ベクトルを記述
4. あらかじめ最適化された行列によって特徴ベクトルを256または512次元に拡張し、2値化

SIFTでの特徴ベクトルの記述方法(^[2]より引用)

1~3の操作は上図にあるように、従来のSIFTと同じ操作です。

CUDAによる高速化

cv::cudaライブラリの使用

グレイスケール変換や行列演算など基本的な処理はすでにopenCVに高速化が実装されています。これを用いれば、CUDAの知識がなくとも簡単に高速化ができます。

特徴ベクトル作成処理の並列化

HashSIFTにおいて大部分を占める特徴ベクトル計算は、まず各特徴点ごとに並列化できます。また、1特徴点について32×32の周辺画素を処理しますが、これも画素ごとに並列化できます。

1つの特徴点に対するパッチ(周辺画素データ)、ヒストグラム、特徴ベクトルをshared memory上で共有するため、下図のように1特徴点に1ブロックを、1画素に1スレッドを割り当てることにしました。 (実際は、32×32画素に同数のスレッドを割り当てるとメモリや同期といった点でパフォーマンスが低下するので、一番効率が良くなるようにスレッドの次元は調節しました。)

各グリッド、スレッドに対する特徴点、画素の割り当て

リダクション命令における並列化

作成した特徴ベクトルは結果として返す前に正規化するのですが、この際に特徴ベクトルの2乗和を計算する必要があります。

ここでshared memoryに置かれた特徴ベクトルの要素を単に集計(リダクション)しようとすると、1つのメモリに対して大量のスレッドがアトミックなアクセスをする必要が出てきて非効率です。そこで

1. 特徴ベクトルの各要素を２乗しつつ、1つのwarp内にまとめる
2. インターリーブペア方式の2乗和計算を、warp shuffle命令を用いて実行する

これによって各特徴点に対する計算を効率化しました。

2乗和計算の並列化

cudaAPIの呼び出し回数削減

SfMは大量の画像を処理する前提なので、今回実装した特徴記述は繰り返し呼び出されることを想定しています。

このとき、呼び出しの度にデバイスメモリを確保/解放するのはオーバーヘッドとなり、できれば避けたいです。

そこで、それらのメモリを関数内のローカル変数からクラス変数にし、そのクラスのインスタンス自体は処理の最初から最後まで保持し続けるようにしました。これにより画像間でのcudaAPIの処理時間を大幅に減少させることができました。

しかしプロファイルを見るとcv::cuda::gemmという行列演算を実行した後に、まだcudaAPIが呼び出されてしまっています。そこでopenCVの実装を確認したところ、cublasのハンドル作成が原因でした。

cublasHandle_t handle;
cublasSafeCall( cublasCreate_v2(&handle) ); 

handleについてもクラス変数にしたところ完全に解消。良かったです。

ヒストグラム作成時のatomicAddの衝突回数削減

ここまで高速化を進めて、まだ処理が重い部分はヒストグラムの作成でした。

今回は64スレッドが4×4のヒストグラムに対して同時に書き込みを行っていくので、アトミック演算であるatomicAddが必要でした。これはreduction命令におけるwarp shuffleで言及した通り非効率です。

理想としてはatomicAddを完全になくしたかったのですが、そのためにはどうしてもスレッド数を減らす必要があり、結果として実行時間が長くなってしまいます。そこで、各スレッドがアクセスするヒストグラムの場所を分散させることで待機状態になる時間をなるべく減らすようにしました。

CPU実装におけるヒストグラム作成時のメモリアクセスパターン

最初は上図のように青色のスレッドが黄色の枠で囲まれたヒストグラム内の画素を上から下に走査していく実装でしたが、これだとヒストグラムの1セル内に16スレッドが衝突する可能性があります。そこで以下のようにスレッドを分散させてやることによって1セル内のスレッド数を4スレッドに下げることができ、結果として高速化に成功しました。

GPUに対して最適化したメモリアクセスパターン

実行結果比較

以上のアルゴリズムを実装した後、処理時間の結果をもとのCPU実装と比較しました。
計測環境は以下の通りです。

計測環境

Key	Value
CPU	Core i9-10900(2.80GHz) 10Core/20T
GPU	RTX 3060 Ti
入力画像	SceauxCastleに含まれる11枚の画像
特徴点数	10,000点に固定
計測方法	各画像について100回実行の平均処理時間を計測し、最後に11枚の平均をとる

特徴記述の処理時間比較 (単位は[msec])

実装＼ディスクリプタ	HashSIFT256	HashSIFT512
リファレンスCPU Single Thread	604.9	740.1
リファレンスCPU Multi Thread	174.6	315.6
cuda-efficient-features	0.89	0.99

CPUのマルチスレッド実装と比較しても200~300倍程度高速にすることができました。

さいごに

hashSIFTのアルゴリズムは並列化と非常に相性が良かったようで、特徴記述の処理時間をほぼ気にする必要がないところまで高速化することができたのではないかと思います。

3、4か月ほど実装に取り組ませてもらっていましたが、様々な高速化技法に触れることができた貴重な期間でした。

最後に、メンターとして日々ご指導いただいている高木章洋さんに感謝申し上げます。

参考文献

脚注

1. I. Suárez, J. M. Buenaposada and L. Baumela, “Revisiting Binary Local Image Description for Resource Limited Devices,” in IEEE Robotics and Automation Letters, vol. 6, no. 4, pp. 8317-8324, Oct. 2021

2. 藤吉弘亘、“Gradientベースの特徴量抽出-SIFTとHOG-”、情報処理学会研究報告、CVIM160、pp. 211-224、2007