コンピュータビジョンセミナーvol5～Multi-View StereoのCUDA高速化～（2024/8/7)

9.8K Views

August 13, 24

#cuda #機械学習 #deeplearning #深層学習 #visualslam #画像解析 #画像処理 #cuda高速化 #自動運転 #コンピュータビジョン #コンピュータビジョンシリーズ #sfm #3次元復元 #Multi-View Stereo #CUDA #PatchMatch

スライド概要

画像による3次元計測・モデリングのパイプラインにおいて、Multi View Stereo(MVS)は、SfMによって得られた疎な点群とカメラ姿勢をもとに、密な点群を生成する技術です。

PatchMatch MVSはデプスマップをランダムな状態から繰り返し改善するMVSの1手法であり、GPUによる並列化と相性が良いことが知られています。

今回は、PatchMatch MVSのアルゴリズムと、GPUによる並列化の手法を解説します。

＜コンピュータビジョンシリーズの過去資料＞
・vol.1　OpenCV活用（OpenCVでCUDAを活用するためのGpuMat解説）：
　　https://www.docswell.com/s/fixstars/ZRXQ72-20220805
・vol.2　視差計算ライブラリ libSGM のアルゴリズム解説と CUDA高速化：
　　https://www.docswell.com/s/fixstars/5ENE7J-20221220
・vol.3　視差計算を利用した障害物検出：
　　https://www.docswell.com/s/fixstars/ZQ8447-20231027
・vol.4　Visual SLAM / SfMで行われる局所特徴量計算のCUDA高速化
　　https://www.docswell.com/s/fixstars/ZM1DJ8-20240327

・フィックスターズの画像処理アルゴリズム開発支援：
　　https://www.fixstars.com/ja/services/image-processing

株式会社フィックスターズ

@fixstars

スライド一覧

フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。　　　　　　・開催セミナー一覧：https://www.fixstars.com/ja/seminar　　・技術ブログ：https://proc-cpuinfo.fixstars.com/

またはPlayer版

埋め込む »CMSなどでJSが使えない場合

（ダウンロード不可）

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

kaggle 画像処理機械学習深層学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 61.4K

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

arm neon

株式会社フィックスターズ 37K

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

kaggle 画像処理機械学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 28.7K

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

cuda gpu cuda高速化高速化シリーズ

株式会社フィックスターズ 27.8K

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

ros2 gpu ロボット自律走行 ros2シリーズ高速化シリーズ

株式会社フィックスターズ 27K

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

gpu cuda cuda高速化高速化シリーズ cuda高速化シリーズ

株式会社フィックスターズ 25.8K

各ページのテキスト

本講演の位置づけ ⚫ 弊社でサービス展開しているコンピュータビジョン領域での様々な技術情報を、コンピュータビジョンセミナーとして発信しています ⚫ Vol.1 OpenCVでCUDAを活用するためのGpuMat解説 ⚫ Vol.2 視差計算ライブラリ libSGM のアルゴリズム解説 ⚫ Vol.3 視差計算を利用した障害物検出 ⚫ Vol.4 Visual SLAM / SfM で行われる局所特徴量計算の CUDA高速化 ⚫ Vol.5 3次元復元アルゴリズム Multi-View Stereo の CUDA 高速化 ⚫ こんな方に向いています ○ 3次元復元アルゴリズムのトレンドを知りたい ○ 実用的なアルゴリズムの GPU 高速化事例を知りたい Copyright © Fixstars Group 3

発表者紹介冨田明彦高木章洋ソリューションカンパニー営業企画ソリューション第二事業部エグゼクティブエンジニア 2008年に入社。金融、医療業界において、ソフトウェア高速化業務に携わる。その後、新規事業企画、半導体業界の事業を担当し、現職。 2014年に新卒入社。自動運転向けの画像認識アルゴリズムの開発、高速化などに携わる。 OSS活動として、libSGMの開発にも携わる。 Copyright © Fixstars Group 4

フィックスターズの強みコンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団ハードウェアの知見アルゴリズム実装力各産業・研究分野の知見目的の製品に最適なハードウェアを見抜き、その性能をフル活用するソフトウェアを開発します。ハードウェアの特徴と製品要求仕様に合わせて、アルゴリズムを改良して高速化を実現します。開発したい製品に使える技術を見抜き、実際に動作する実装までトータルにサポートします。 Copyright © Fixstars Group 6

サービス概要お客様専任のエンジニアが直接ヒアリングを行い、高速化を実現するために乗り越えるべき課題や問題を明確にしていきます。高速化のワークフローお客様オリジナルソースコードのご提供高速化したコードコンサルティング高速化サポート先行技術調査アルゴリズムの改良・開発レポートやコードへのQ&A 性能評価・ボトルネックの特定ハードウェアへの最適化実製品への組込み支援レポート作成 Copyright © Fixstars Group 7

サービス提供分野半導体産業機器金融自動車 ● NAND型フラッシュメモリ向けファームウェア開発 ● 次世代AIチップの開発環境基盤生命科学 ● Smart Factory実現への支援 ● マシンビジョンシステムの高速化 ● 自動運転の高性能化、実用化 ● ゲノム解析の高速化 ● 次世代パーソナルモビリティの研究開発 ● 医用画像処理の高速化 Copyright © Fixstars Group ● デリバティブシステムの高速化 ● HFT(アルゴリズムトレード)の高速化 ● AI画像診断システムの研究開発 8

10.

11.

