【KMC春合宿2024】実装視点で見るNeural Radiance Fields

KMC 春合宿2024の講座として発表したスライド資料です．内容はNeural Radiance FieldsをCUDA C++でフルスクラッチ実装してみたので，そのうえで実装に近い視点でNeural Radiance Fieldsを説明していくものとなります．実装に完全に特化した説明を読みたい場合は私が先日書いた記事を参考にしていただけると嬉しいです．CUDAによる実装（ソースコード）の一部とそれに対する説明，そしてそのプログラムを動かした結果などが記述されています．
[URL]
Part1: MLP編: https://tetrisyoshi.hatenablog.com/entry/2024/01/21/213713
Part2: エンコーダー編: https://tetrisyoshi.hatenablog.com/entry/2024/02/13/141046
Part3: NeRF編: https://tetrisyoshi.hatenablog.com/entry/2024/02/24/155515

Transcript

1. 実装視点で見るNeural Radiance Fields Yoshi’s Dream KMCID: tetrisyoshi 1
2. 自己紹介 本名: (deleted) / HN: Yoshi’s Dream / KMCID: tetrisyoshi 電気電子工学科 B2 / KMC47th 普段: 主に3DCGのレンダリング処理を書いてます． 興味: 物理現象のシミュレーション，画像処理，計算機のハードウェア最適化など けものフレンズはいいぞ 2
3. 本講座の目的 CUDAプログラミングにおける簡単な概念を得る ニューラルネットワークの基本的な理解をする NeRFに対する概形のみの理解にとどまらず，実装に必要な具体的な処理について理解する 3
4. 目次の雰囲気 NeRFに関するぼんやりした説明，そして実装したもののデモ紹介 CUDAプログラミングの基礎 パラメーター空間の探索の基礎 NeRFに関する少し具体的な説明 ボリュームレンダリングのお気持ち NeRFの誤差逆伝播 Occupancy Gridの説明 Neural Networkの基本 Neural Networkの実装の説明，NVIDIA GPUによる行列演算 エンコーダーについての説明 4
5. Neural Radiance Fields [Ben M. et.al. ECCV 2020] 略して“NeRF（ナーフ）”と書かれることが多い 2次元の色んな視点から見た空間の画像をもとに，対象の3次元空間を推定 5
6. Instant NeRF [NVIDIA, 2022] 6
7. 実装したもの 7
8. 実装したもの 8
9. 実装したもの 9
10. 準備 CUDAプログラミングの基礎 修飾子と記述 メモリの扱い メモリの種類 パラメータ空間の探索 勾配降下法 画像近似 2次元平面の再現と3次元空間の再現の狭間 10
11. 準備: CUDAについて NVIDIA GPU上でプログラムを動かすためのプラットフォーム的なもの CUDA C++は基本的にはC, C++と同様の記述が可能 GPU上で並列に実行できるため，並列性の高いプログラムに有効 ただし，並列処理を書く上で気に掛けることは多い 11
12. CUDAプログラムの記述: 修飾子 __global__ __host__ __device__の3種類 CPUをホスト，GPUをデバイスと呼ぶことが多い． __host__ ホスト側で呼び出され，ホスト側で実行される関数 ただのC++ __host__の関数Add(), PrintAdd() 出力結果 12
13. __device__ デバイス(GPU)から呼び出され，デバイス(GPU)で動く関数 GPU上で動くだけで，記述はCPU側と基本的に同じ ホスト(CPU)から呼び出すことは不可 __host__ __device__ 13
14. __global__ カーネル(Kernel)とも呼ばれる ホスト(CPU)から呼び出され，デバイス(GPU)上で動く関数 （GPUコードの入り口） 返り値は取れない（void型のみ許容される） <<<GridSize, BlockSize>>>というCUDA特有の指定が必要（後述） 14
15. CUDAプログラムのイメージ main() __host__ __global__ __device__ GPU CPU __host__ __device__ 15
16. スレッドの単位 (CUDA) Warp: 32スレッドの「塊」 Warp内のスレッドは同時に同じ命令を実行しようとする if分岐などによってスレッドごとに実行する命令が異なる場合，異なる命令のスレッドは眠る 並列性の低下（Branch Divergence） Warpを意識したコード設計が高速動作に必須 1 2 3 4 5 6 7 8 if (id <= 4) { A(); } else { B(); } 1 2 3 4 z A() z z z 5 6 7 8 z B() z z z 1 2 3 4 5 6 7 8 diverge converge 16
17. スレッドの単位 (CUDA) Block: スレッドのグループ Streaming Multiprocessor (SM) 内部で動くスレッド ブロック内部ではデータを高速に共有することが可能（後述） CUDA実行時にBlockSizeとしてスレッド数を指定 CUDA実行中に実行中のスレッド番号をthreadIdxで受け取れる CUDA実行中にスレッド数をblockDimで受け取れる 17
18. スレッドの単位 (CUDA) Grid: Blockの集まり 最大のスレッド単位 CUDA実行時にGridSizeとしてBlock数を指定 CUDA実行中に実行中のブロック番号をblockIdxで受け取れる 右図の例では1Blockにつき32スレッドでそれを4Grid用意している すなわち128スレッド並列 ブロック番号にブロック内のスレッド数をかけ，それにブロック内での スレッド番号を足すことで実行中のスレッド番号が分かる！ <<<GridSize, BlockSize>>> 18
19. スレッドの単位のイメージ … Warp Warp Warp Warp Warp … Thread Warp Block Grid Block Block Block Block … Block 19
20. メモリの扱い (CUDA) ホスト(CPU)とデバイス(GPU)は別々のメモリを使用する CPUプログラムはGPUメモリから直接データを読み出すことはできない（逆も然り） CPU-GPUメモリ転送(Host-to-Device, Device-to-Host)が必要 メモリ転送には時間がかかるため，ここの設計は重要 C言語のmalloc, memset, memcpyみたいなアレをやる Thrustライブラリの使用により，安全かつ簡単に実装できる！（詳細は略） 20
21. No Thrust Thrust 21
22. 主なメモリの種類 (CUDA) ホスト(CPU)から転送する先のデバイス(GPU)メモリはglobal memoryという領域 global memory いわゆるGPUメモリ 大容量！RTX3080は10GBあります．A100なら80GBあります． 遅い！！！ shared memory Blockごとに共有されているメモリ領域 On-Chipメモリであり，読み書きが非常に高速 容量が小さい！RTX3080だと48KB． __shared__ 修飾子で確保する 22
23. 準備: CUDAについて さいごに Warpを意識した設計を行うことにより，実行時のDivergenceを避ける Warp内部は共同作業をする意識で，なるべく分岐はさせない アルゴリズムに問題がある可能性も考慮する Blockとshared memoryを意識し，メモリ読み書きの高速化を図る なるべくglobal memoryとの読み書きを減らし，shared memoryを使用する CPU-GPU間メモリ転送はマジで最小限に 23
24. 準備: パラメーターの探索 関数が次で与えられる 目的の値を次で定める どんなxがこれを満たすかを求めたい．「誤差の大きさ」を次の式で定義する 言うまでもなく，x = -1のとき，L(x)は最小(= 0)となる． 𝑓(𝑥) = 𝑥2 + 2𝑥 + 2 g(𝑥) = 1 L 𝑥 = 1 2 𝑓 𝑥 − 𝑔 𝑥 2 0 1 2 3 4 5 -3 -2 -1 0 1 24
25. 準備: パラメーターの探索 L(x)を最小化するパラメーターxはどうやって求めるか？ 右の例では解析的に簡単に求まる．x = -1 複雑な関数だと厳しい 数値計算で頑張るにはどうする？ 勾配降下法を使用してみると可能．次の手続きを行う． xを𝑥0 で初期化する. 収束するまで次を行う 𝑑𝐿 𝑑𝑥 を計算する 𝑥 = 𝑥 − 𝜂 𝑑𝐿 𝑑𝑥 𝑓(𝑥) = 𝑥2 + 2𝑥 + 2 g(𝑥) = 1 L 𝑥 = 1 2 𝑓 𝑥 − 𝑔 𝑥 2 25
26. 準備: パラメーターの探索 xの初期値を2とする．ηを0.1とする 𝑑𝐿 𝑑𝑥 = 𝑑𝐿 𝑑(𝑓−𝑔) 𝑑 𝑓−𝑔 𝑑𝑥 = 𝑓 − 𝑔 2𝑥 + 2 x0 = 2.0 x1 = -3.4 x2 = -0.6352 x3 = -0.644909 x4 = -0.653864 … x1e3 = -0.950463 x1e5 = -0.994703 (1) xを𝑥0 で初期化する. (2) 収束するまで次を行う 2.1 𝑑𝐿 𝑑𝑥 を計算する (2.2) 𝑥 = 𝑥 − 𝜂 𝑑𝐿 𝑑𝑥 𝑓(𝑥) = 𝑥2 + 2𝑥 + 2 g(𝑥) = 1 L 𝑥 = 1 2 𝑓 𝑥 − 𝑔 𝑥 2 0 3 6 9 12 15 -4 -3 -2 -1 0 1 2 3 26
27. 準備: パラメーターの探索 先ほどの議論はパラメータ空間が高次元になっても行える そもそも勾配降下法とは何か？ あるパラメーター空間において，関数が同じ値をとる「面」がある その面に対して垂直に伸びるベクトルが勾配 誤差を小さくする方向に延びるベクトルを求め，それをもとにパラメーターを更新している ニューラルネットワークは，いわばクソデカ関数 目的関数にそのクソデカ関数を近づけたい 関数のパラメーターに対する勾配を計算し，それをもとにパラメーターを更新すればよい 誤差逆伝播によりこの計算が可能になる 27
28. 概念の準備: 画像近似 画像は(x, y)を与えると(R, G, B)を返す関数としてみることが出来る (x,y)を入力としたニューラルネットワークで(R,G,B)を推定 x y … … …… … … R G B 28
29. 3次元空間の推定？ 画像近似は2次元平面の推定であった 推定結果(R,G,B)は(x,y)におけるピクセルの色として直ちに知覚される（画像となる） では3次元空間はどうか？3次元空間を2次元画像に落とし込むには？ レンダリングを行う必要がある 3次元空間を表現する何かしらの情報が必要 ボリューム（媒質）として扱う 3DのGaussianを空間に置きまくる [Bernhard et.al. SIGGRAPH 2023] NeRF 29
30. Neural Radiance Fields 言うまでもなく，3次元空間を推定する（例えば単位立方体の内部の空間で考える） 3次元空間をボリューム（媒質）で表現する 媒質をある方向から見た際の，ある場所での媒質のパラメーターを推定する パラメーター: 「色」，「密度」 そのパラメーターをもとに，ボリュームレンダリングを行い，画像を得る！ Position Direction Neural Network Volume Parameter Volume Rendering R G B 30
31. 3D空間のレンダリング？ 3D空間はカメラ等の視覚デバイスにより2次元の画像情報となる 画像上の色となるのは，カメラに入ってきた光のエネルギー 厳密な話はさておき，放射輝度(Radiance)として扱うことが多い 光の波長より十分大きな空間を考えると，光は光線（レイ）として扱える（幾何光学） カメラに入ってくるレイに対する放射輝度の変化を考える 31
32. ボリュームにおける放射輝度の変化 媒質内部では粒子が飛び回っている 粒子にぶつかることにより (1) レイが散乱し，放射輝度が減衰する (2) レイが吸収され，放射輝度が減衰する (3) レイが散乱し，放射輝度が増加する 粒子は発光していることもある（ろうそくの炎とか） (4) 粒子の発光により，放射輝度が増加する (1) (2) (3) (4) 32
33. 先ほど示した放射輝度の変化を計算したい t_nearからt_farまでに充満するボリュームの放射輝度への関与を考える C: 「色」 / σ: 「密度」 / r(t): カメラから距離tの座標 / dir: カメラの視線ベクトル ボリュームにおける放射輝度の計算 𝑡𝑛𝑒𝑎𝑟 𝑡𝑓𝑎𝑟 レイ ボリューム 𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒 𝑟𝑎𝑦 = න 𝑡𝑛𝑒𝑎𝑟 𝑡𝑓𝑎𝑟 𝑇 𝑡 𝜎 𝑟 𝑡 𝐶 𝑟 𝑡 , 𝑑𝑖𝑟 𝑑𝑡 𝑇 𝑡 = exp − න 𝑡𝑛𝑒𝑎𝑟 𝑡 𝜎 𝑟 𝑠 𝑑𝑠 33
34. ボリュームにおける放射輝度の計算 計算機上で積分を行うのは厳しい レイ上にサンプル点を取り，離散的に計算する 細かい話は置いといて，下に示す式となる [“Optical Models for Direct Volume Rendering", Nelson Max, 1995] 𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒 𝑟𝑎𝑦 = න 𝑡𝑛𝑒𝑎𝑟 𝑡𝑓𝑎𝑟 𝑇 𝑡 𝜎 𝑟 𝑡 𝐶 𝑟 𝑡 , 𝑑𝑖𝑟 𝑑𝑡 𝑇 𝑡 = exp − න 𝑡𝑛𝑒𝑎𝑟 𝑡 𝜎 𝑟 𝑠 𝑑𝑠 Camera ボリューム C1, σ1 C2, σ2 C3, σ3 C4, σ4 C5, σ5 δ1 δ2 δ3 δ4 Radiance 𝑟𝑎𝑦 = σ𝑖=1 𝑁 𝑇𝑖 (1 − exp −𝜎𝑖 𝛿𝑖 )𝐶𝑖 , where 𝑇𝑖 = exp(− σ 𝑗=1 𝑖−1 𝜎𝑗 𝛿𝑗 ) 34
35. Neural Radiance Fields 一旦整理 レイの上の点をいくつかサンプリングし，PositionとDirectionを得る PositionとDirectionをNNに入力し，「色」と「密度」，つまりCとσを得る ボリュームレンダリングの式に代入し，そのレイの放射輝度をRGBで求める Position Direction Neural Network Volume Parameter Volume Rendering R G B Radiance 𝑟𝑎𝑦 = σ𝑖=1 𝑁 𝑇𝑖 (1 − exp −𝜎𝑖 𝛿𝑖 )𝐶𝑖 , where 𝑇𝑖 = exp(− σ 𝑗=1 𝑖−1 𝜎𝑗 𝛿𝑗 ) 35
36. Neural Radiance Fields 誤差逆伝播 ボリュームレンダリングにより，ある視点から見た空間の画像が得られた(正確な表現は後程) 教師画像データと比較し，誤差を計算し，Volume Parameterの最適化を行いたい 正確には，Neural Networkの最適化を行いたい 36
37. Neural Radiance Fields 誤差逆伝播 画像との誤差はイイ感じに計算してあげると良いが，問題はそこではない Volume Renderingを誤差逆伝播させる必要がある そもそも微分可能なのか？ （これにより微分可能なレンダリング手法がここに求められる） Position Direction Neural Network Volume Parameter Volume Rendering R G B 𝜕𝐿 𝜕𝑅𝐺𝐵 37
38. ボリュームレンダリング 誤差逆伝播 順方向は以下の式で計算できた やはり細かい計算はおいといて，次の式で微分が求まる． Radiance_k: カメラに入射するレイの放射輝度の内，k番目のサンプル点による放射輝度 つまりS_iはi+1番目以降のサンプル点による寄与の量を表す Radiance 𝑟𝑎𝑦 = σ𝑖=1 𝑁 𝑇𝑖 (1 − exp −𝜎𝑖 𝛿𝑖 )𝐶𝑖 , where 𝑇𝑖 = exp(− σ 𝑗=1 𝑖−1 𝜎𝑗 𝛿𝑗 ) 𝜕𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒 𝜕𝐶𝑖 = 𝑇𝑖 (1 − exp(−𝜎𝑖 𝛿𝑖 )) 𝜕𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒 𝜕𝜎𝑖 = 𝜎𝑖 (𝑇𝑖+1 𝐶𝑖 − 𝑆𝑖 ) 𝑆𝑖 = 𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒 − ෍ 𝑘=1 𝑖 𝑅𝑎𝑑𝑖𝑎𝑛𝑐𝑒𝑘 1 2 3 4 5 38
39. Neural Radiance Fields 誤差逆伝播 ボリュームレンダリングの誤差逆伝播が終わり，NNの勾配が得られた NNの処理は現状ブラックボックスにしておきます． いずれにせよこの状態からNNを最適化することができます 実はNeRFのやっていることはこれだけです これによりVolume Parameterが最適化され，Volumeの状態，即ち3次元空間の推定ができます Position Direction Neural Network Volume Parameter Volume Rendering R G B 𝜕𝐿 𝜕𝐶 𝜕𝐿 𝜕𝜎 39
40. 概念としてのNeRF: まとめ (1): カメラから画像のピクセルに対応する方向にレイ（光線）を飛ばす (2): レイの経路上の点をサンプリング (3): サンプル点の座標，そしてレイの向きをニューラルネットワーク(NN)に入力 (4): NNを実行し，サンプル点における「色」と「密度」を計算 (5): 得られたサンプル点の情報からボリュームレンダリング方程式を計算 (6): 計算結果と対応するピクセルの色を比較し，誤差を計算 (7): ボリュームレンダリングの処理を逆方向に行う (8): NNを誤差逆伝播 (9): NN最適化 (10): Occupancy Gridの最適化（16イテレーションごと） 40
41. カメラからレイを飛ばす？ 結果的には，カメラの「スクリーン上」に画像が現れるというイメージ ピンホールカメラモデル 理想的なカメラ 現実： レンズによる歪み OpenCVカメラなど（説明は省略） 41
42. Occupancy Grid 原著論文にはない，Instant NeRFにて導入された概念 「見える」物質は「密度」が大きい媒質としてみることが出来る 一般的に，「空間を近似する」といえど，近似対象は「空間における物体」である 物体がない空間をサンプリングしても嬉しくない どこに物体があるかを把握し，効率の良いサンプリングを行いたい 「密度」の高い物質より向こう側のサンプル点は見えない（放射輝度への寄与の度合いが小さい） ある一定の寄与の度合いを下回ったらサンプル点を無視したい 無視された点 採用されたサンプル点 42
43. Occupancy Gridの構築 空間を一辺128マスのブロックに分割する （Morton Codeで陰に行う） 2種類のOccupancy Gridを構築する (1): Bit Occupancy Grid: 「物体」の有無を管理 (2): Float Occupancy Grid: 「密度」の程度を管理 NNの最適化時に(2)を更新し，その結果得られた(2)に基づいて(1)を更新する 手続き: (i): (2)の全てのブロックの値を0.95倍 (ii): (2)からM個のブロックを選択し，中身を現在の値，ブロック内のある一点における「密度」 の大 きい方で更新する (iii): (2)がある閾値を超えていれば(1)の対応するブロックをtrueに，そうでなければfalseにする． 43
44. Occupancy Gridの実装に関する処理 Morton Codeとは？ 空間充填曲線の一種 3次元で表現される座標を再帰構造で1次元で表現する 空間的に近い→メモリ上で近い メモリアクセスの局所性向上 例えば解像度を1辺4bitで表現するとすると（表現法の一例） x: 𝑥1 𝑥2 𝑥3 𝑥4 y: 𝑦1 𝑦2 𝑦3 𝑦4 z: 𝑧1 𝑧2 𝑧3 𝑧4 Morton Code: 𝑥1 𝑦1 𝑧1 𝑥2 𝑦2 𝑧2 𝑥3 𝑦3 𝑧3 𝑥4 𝑦4 𝑧4 Morton Codeでソートすると右図のZみたいなカーブに沿ってindexが振られる 44 Z階数曲線 Wikipedia
45. Occupancy Gridの実装に関する処理 3次元空間のMortonCodeでの表現 45
46. Occupancy Gridの実装に関する処理 M個のサンプリング 256itersまではM=128*128*128 それ以降はM = 128*128*128/2 M/2個は全部のブロックから一様サンプリング 残りのM/2個は「Bit_OccupancyGridがtrueであるブロック」から一様サンプリング 棄却法によりサンプリング 別にしなくても学習ではできますｗ 46
47. OccupancyGridを考慮したサンプリング処理 47
48. OccupancyGridを考慮したサンプリング処理 48 Radiance 𝑟𝑎𝑦 = σ𝑖=1 𝑁 𝑇𝑖 (1 − exp −𝜎𝑖 𝛿𝑖 )𝐶𝑖 , where 𝑇𝑖 = exp(− σ 𝑗=1 𝑖−1 𝜎𝑗 𝛿𝑗 )
49. サンプリング処理の実装上の苦難 サンプル点の数は動的である では，サンプル点数に上限を設けてサンプリングすればよいのでは？ メモリ的に効率がかなり悪い 実は，学習序盤では大量のサンプル点が必要となる しかし，学習が進むにつれてサンプル点の数は減少する 上限分だけ静的に確保する実装が考えられるが，無駄が多くなってしまう メモリ効率を得るために，実行時間を多少犠牲にすることを考える 1回目のサンプリングを行い，そのレイのサンプル点の数を得る 2回目に全く同じサンプリングを行い，イイ感じにメモリ領域に格納していく 49
53. 実装的にはこんなことをする 53
54. NeRF: まとめ (1): カメラから画像のピクセルに対応する方向にレイ（光線）を飛ばす (2): レイの経路上の点をOccupancy Gridを参照しつつサンプリング (3): サンプル点の座標，そしてレイの向きをニューラルネットワーク(NN)に入力 (4): NNを実行し，サンプル点における「色」と「密度」を計算 (5): 得られたサンプル点の情報からボリュームレンダリング方程式を計算 (6): 計算結果と対応するピクセルの色を比較し，誤差を計算 (7): ボリュームレンダリングの処理を逆方向に行う (8): NNを誤差逆伝播 (9): NN最適化 (10): Occupancy Gridの最適化（16イテレーションごと） 54
55. Neural Network 機械学習で使用されている，神経回路を云々（MLP） 今回使用するのは線形層（全結合層）と活性化層の繰り返し 難しく考えずに単なるクソデカ関数だと考える． … … …… … … 55
56. 線形層（全結合層） 入力ベクトルの基底の線形結合を出力する W: パラメーター行列 / b: バイアス項 / x: 入力ベクトル / y: 出力ベクトル 特に，今回はバイアス項を使用しないので 𝒚 = 𝑊𝒙 + 𝒃 𝒚 = 𝑊𝒙 x1 x2 W11 W12 W21 W22 W31 W32 y1 y2 y3 W11*x1 + W12*x2 W21*x1 + W22*x2 W31*x1 + W32*x2 = = 56
57. 活性化層 線形層（全結合層）の出力ベクトルの各要素に対して，非線形な関数を与える 色んな関数があるが，今回はSigmoid(Logistic)とReLUが主 Sigmoid(Logistic): ReLU 𝑦 = 1 1 + 𝑒−𝑥 𝑦 = max(0, 𝑥) 0 0.5 1 -5 -4 -3 -2 -1 0 1 2 3 4 5 0 0.5 1 -1 0 1 57
58. 逆伝播 計算で簡単に示せる 線形層（全結合層） Sigmoid (Logistic) ReLU 𝜕𝐿 𝜕𝒙𝑇 = 𝜕𝐿 𝜕𝒚𝑇 𝑊 𝜕𝐿 𝜕𝑊 = 𝜕𝐿 𝜕𝒚 𝒙𝑇 𝑑𝐿 𝑑𝑥 = 𝑑𝐿 𝑑𝑦 𝑦(1 − 𝑦) 𝑑𝐿 𝑑𝑥 = ൞ 𝑑𝐿 𝑑𝑦 (𝑥 > 0) 0 (𝑒𝑙𝑠𝑒) 𝒚 = 𝑊𝒙 𝑦 = 1 1 + 𝑒−𝑥 𝑦 = max(0, 𝑥) 58
59. 最適化処理 基本は勾配降下法 Adam: 勾配降下法の発展形の最適化関数（かなり強力） 59 𝑤 = 𝑤 − 𝜂 𝜕𝐿 𝜕𝑤 𝑚 = 𝛽1 𝑚 + 1 − 𝛽1 𝜕𝐿 𝜕𝑤 𝑣 = 𝛽2 𝑣 + 1 − 𝛽2 𝜕𝐿 𝜕𝑤 2 ෝ 𝑚 = 𝑚 (1 − 𝛽1 𝑡) ො 𝑣 = 𝑣 (1 − 𝛽2 𝑡) 𝑤 = 𝑤 − 𝜂 ෝ 𝑚 ො 𝑣 + 𝜖
60. バッチ処理 実際の入力データは大量にある 入力データの塊を一部取り出して学習させることが多い 活性化層は要素ごとの処理なのでやることは同じ 線形層（全結合層）は？ x Neural Network y x Neural Network x x x x y y y y y 60
61. バッチ処理: 線形層（全結合層） 実は，ベクトルが行列となるだけ 逆伝播も同じくして WX = 𝑊 𝒙𝟏 𝒙𝟐 𝒙𝟑 … = 𝑊𝒙1 𝑊𝒙2 𝑊𝒙𝟑 … = 𝒚𝟏 𝒚𝟐 𝒚𝟑 … = 𝑌 𝜕𝐿 𝜕𝑋𝑇 = 𝜕𝐿 𝜕𝑌𝑇 𝑊 𝜕𝐿 𝜕𝑊 = 𝜕𝐿 𝜕𝑌 𝑋𝑇 y y x x y w w w w w w = y y x x y w w w w w w = y y y y y y y y y y y y x x x x x x x x 61
62. Neural Networkの実装 花子さん: なんか機械学習やりたい気分！ 太郎くん: いいね．PyTorchとかcuDNNとか使ったらいいんじゃない？ 花子さん: 私に既存フレームワークを使えっていうの？ふざけないで！ 太郎くん : ！？ 花子さん: 何よその顔．フルスクラッチ実装が正義でしょ！ 太郎くん？: それはとても素晴らしい 62
63. Fully Fused MLP [NVIDIA 2021] NVIDIA GPUに最適化されたMLPの実装 shared memoryを最大活用し，global memoryとのやり取りを最小限に抑える 各Blockが128バッチを処理する設計とする 入力層～出力層の次元は大きくても128次元とする 入力層から流れていくベクトルをshared memoryに載せたまま処理できる！ bank conflictという現象を避けるためにゼロパディングすることもある 線形層（全結合層）の処理にTensorコアを使用する 行列演算にハードウェア最適化の恩恵を受けられる 63
64. Tensorコアを使用した行列演算 大きな行列の演算は，各々を構成する小行列の演算を繰り返すことで行える 下の図では64x64と64x128の行列を乗算している（1マス16次元） この行列演算は結局，16x16の行列同士の掛け算の繰り返しで完了する 色が違う行列演算は並列化可能 1Blockにおいて4並列可能 64
65. Tensorコアを使用した行列演算 WMMA(Wave Matrix Multiply-accumulate) のAPIでTensorコアを使用できる このAPIで行える行列演算は次元が決められている そのうちの一つに，16x16行列同士の演算がある 設計としては16x16の行列の要素をWarpに含まれる32スレッドにばらしている バラされた要素はfragmentという構造体として持たれる 乗行列，被乗行列，出力行列はすべて構造体が異なる (matrix_a, matrix_b, accumulator) そして，それらは互換性がないことに注意 1Blockで4並列で計算できる場合 1Blockに4Warp使用する（つまりBlockSize = {32, 4, 1}） 65
66. Tensorコアを使用した行列演算 load_matrix_syncで通常のfloat配列などからロードできる 66
67. Tensorコアを使用した行列演算 mma_syncで行列の乗算を行える: C = AB + C 未初期化のfragmentはゴミ値が入っているので注意(fill_fragmentで0埋めを行う) 67
68. Tensorコアを使用した行列演算 store_matrix_syncでfragmentを通常のfloat配列などに保存できる 68
69. 私の実装では 学習処理 隠れ層の数と1epoch処理時間(バッチサイズ128) バッチサイズと1epoch処理時間(隠れ層4層) 69
70. 私の実装では 推論処理 隠れ層の数と1epoch処理時間(バッチサイズ128) バッチサイズと1epoch処理時間(隠れ層4層) 70
71. エンコーダー NeRFではPosition(3D), Direction(3D)を入力としていた たった6次元のデータ量から精度よく推定することは難しい (1) 入力の次元を上げよう！ (2) 入力ベクトルの基底を変更して次元数を増やそう エンコーダーの出番 Position Direction Neural Network Volume Parameter 71
72. Frequency Encoding 原著NeRFが使用していたエンコーダー 入力pを高次元のベクトルγにエンコードする． 1次元のベクトルはこれにより2L次元となる 𝛾 𝑝 = (sin 20𝜋𝑝 , cos 20𝜋𝑝 , sin 21𝜋𝑝 , cos 21𝜋𝑝 , … , sin 2𝐿−1𝜋𝑝 , cos(2𝐿−1𝜋𝑝)) 72
73. 原著NeRF - Network Green: Input Red: Output Blue: Fully-Connected Layer (Number is the dimension) Array: Black = ReLU / Orange = None/ Dashed = Sigmoid NeRF paper – Fig. 7 73
74. 原著NeRF 重い！！！！！ 学習: 数日 描画: 30s/frame (NVIDIA V100) NeRF paper – Fig. 7 74
75. Multiresolution Hash Encoding Positionをエンコードする 座標が存在する領域を「様々な解像度で解釈する」 これによりNeRFが非常に高速化された 元々は学習に数日かかっていたのが数分まで短縮された」 （数学的な意味は私はよく分からない） 75
76. Multiresolution Hash Encoding (2D) 座標が存在する領域を様々な解像度に分割する 76
77. Multiresolution Hash Encoding (2D) 座標が存在する領域を様々な解像度に分割する 分割されたどの小正方形内部にサンプル点があるかを求める 小正方形内部のどこにサンプル点があるかを求める 77
78. Multiresolution Hash Encoding (2D) 各格子点は特徴ベクトルを持つ インデックスは頂点座標をハッシュ関数に入れることで求まる それをもとにテーブル上の特徴ベクトルにアクセス サンプル点における特徴ベクトルをバイリニア補完で求める (3Dの時のハッシュ関数) 78
79. Multiresolution Hash Encoding (2D) 各レベルごとにF次元の特徴ベクトルが得られる 全てを結合してLF次元とし，これをエンコード結果とする 79
80. Multiresolution Hash Encoding (3D) 80
81. MHEの誤差逆伝播 81
82. Spherical Harmonic Encoding 球面調和関数を用いてDirectionをエンコードする． ラプラス方程式をうみゃみゃして得られる次の式は 完全正規直交性を持つ．（記事で正規直交性は示しましたが完全性は示せてません） つまり，球面上で定義される連続で滑らかな関数は球面調和関数系の基底で展開できる． Directionは半径1の球面上の点とみることが出来る（長さ1のベクトルであるため） Directionを(x, y, z)から球面調和関数の基底へ変換できる！ 82
83. The Network (𝑥, 𝑦, 𝑧) (𝐷𝑖𝑟𝑋, 𝐷𝑖𝑟𝑌, 𝐷𝑖𝑟𝑍) Density Color MHE FC ReLU FC ReLU FC ReLU FC ReLU FC SH Projection Input Exp Sigmoid V[0] [15:0] [31:16] 3 32 64 16 16 3 64 64 3 83
84. 最後にもういちど整理 (1): カメラから画像のピクセルに対応する方向にレイ（光線）を飛ばす (2): レイの経路上の点をOccupancy Gridを参照しつつサンプリング (3): サンプル点の座標，そしてレイの向きをニューラルネットワーク(NN)に入力 (4): NNを実行し，サンプル点における「色」と「密度」を計算 (5): 得られたサンプル点の情報からボリュームレンダリング方程式を計算 (6): 計算結果と対応するピクセルの色を比較し，誤差を計算 (7): ボリュームレンダリングの処理を逆方向に行う (8): NNを誤差逆伝播 (9): NN最適化 (10): Occupancy Gridの最適化（16イテレーションごと） 84
85. 参考文献 NeRF: Representing Scenes as Neural Radiance Fields for View Synthesis [Ben Mildenhall et al. ECCV2020] Instant Neural Graphics Primitives with a Multiresolution Hash Encoding [Thomas Müller et al. SIGGRAPH2022] Optical Models for Direct Volume Rendering [Nelson Max. 1995] memoRANDOM ボリュームレンダリング方程式 (Volume Rendering Equation) 1 URL: https://rayspace.xyz/CG/contents/VLTE1/ NVIDIA Developer Blog, Thinking Parallel, Part III: Tree Construction on the GPU https://developer.nvidia.com/blog/thinking-parallel-part-iii-tree-construction-gpu/ Ushioさんのredpillの実装と解説スライド https://github.com/Ushio/redpill Optim Tech Blog. Instant NeRF の心臓、Multiresolution Hash Encoding をシンプルに実装しつつ2次元画像で試してみる https://tech-blog.optim.co.jp/entry/2022/04/22/100000 Real-time Neural Radiance Caching for Path Tracing tiny-cuda-nnのgitリポジトリ https://github.com/NVlabs/tiny-cuda-nn Tensorコアを使ってみた https://proc-cpuinfo.fixstars.com/2018/10/tensorcore/ 85
86. ご清聴ありがとうございました 大体の実装は私のはてなブログに載せてありますので，実装に興味がある方はそちらを参照 してください．https://tetrisyoshi.hatenablog.com/about 質問等あればどうぞ．Slackでも受け付けております． フルスクラッチはいいぞ けものフレンズはいいぞ 86