画像処理アルゴリズム開発高速な画像処理需要に対して、経験豊富なエンジニアが責任を持って製品開発をご支援します。お客様の課題ご支援内容高度な画像処理や深層学習等のアルゴリズムを開発できる人材が社内に限られているアルゴリズム調査・改変課題に合ったアルゴリズム・実装手法を調査製品実装に向けて適切な改変を実施機能要件は満たせそうだが、ターゲット機器上で性能要件までクリアできるか不安深層学習ネットワーク精度の改善様々な手法を駆使して深層学習ネットワークの精度を改善製品化に結びつくような研究ができていない論文調査・改善活動論文調査から最先端の手法の探索性能向上に向けた改善活動を継続 Copyright © Fixstars Group 10

12.

GPU向け高速化高性能なGPUの本来の性能を十分に引き出し、ソフトウェアの高速化を実現します。お客様の課題ご支援内容 GPUで計算してみたが期待した性能が出ない GPU高速化に関するコンサルティング GPU/CPUを組み合わせた全体として最適な設 CPU・GPU混在環境でのシステム設計計がしたいアルゴリズムのGPU向け移植原価を維持したまま機能を追加するため、もう少し処理を速くしたい GPUプログラム高速化品質確保のため、精度を上げたく演算量は増えるが性能は維持したい Copyright © Fixstars Group 継続的な精度向上 11

13.

14.

Multi-View Stereoとは ⚫ 複数枚の画像から被写体の3次元モデルを復元する技術 ○ 特にカメラ姿勢の推定までをSfM、それ以降をMVSとして扱うことが多い Structure from Motion(SfM) • 入力画像群から、元のカメラ姿勢を推定 Multi View Stereo(MVS) • 入力画像群と、対応するカメラ姿勢をもとに、被写体の3次元構造を復元 • 必要に応じて被写体の質感も再現 SfM/MVSのパイプライン (図は[1]より引用) Copyright © Fixstars Group 13

15.

狭義のMulti-View Stereo ⚫ MVSのパイプライン前段では密な点群生成を行う処理があり、これを指してMVSと呼ぶこともある狭義のMVS 入力画像群密な点群生成 Feature Extraction メッシュ化 Feature Matching SfM MVS レンダリング Relateve Poses 精緻な3Dモデル Reconstruction Bundle Adjustment カメラ姿勢+疎な点群 SfM/MVSのパイプライン詳細 Copyright © Fixstars Group 14

16.

SfM/MVSのライブラリとPatchMatch ⚫ SfMやMVS向けに様々なライブラリが公開されている ⚫ MVSライブラリの多くが、密な点群生成においてPatchMatchを採用 ⚫ 本セミナーでは、PatchMatch MVSについて解説 SfM/MVSのライブラリとサポート機能ライブラリ URL SfM openMVG https://github.com/openMVG/openMVG openMVS https://github.com/cdcseacave/openMVS COLMAP https://github.com/colmap/colmap OpenSfM https://github.com/mapillary/OpenSfM geogram https://github.com/BrunoLevy/geogram MVS (dense | mesh) 〇〇 (PM,SGM) 〇〇〇 (PM) 〇〇〇 (PM) 〇〇 PM：PatchMatch SGM：Semi-Global Matching Copyright © Fixstars Group 15

17.

【参考】深層学習系のアプローチ ⚫ 深層学習ベースのMVSも近年研究されている ⚫ 特徴抽出の能力に優れており、精度面では古典手法を上回るとされる ○ 一方で大量のGPUメモリを必要とし、高解像度データの扱いが困難 ETH3D：高解像度データセット深層学習系 PatchMatchなどの古典系 MVSアルゴリズムと性能分布 (図はAPD[6]より引用) Copyright © Fixstars Group 16

18.

【参考】3D Gaussian Splatting ⚫ SfMの結果から、精緻かつ高速なレンダリングを行う技術 ○ https://repo-sam.inria.fr/fungraph/3d-gaussian-splatting/ ○ MVSと似ているが、点群やメッシュを経由せずにレンダリングする点が異なる ⚫ 近年、3D Gaussian Splattingを3D復元に応用する手法も登場 ○ SuGaR[10]：https://anttwo.github.io/sugar/ A B C SuGaRのデモ(図は[10]より引用) A,B,Cの各シーンを3D復元し、B,CをAに合成 (ロボットのポーズも編集) Copyright © Fixstars Group 17

19.

20.

21.

デプスマップ推定の原理 1. 参照視点の注目画素における点を、その画素のデプス値を用いて逆投影 2. 1.の点を近傍視点に投影し、画像上でマッチングコストを計算 3. コスト最小化により最適なデプス値を推定 • • • 参照視点：デプスマップ推定の対象である視点近傍視点：参照視点と一定の視野を共有する視点マッチングコスト：画素間の見た目の相違度注目画素逆投影参照視点のデプスマップ投影視点間の座標変換 Copyright © Fixstars Group マッチングコスト計算 20

22.

PatchMatch MVS ⚫ PatchMatch：ランダム性を利用した、画像マッチング(対応点探索)手法 ⚫ MVSでは次のような処理から構成されるステップ1：ランダム初期化 • デプスマップを乱数で初期化 • 各デプスに対応する初期コストを計算ステップ2：モデル伝搬 • 注目画素に対し、ある近傍画素を参照 • 近傍画素のデプス値を用いて、注目画素におけるコストを計算 • 注目画素のコストと比較して小さい場合、注目画素のデプスを近傍の値に置き換えデプスマップコストマップ Z11 Z12 Z13 Z14 Z15 C11 C12 C13 C14 C15 Z11 Z12 Z13 Z14 Z15 Z21 Z22 Z23 Z24 Z25 C21 C22 C23 C24 C25 Z21 Z22 Z23 Z24 Z25 Z31 Z32 Z33 Z34 Z35 C31 C32 C33 C34 C35 Z31 Z32 Z33 Z34 Z35 Z41 Z42 Z43 Z44 Z45 C41 C42 C43 C44 C45 Z41 Z42 Z43 Z44 Z45 Z51 Z52 Z53 Z54 Z55 C51 C52 C53 C54 C55 Z51 Z52 Z53 Z54 Z55 Zij ← rand(Zmin,Zmax) Cij ← cost(i,j,Zij) C’33 ← cost(3,3,Z22) if C’33 < C33 then Z33 ← Z22 C33 ← C’33 endif ステップ3：ランダム改善 • 注目画素のデプスに対し、乱数を用いて小さい変更を与える • 変更後のデプスに対するコストを評価し、コストが減少したらデプス値を置き換える Z11 Z12 Z13 Z14 Z15 Z21 Z22 Z23 Z24 Z25 Z31 Z32 Z33 Z34 Z35 Z41 Z42 Z43 Z44 Z45 Z51 Z52 Z53 Z54 Z55 Z’33 ← Z33 + rand(-ΔZ,+ΔZ) C’33 ← cost(3,3,Z’33) if C’33 < C33 then Z33 ← Z’33 C33 ← C’33 endif ステップ2,3を指定回数繰り返し ※以降はステップ2,3をまとめてモデル更新と呼ぶことがあります Copyright © Fixstars Group 21

23.

24.

25.

26.

マッチングコストの基本 ⚫ 視点間の座標変換により、2視点の対応点を算出 ⚫ 対応点におけるパッチ(小領域)間の類似度をもとにコストを計算 ○ PatchMatch MVSでは、入力画像をグレースケールに変換した上で、パッチ間のゼロ平均正規化相互相関(ZNCC)を計算し、マッチングコストに利用することが多い逆投影投影視点間の座標変換マッチングコスト計算 Copyright © Fixstars Group 対応点におけるパッチ 25

27.

28.

Homography変換 ⚫ パッチ内の各点は、局所平面 𝑿, 𝒏 上にあるとしてモデル化 ⚫ 平面から導かれるHomography変換※により、画素位置を対応付ける 𝒕𝒏𝑇 𝑲−1 ○ 𝑯 = 𝑲′ 𝑹 − 𝑇 ○ 𝒑′′ = 𝑯𝒑 = 𝑥′′, 𝑦′′, 𝑤′′ 𝑇 ○ 𝒑′ = 𝑥′, 𝑦′, 1 𝑇 = 𝒏 𝑿 1 𝑤′′ 𝒏：法線ベクトル局所平面 𝑿 = 𝑋, 𝑌, 𝑍 ：3D点 𝒑′′ 近傍視点画像参照視点画像変数意味 𝑪, 𝑪′ 参照視点および近傍視点のカメラ 𝑲, 𝑲′ 𝑪, 𝑪′のカメラ行列(内部パラメータ) 𝑿, 𝒏 注目画素に対応する𝑪座標系の3D点と法線ベクトル 𝒑, 𝒑′ 平面 𝑿, 𝒏 上の3D点を𝑪, 𝑪′に投影した点 𝑹, 𝒕 𝑪座標系の3D点を𝑪′座標系に移す剛体変換(回転+並進) 𝑿′ = 𝑹𝑿 + 𝒕 𝒑 = 𝑥, 𝑦, 1 T 𝒑′ = 𝑥′, 𝑦′, 1 T 𝑯 𝑪 𝑪′ ※これをplane-induced Homographyと呼ぶ導出はHomography (computer vision)[11]などを参照のこと Copyright © Fixstars Group 27

https://en.wikipedia.org/wiki/Homography_(computer_vision)

29.

30.

マッチングコスト計算の流れを整理 ⚫ 注目画素における平面モデル𝜽 = 𝑍: 𝒏 を評価したい ⚫ 𝑿 = 𝑋, 𝑌, 𝑍 と𝒏からHomography行列𝑯を計算 ⚫ 参照視点のパッチ内の各点𝒑𝑖 について、近傍視点中の対応点𝒑′𝑖 を計算 ⚫ 対応点における輝度値をサンプリング ○ 𝒔 = 𝐼 𝒑1 , 𝐼 𝒑2 , … , 𝐼 𝒑𝑘 ○ 𝒔′ = 𝐼′ 𝒑′1 , 𝐼′ 𝒑′2 , … , 𝐼′ 𝒑′𝑘 ⚫ 輝度ベクトル間のZNCCからマッチングコストを計算 ○ 𝑚 𝜽 = 1 − ZNCC 𝒔, 𝒔′ ○ ↑相関係数+1のときコスト0となるように Copyright © Fixstars Group 29

31.

Bilateral重みづけによるZNCC ⚫ パッチ内には注目画素が属する対象物の他、それとは関係のない背景などが含まれることがあり、これらは類似度計算においてノイズとなる ⚫ そこで、ZNCCを計算する際に画素位置毎にBilateral重みを設定することで、対象物の寄与度を向上させる σ𝑖 𝑤𝑖 𝑠𝑖 −𝑠ҧ 𝑠′𝑖 −𝑠′ ഥ ○ ZNCC = ○ 𝑤𝑖 = exp − σ𝑖 𝑤𝑖 𝑠𝑖 −𝑠ҧ 2 σ𝑖 𝑤𝑖 𝑠′𝑖 −𝑠′ ഥ 2 ∆𝑔2 ∆𝑓2 2𝜎𝑔 2𝜎𝑓2 ∆𝑔：パッチ中心からの距離 ∆𝑔 = 𝑑𝑥 2 + 𝑑𝑦 2 2 − ∆𝒇 ：パッチ中心における輝度𝒔𝒄と𝒔𝒊 との差 ∆𝑓 = 𝑠𝑖 − 𝑠𝑐 Copyright © Fixstars Group 30

32.

[beta]

Bilateral重みづけによるZNCC
⚫ Bilateral重みづけによるZNCCの実装例
const float fc = I1(xc, yc);
float s12 = 0, s11 = 0, s22 = 0, s1 = 0, s2 = 0, sw = 0;
for (int dy = -PATCH_RADIUS; dy <= PATCH_RADIUS; dy += SAMPLE_STEP)
{
for (int dx = -PATCH_RADIUS; dx <= PATCH_RADIUS; dx += SAMPLE_STEP)
{
const int x1 = cx + dx;
const int y1 = cy + dy;
const auto [x2, y2] = homographyTransform(x1, y1, H);
const float f1 = I1(x1, y1);
const float f2 = I2(x2, y2);
const float w = calcBilateralWeight(dx, dy, f1 - fc);
s12 += w * f1 * f2;
s11 += w * f1 * f1;
s22 += w * f2 * f2;
s1 += w * f1;
s2 += w * f2;
sw += w;
}
}
...以下略...

Copyright © Fixstars Group

31

33.

34.

近傍視点とマッチングコストの集約 ⚫ これまで近傍視点を1つに限定して説明してきたが、実際には一つの参照視点につき複数の近傍視点を割り当てる ○ マッチングの信頼性を向上させるため ○ 近傍視点の割り当て方法としては、前段のSfMで得られた疎な点群(特徴点)を利用し、一定数の特徴点を共有する他の視点を最大N個まで採用する、などがある ⚫ 複数の近傍視点に対して計算したマッチングコストは総和や平均をとるなどして、集約(aggregation)する 1 σ𝑖∈𝑁 𝑚𝑖 𝜽 ○ 𝑚agg 𝜽 = ○ 𝑁は近傍視点のインデックス集合 𝑁 Copyright © Fixstars Group 一定の視野を共有する視点の集合 33

35.

視点選択(View Selection) ⚫ ある注目画素では、対応するパッチが別の物体で隠れたり(オクルージョン)、画像範囲外となることで、近傍視点中に存在しないことがある ⚫ このような場合のマッチングコストは、集約の際にノイズとなる ⚫ これを避けるため、画素ごとに有効な近傍視点を選択することを、視点選択(View Selection)と呼ぶオクルージョンの発生 Copyright © Fixstars Group 34

36.

視点選択手法の例 ⚫ K-best selection (Gipuma[2]) ○ 各視点に対するマッチングコストを小さい順にソートし、上位K件のコストで平均をとる ○ 𝑚′1 , 𝑚′2 , … , 𝑚′ 𝑁 = sort 𝑚1 , 𝑚2 , … , 𝑚𝑁 ○ 𝑚agg 𝜽 = σ𝐾 𝑖=1 𝑚′𝑖 𝜽 1 𝐾 ⚫ Multi-Hypothesis Joint View Selection (ACMH[3]) ○ 注目画素近傍から、8個の平面モデル𝜽1 , 𝜽2 , … , 𝜽8 を取得 ○ 8 × 𝑁のコスト行列を作成し、各列に属するコストの値から、各視点の重み𝑤𝑖 を決定 ○ 𝑚agg 𝜽 = σ 1 𝑤𝑖 σ𝑁 𝑖=1 𝑤𝑖 𝑚𝑖 𝜽 視点の次元 𝑚1,1 ⋮ 𝑀= 𝑚8,1 ⋯ 𝑚1,𝑁 ⋱ ⋮ ⋯ 𝑚8,𝑁 𝑤1 𝑤𝑁 コスト行列 Multi-Hypothesis Joint View Selectionの考え方 • 良い視点では、いくつかのモデルに対して低いコストをとる • 悪い視点では、どんなモデルに対しても高いコストをとる Copyright © Fixstars Group 平面候補(Hypothesis)の次元 35

37.

38.

Propagation Scheme ⚫ 次の方針をPropagation Schemeと呼ぶ ○ 近傍画素としてどの位置を参照するか ○ 全画素をどのような順序で更新していくか ⚫ Gipuma以前の手法では、下図のようなシーケンシャルな伝搬が主流だった ○ 現在の注目画素における伝搬は、直前の更新の結果に依存 ○ 部分的に並列化できる手法はあっても、GPUの性能を十分には発揮できていなかったシーケンシャル伝搬 (図は[8]より引用) Copyright © Fixstars Group 37

39.

40.

41.

Adaptive Checkerboard Sampling (ACMH[3]) ⚫ 伝搬に用いる近傍画素を動的に探索 (下図c) ○ 注目画素周辺に8つの小領域を定義 (薄赤色領域、黒画素は含まない) ○ コストマップ内を探索し、小領域内で最もコストの小さい画素を選択 (黄色画素) ⚫ よい平面モデルを効率的に伝搬 (a) シーケンシャル (b) Symmetric Checkerboard (c) Adaptive Checkerboard 伝搬方法の比較 (図は[3]より引用) Copyright © Fixstars Group 40

42.

ACMHにおけるモデル更新方法のまとめ ⚫ モデル伝搬 ○ Adaptive Checkerboard Samplingにより、8つの候補モデル𝜽1 , 𝜽2 , … , 𝜽8 を取得 ○ 各モデルと各視点の組に対しマッチングコストを計算し、コスト行列を作成 ○ コスト行列から各視点の重みを計算し、マッチングコストを集約 ■ ○ 1 𝑚agg 𝜽𝑗 = σ 𝑤 σ𝑁 𝑖=1 𝑤𝑖 𝑚𝑖 𝜽𝑗 𝑖 注目画素のモデル𝜽0 を含む全候補モデルの内、最小コストのものを伝搬 ■ 𝜽best = argmin 𝑚agg 𝜽𝑗 𝑗∈ 0,1,2,…,8 ⚫ ランダム改善 ○ 候補モデルをランダムに5つ生成 ○ 5つの候補についてマッチングコストを計算し、コストが減少したら更新 Copyright © Fixstars Group 41

43.

44.

45.

Geometric Consistency ⚫ 事前に輝度ベースのPatchMatchで、各視点の中間的なデプスマップを作成 ⚫ 2視点のデプスマップを利用し、視点間の再投影誤差を計算 1. 𝒑2 = proj2←1 𝒑1 , 𝑍1 ← 𝜽 2. 𝒑′1 = proj1←2 𝒑2 , 𝑍2 ← 𝐷2 𝒑2 3. 𝑚geo 𝜽 = 𝒑1 − 𝒑′1 ⚫ 輝度ベースのマッチングコストと組み合わせる ○ 𝑚 𝜽 = 𝑚potho 𝜽 + 𝑤geo 𝑚geo 𝜽 ① ③ Geometric Consistencyの重み ② 参照視点 Copyright © Fixstars Group 近傍視点 44

46.

Multi-Scale PatchMatch ⚫ 画像を低解像度にダウンサンプルすることで、相対的にパッチ(視野)を拡大、パッチ内に有効なテクスチャを含まれやすくする ⚫ 低解像度で推定したデプスマップは、アップサンプルして高解像度側のデプスマップ推定における初期値として利用 Multi-Scale PatchMatchの概略 (図は[3]より引用) Copyright © Fixstars Group 45

47.

48.

49.

Planar Priorの導入 ⚫ Textureless領域が広範囲になると、ACMMでも対応しきれない ⚫ このような領域は壁や床など、平面と仮定できるケースが多く、これを利用した改善手法がいくつか提案されている ⚫ 基本的なアプローチとしては以下の通り ○ 各画素を、 Reliable(周辺にTextureがある)か、Unreliable(Textureless)かのどちらかに、コストの値などを用いて分類 ○ Unreliableな画素における平面モデルを、周囲のReliableな画素から補間 ○ 補間した平面モデルをガイドとして、輝度ベースのマッチングコストと組み合わせる Copyright © Fixstars Group 48

50.

ACMP[4] ⚫ 従来型のPatchMatchを実行後、コストの値をもとにReliableな点を抽出 ⚫ Reliableな点群に対しドロネー三角形分割を適用 ○ 三角形の頂点はReliable、三角形の内部はUnreliableという関係になっている ⚫ 頂点におけるデプスの値から三角形内部のデプス+法線を推定し、これを PatchMatchのガイドとして利用デプス補間の概要 (図は[4]より引用) Copyright © Fixstars Group 49

51.

Adaptive Patch Deformation (APD-MVS[6]) ⚫ 深層学習におけるDeformable Convolutionの考え方を輸入 ○ Deformable Convolution：Convolution時の参照ユニットを動的に学習するCNN ⚫ Unreliableな画素に対し、周囲8点のReliableな画素を探索 ⚫ 同一の平面パラメータを8か所のReliableな位置におけるパッチで評価し、これを従来のマッチングコストに組み込む Adaptive Patch Deformationの概要 (図は[6]より引用) Copyright © Fixstars Group 50

52.

53.

54.

55.

ETH3D benchmark[9] ⚫ Laser Scannerで計測した点群をGT(ground truth)として提供 ⚫ 評価指標は accuracy,completeness,F1 score の3種類 ○ ○ ○ accuracy ■ 復元点を3種類の領域のいずれかに分類 ● 緑：accurate、赤：inaccurate、青：unobserved (Laser視野外のため不問) ■ accuracy = accurate点数 / (accurate点数 + inaccurate点数) completeness ■ あるGT点qiをクエリとし、最近傍の復元点piを検索 ■ completeness = dist(pi,qi) ≦ threshold※を満たすqiの割合 F1 score ■ accuracy (precision) pとcompleteness (recall) r の調和平均 accuracy評価用の領域 (図は[9]より引用) ■ F1 = (p×r)/(p+r) ※MVS手法の論文ではthresholdとして2[cm]や10[cm]が用いられる Copyright © Fixstars Group 54

56.

F1 scoreを上げる方針 ⚫ accuracyの向上 ○ Reliableな画素については、十分に伝搬を繰り返せばいずれ最適解に収束 ○ Unreliableな画素はできるだけ復元しないようにする ⚫ completenessの向上 ○ Textureless領域の復元性能を上げる ○ デプスマップの解像度を大きくする ■ GT点の数が多いため、復元する点群の密度を増やすだけで、ある程度completenessが向上する ■ ETH3Dの入力画像サイズは24メガピクセル(6221×4146)程度で、原解像度のままデプスマップを推定するのが理想的だが、処理時間やメモリの制約から、ある程度縮小するのが一般的 ● ● PatchMatch系の論文では、6.5メガピクセル(3200×2130)程度に縮小深層学習系ではGPUメモリを大量消費するため、さらに解像度を落とさざるを得ない Copyright © Fixstars Group 55

57.

58.

59.

PatchMatch MVSのCUDA実装 ⚫ Gipumaをはじめとして、様々なCUDA実装のPatchMatch MVSが存在 ⚫ これらを参考にしつつも、基本的には1から実装を行った ○ モデル更新はACMHをベースに実装既存のPatchMatch MVS実装手法発表年実装言語 URL Gipuma 2015 C++/CUDA https://github.com/kysucix/gipuma ACMH/ACMM 2019 C++/CUDA https://github.com/GhiXu/ACMM ACMP 2020 C++/CUDA https://github.com/GhiXu/ACMP ACMMP 2022 C++/CUDA https://github.com/GhiXu/ACMMP HPM-MVS 2023 C++/CUDA https://github.com/CLinvx/HPM-MVS APD-MVS 2023 C++/CUDA https://github.com/whoiszzj/APD-MVS Copyright © Fixstars Group 58

60.

CUDAの簡単なおさらい ⚫ スレッドおよびメモリの階層スレッドの階層構造 • スレッド間に階層構造がある • 近いスレッド同士はより密に通信・同期を行うことができるメモリの階層構造 • おおむねスレッドの階層構造と対応 Fixstars CUDA高速化セミナーvol.1 ～画像処理アルゴリズムの高速化～より引用 https://www.docswell.com/s/fixstars/K24MYM-20220527 Copyright © Fixstars Group 59

https://www.docswell.com/s/fixstars/K24MYM-20220527

61.

62.

スレッド割り当ての方針 ⚫ Red-Black伝搬を採用、1つの赤(黒)画素に1スレッドを割り当て ⚫ Thred Blockのサイズを16×16に設定 → 画像中の32×16の領域を担当 ○ x方向には1画素おきにスレッドを割り当てるため、処理領域の幅は倍になる x方向には1画素おきにスレッドを割り当て 16 32 16 16 Thread Block 入力画像 Copyright © Fixstars Group 61

63.

[beta]

モデル更新のCUDA実装例
⚫ カーネル関数及び、カーネル関数呼び出し

__global__ void propagateAndRefineACMHKernel(float* Z, Vec3f* N, float* C,
int w, int h, int nviews, int color, float minZ, float maxZ, bool geom)
{
// checkerboardパターン(市松模様)となるようx,yを計算
const int x = 2 * (blockIdx.x * blockDim.x + threadIdx.x) + ((y % 2) ^ color);
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= w || y >= h)
return;

void propagateAndRefineACMH(GpuMat& depths, GpuMat& normals, GpuMat& costs,
int nviews, int color, float minZ, float maxZ, bool geom)
{
const int w = depths.cols;
const int h = depths.rows;
const int npixelsX = w / 2; // x方向はw/2個の画素を更新
const int npixelsY = h;

MatchingCost MC(nviews, geom); // マッチングコスト計算クラス
float viewWeights[MAX_VIEWS]; // コスト集約時の視点重み

constexpr int BLOCK_SIZE = 16;
const dim3 block(BLOCK_SIZE, BLOCK_SIZE);
const dim3 grid(divUp(npixelsX, block.x), divUp(npixelsY, block.y));

// モデル伝搬
propagateACMH(MC, Z, N, C, w, h, x, y, minZ, maxZ, viewWeights);

auto Z = depths.ptr<float>();
auto N = normals.ptr<Vec3f>();
auto C = costs.ptr<float>();

// ランダム改善
Xorshift random(getRandomSeed(x, y, w)); // 乱数生成器
refineACMH(MC, Z, N, C, w, h, x, y, minZ, maxZ, viewWeights, random);

propagateAndRefineACMHKernel<<<grid, block>>>(Z, N, C, w, h,
nviews, color, minZ, maxZ, geom);
CUDA_CHECK(cudaGetLastError());

}

}

モデル伝搬(ACMH)のカーネル関数

カーネル関数呼び出し

Copyright © Fixstars Group

62

64.

1スレッドのタスク (ACMH再掲) ⚫ モデル伝搬 ○ Adaptive Checkerboard Samplingにより、8つの候補モデル𝜽1 , 𝜽2 , … , 𝜽8 を取得 ○ 各モデルと各視点の組に対しマッチングコストを計算し、コスト行列を作成 ○ コスト行列から各視点の重みを計算し、マッチングコストを集約 ■ ○ 1 𝑚agg 𝜽𝑗 = σ 𝑤 σ𝑁 𝑖=1 𝑤𝑖 𝑚𝑖 𝜽𝑗 𝑖 注目画素のモデル𝜽0 を含む全候補モデルの内、最小コストのものを伝搬 ■ 𝜽best = argmin 𝑚agg 𝜽𝑗 𝑗∈ 0,1,2,…,8 ⚫ ランダム改善 ○ 候補モデルをランダムに5つ生成 ○ 5つの候補について伝搬と同様にコストを計算し、コストが減少したら更新 Copyright © Fixstars Group 63

65.

処理のボトルネック ⚫ マッチングコスト計算(ZNCC)が最大のボトルネック ○ モデル伝搬～ランダム改善で、多数の候補モデル𝜽を評価 ⚫ さらに要因を分類すると… ○ ○ ○ メモリアクセス ■ 輝度値のサンプリング時に、グローバルメモリへのアクセスが発生 ■ 特に近傍視点側は参照位置のランダムアクセス性が高く、キャッシュヒット率も低 ■ ↑に加えて参照位置が浮動小数点数となるため、バイリニア補間が必要演算 ■ パッチ内の各画素に対するHomography変換には除算が含まれ、演算コスト大 Warp Divergence ■ 画像の内外判定などの条件分岐によって有効な演算を行わないスレッドが発生 ■ ワープ内での異なる方向への分岐は性能劣化につながる Copyright © Fixstars Group 64

66.

[beta]

処理のボトルネック
⚫ ZNCC計算部のナイーブな実装例

Warp Divergence
• ワープ内のスレッドが異なる
方向に分岐する可能性

for (int dy = -PATCH_RADIUS; dy <= PATCH_RADIUS; dy += SAMPLE_STEP)
{
for (int dx = -PATCH_RADIUS; dx <= PATCH_RADIUS; dx += SAMPLE_STEP)
{
const int x1 = cx + dx;
const int y1 = cy + dy;
if (x1 < 0 || x1 >= w1 || y1 < 0 || y1 >= h1)
continue;
const auto [x2, y2] = homographyTransform(x1, y1, H);
if (x2 < 0 || x2 >= w2 || y2 < 0 || y2 >= h2)
continue;
const float f1 = I1[y1 * w1 + x1];
const float f2 = sampleBilinear(I2, w2, h2, x2, y2);

Homography変換
• 除算が発生
メモリアクセス
• グローバルメモリへのアクセスが発生
• 近傍視点側はバイリニア補間が必要

const float w = calcBilateralWeight(dx, dy, f1 - fc);
s12 += w * f1 * f2;
s11 += w * f1 * f1;
s22 += w * f2 * f2;
s1 += w * f1;
s2 += w * f2;
sw += w;
}
}

Copyright © Fixstars Group

65

67.

Shared Memoryの使用 ⚫ Shared Memory Thread Block内で共有可能な、64KB～100KB程度の高速なメモリ ○ ⚫ 参照画像については、同じ画素に繰り返しアクセスするため、データをあらかじめShared Memoryに保持 ⚫ 近傍画像については、アクセス位置が事前に分からないため、利用を断念 16 32+Δ 16+Δ 16 Thread Block パッチサイズの分、拡張した領域をコピー参照画像 Copyright © Fixstars Group 66

68.

[beta]

Texture Fetchingの使用
⚫ Texture Fetchingとは
○

GPUのテクスチャユニットを使用した、画像データへのアクセス方法 (詳細は次頁)

⚫ CUDAからはTexture Objectを介して利用
// 画像サイズやデバイスメモリの情報を設定
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(cudaResourceDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = image.data;
resDesc.res.pitch2D.height = image.rows;
resDesc.res.pitch2D.width = image.cols;
resDesc.res.pitch2D.pitchInBytes = image.step;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<uchar>();
// Texture Fetchingの設定
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(cudaTextureDesc));
texDesc.addressMode[0] = addressMode;
texDesc.addressMode[1] = addressMode;
texDesc.addressMode[2] = addressMode;
texDesc.filterMode = filterMode;
texDesc.readMode = readMode;
texDesc.normalizedCoords = 0;

__device__ inline float sampleTex(cudaTextureObject_t obj, float x, float y)
{
return tex2D<float>(obj, x + 0.5f, y + 0.5f);
}

Texture Fetching (デバイス側処理)
• tex2D<T>を用いてデータを取得、座標はfloatで指定
• 整数座標における1ピクセルはfloat座標で0~1の範囲をとる
• 整数座標上の(x,y) は、float座標上で(x+0.5,y+0.5) と対応

// Texture Objectの作成
cudaTextureObject_t textureObject;
cudaCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);

Texture Objectの作成 (ホスト側処理)
Copyright © Fixstars Group

67

69.

Texture Fetchingの使用 ⚫ Texture Fetchingのメリット ○ Textureキャッシュが働き、局所的な2Dアクセスの効率が向上 ■ ただし最近のGPUではL1キャッシュとTextureキャッシュは統合されており(unified L1/Texture cache)、同等の効果を自動的にL1キャッシュでまかなえることも多い ○ ○ Filter Mode機能 ■ データ取得時の補間方法を指定可能 ■ cudaFilterModeLinearを設定することで、バイリニア補間を利用可能 Address Mode機能 ■ 画像範囲外アクセス時のアドレッシング方法を指定可能 (Wrap,Clamp,Mirror,Border) ■ 画像の内外判定が不要に ⚫ これらメリットから、近傍画像へのアクセスにはTexture Fetchingを使用 Copyright © Fixstars Group 68

70.

Homography変換の近似 ⚫ Homography変換式 ○ 𝑥′ = ℎ11 𝑥+ℎ12 𝑦+ℎ13 ℎ31 𝑥+ℎ32 𝑦+ℎ33 , 𝑦′ = ℎ21 𝑥+ℎ22 𝑦+ℎ23 ℎ31 𝑥+ℎ32 𝑦+ℎ33 ⚫ ここでは𝑥 ′ の計算に着目、𝑦を固定し変換式を𝑥の関数とみなす ○ 𝑥′ = 𝑓 𝑥 = 𝑢 𝑥 𝑤 𝑥 ⚫ これをパッチ中心𝑐𝑥 の周りで1次近似 ○ 𝑓 𝑥 ≃ 𝑓 𝑐𝑥 + 𝑓 ′ 𝑐𝑥 𝑥 − 𝑐𝑥 = 𝑏𝑥 + 𝑎𝑥 𝑑𝑥 ○ 𝑏𝑥 = 𝑓 𝑐𝑥 , 𝑎𝑥 = 𝑓 ′ 𝑐𝑥 = 𝑢′ 𝑐𝑥 𝑤 𝑐𝑥 −𝑢 𝑐𝑥 𝑤 ′ 𝑐𝑥 𝑤 𝑐𝑥 2 , 𝑑𝑥 = 𝑥 − 𝑐𝑥 ⚫ 𝑥に関するループ内の除算を回避 Copyright © Fixstars Group 69

71.

[beta]

ZNCC計算の実装例
⚫ 高速化前後の比較
for (int dy = -PATCH_RADIUS; dy <= PATCH_RADIUS; dy += SAMPLE_STEP)
{
for (int dx = -PATCH_RADIUS; dx <= PATCH_RADIUS; dx += SAMPLE_STEP)
{
const int x1 = cx + dx;
const int y1 = cy + dy;
if (x1 < 0 || x1 >= w1 || y1 < 0 || y1 >= h1)
continue;

Sharedメモリの使用

const auto [x2, y2] = homographyTransform(x1,• y1,
H);
参照画像の読み込みを
if (x2 < 0 || x2 >= w2 || y2 < 0 || y2 >= h2)
高速化
continue;

for (int dy = -PATCH_RADIUS; dy <= PATCH_RADIUS; dy += SAMPLE_STEP)
{
const float ax, bx = ...;
const float ay, by = ...;
for (int dx = -PATCH_RADIUS; dx <= PATCH_RADIUS; dx += SAMPLE_STEP)
{
const int x1 = cx + dx;
Homography変換の近似
const int y1 = cy + dy;
• ループ内の除算を回避
const float x2 = ax * dx + bx;
const float y2 = ay * dy + by;
const float f1 = fetchFromShared(x1, y1, ...);

const float f1 = I1[y1 * w1 + x1];
Texture
const float f2 = sampleBilinear(I2, w2, h2, x2,
y2); Fetching

•

const float f2 = fetchFromTexture(x2, y2, ...);

バイリニア補間

const float w = calcBilateralWeight(dx, dy, f1 - fc);
s12 += w * f1 * f2;
s11 += w * f1 * f1;
s22 += w * f2 * f2;
s1 += w * f1;
s2 += w * f2;
sw += w;

const float w = calcBilateralWeight(dx, dy, f1
fc);
• -内外判定の除外
s12 += w * f1 * f2;
s11 += w * f1 * f1;
s22 += w * f2 * f2;
s1 += w * f1;
s2 += w * f2;
sw += w;
}

}
}

}

高速化前

高速化後
Copyright © Fixstars Group

70

72.

処理時間計測 ⚫ 計測環境 GPU GeForce RTX 3080 入力画像 ETH3Dデータセットのofficeシーン (参照画像はDSC_0219.JPG) 画像サイズ 1600×1066にリサイズ視点数 21視点(参照1 + 近傍20) 計測区間ランダムな初期状態※1から、モデル更新(赤画素+黒画素)×3反復の時間を計測 ※1 デプスがランダムの時、画像内外判定による分岐が発生しやすい ⚫ 計測結果計測項目処理時間 ※2 ナイーブ実装高速化後 312.3[msec] 110.0[msec] Compute Throughput 68.2[%] 66.3[%] Memory Throughput 45.5[%] 84.1[%] Avg. Divergent Branches 40216.5 7244.2 ナイーブ実装から 3倍程度高速化 ※2 NVIDIA Nsight Computeで取得した各種指標から抜粋 (計6回のカーネル実行の平均) Copyright © Fixstars Group 71

https://developer.nvidia.com/nsight-compute

73.

Compute Throughputが十分に上がらない要因 ⚫ レジスタ大量使用によるOccupancy(GPU利用率)低下 ○ GPUのStreaming Multiprocessor(Thread Blockが割り当てられる演算ユニット)にはレジスタ数の上限がある ○ 1スレッドが多くのレジスタを使用すると、相対的に同時実行可能なスレッド数が減少 ⚫ レジスタ数の制限 ○ CUDAのコンパイルオプションで、 1スレッドが使用するレジスタ数を制限できる ○ Occupancyは改善できるが、レジスタに乗らない変数はLocal Memoryに退避されメモリアクセスが増加するため、トレードオフの関係にある ○ レジスタ数の制限も試してみたが、現時点では効果が見られていない Copyright © Fixstars Group 72

74.

75.

参考文献 [1] Furukawa, Yasutaka, and Carlos Hernández. "Multi-view stereo: A tutorial." Foundations and Trends® in Computer Graphics and Vision 9.1-2 (2015): 1-148. [2] Galliani, Silvano, Katrin Lasinger, and Konrad Schindler. "Massively parallel multiview stereopsis by surface normal diffusion." Proceedings of the IEEE international conference on computer vision. 2015. [3] Xu, Qingshan, and Wenbing Tao. "Multi-scale geometric consistency guided multi-view stereo." Proceedings of the IEEE/CVF conference on computer vision and pattern recognition. 2019. [4] Xu, Qingshan, and Wenbing Tao. "Planar prior assisted patchmatch multi-view stereo." Proceedings of the AAAI conference on artificial intelligence. Vol. 34. No. 07. 2020. [5] Xu, Qingshan, et al. "Multi-scale geometric consistency guided and planar prior assisted multi-view stereo." IEEE Transactions on Pattern Analysis and Machine Intelligence 45.4 (2022): 4945-4963. [6] Wang, Yuesong, et al. "Adaptive patch deformation for textureless-resilient multi-view stereo." Proceedings of the IEEE/CVF Conference on Computer Vision and Pattern Recognition. 2023. [7] Ren, Chunlin, et al. "Hierarchical prior mining for non-local multi-view stereo." Proceedings of the IEEE/CVF International Conference on Computer Vision. 2023. [8] Zheng, Enliang, et al. "Patchmatch based joint view selection and depthmap estimation." Proceedings of the IEEE conference on computer vision and pattern recognition. 2014. [9] Schops, Thomas, et al. "A multi-view stereo benchmark with high-resolution images and multi-camera videos." Proceedings of the IEEE conference on computer vision and pattern recognition. 2017. [10] Guédon, Antoine, and Vincent Lepetit. "Sugar: Surface-aligned gaussian splatting for efficient 3d mesh reconstruction and high-quality mesh rendering." Proceedings of the IEEE/CVF Conference on Computer Vision and Pattern Recognition. 2024. [11] https://en.wikipedia.org/wiki/Homography_(computer_vision) Copyright © Fixstars Group 74

76.

コンピュータビジョンセミナーvol5～Multi-View StereoのCUDA高速化～（2024/8/7)

株式会社フィックスターズ

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

いまさら聞けない！CUDA高速化入門 ～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

各ページのテキスト

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）