はじめに
IntelはRealSenseデバイスのSDKであるlibrealsense2(https://github.com/IntelRealSense/librealsense)をOSSとして公開しています。本記事ではこのlibrealsense2のCUDA対応を紹介します。また、高速化の取り組みを少しだけ行ったのでそちらも合わせて紹介します。
動作確認環境
ハードウェア
- reComputer J4012(Jetson Orin NX 16GB)
- Intel RealSense D455
ソフトウェア
- Jetson Linux 35.3.1
- librealsense 2.53.1
RGB-DキャプチャのCPU使用率
今回、効果測定のためにlibrealsense2のサンプルプログラムの1つであるrs-alignを用いました。
reComputer J4012(Jetson Orin NX 16GB)上でrs-align（CPUのみ）を実行したところ、CPU使用率が112%となりました。このままでもRealSenseデバイスを1台接続する分には問題になりにくいのですが、同時に複数台接続するようなケースだと単純に台数分だけキャプチャ処理のCPU使用率が上がっていくので問題になってきます。
librealsense2のCUDA対応
librealsense2はhttps://github.com/IntelRealSense/librealsense/issues/6302でCUDA対応が既に行われています。https://dev.intelrealsense.com/docs/build-configurationにあるようにCMakeオプションでBUILD_WITH_CUDA=ONとしてlibrealsense2をソースビルドすることでCUDA実装を有効にすることができます。
librealsense2で対応されているCUDAカーネルは以下にあります。
- https://github.com/IntelRealSense/librealsense/tree/v2.53.1/src/cuda
- https://github.com/IntelRealSense/librealsense/tree/v2.53.1/src/proc/cuda
具体的には以下の処理がCUDA対応されています。
- 色変換処理（YUV422->RGBなど）
- pixelからpointへのdeproject処理
- RGBとdepthのalign処理
CPU使用率、処理時間
reComputer J4012(Jetson Orin NX 16GB)上でrs-alignを実行した時のCPU使用率を下表にまとめます。
CUDA対応によりCPU負荷が低減できていることがわかります。
| ビルドオプション、高速化施策 | CPU使用率[%] |
|---|---|
| WITH_CUDA=OFF | 112 |
| WITH_CUDA=ON | 31 |
また、reComputer J4012(Jetson Orin NX 16GB)上でrs-alignを実行した時のYUV422->RGB色変換処理の処理速度（デバイス、ホスト間のデータコピー時間含む）を下表にまとめます。この結果からCUDA対応により処理時間が改善できていることがわかります。
| ビルドオプション、高速化施策 | 処理時間[ms]（デバイス、ホスト間のデータコピー時間含む） |
|---|---|
| WITH_CUDA=OFF | 3.644 |
| WITH_CUDA=ON | 1.952 |
Jetsonデバイス向けにおけるlibrealsense2の課題
Intel主導で開発していることもあり、開発プライオリティから考えて仕方がないところもありますが、現行のlibrealsense2をJetsonデバイス上で使うときにいくつか気になる点があります。
- 一部CUDA実装にて毎フレームで
cudaMallocをコールしている- 画像データなど一度メモリを確保して使い回せるものがあるので毎フレーム呼び出す必要がない
- Jetsonデバイス（iGPU）に向いたメモリ確保になっていない
- Arm CPU向けの高速化対応が少ない
- Intel CPU向け（SSE3、AVX2実装）の高速化対応が入っている
- 一方でArm CPU向けの高速化対応はコンパイルオプション付与くらい https://github.com/IntelRealSense/librealsense/blob/v2.53.1/CMake/unix_config.cmake#L10-L15
高速化施策（YUV422->RGB色変換処理）
ここでは処理がわかりやすいYUV422->RGB色変換処理を題材にして高速化を行なってみます。
cudaMalloc呼び出しを減らす
キャプチャ設定（カラースペース、解像度）が決まれば色変換処理（YUV422->RGBなど）で使われるメモリ使用サイズは固定となります。そのため、一度メモリを確保して使い回すことでcudaMalloc呼び出しを減らすことができます。
コールスタック（変更前）
変更前のコールスタックは以下の通りです。ここからもわかるようにcudaMallocを毎フレーム呼び出しています。
librealsense::unpack_yuy2
rscuda::unpack_yuy2_cuda
rscuda::unpack_yuy2_cuda_helper
alloc_dev
cudaMalloc ... srcのメモリ確保
cudaMemcpy(HostToDevice)
alloc_dev
cudaMalloc ... dstのメモリ確保
kernel_unpack_yuy2_rgb8_cuda ... 色変換処理のCUDAカーネル
cudaGetLastError
cudaDeviceSynchronize
cudaMemcpy(DeviceToHost)
Nsight Systemsで確認したタイムラインは以下の通りです。フレームの処理において、src、dstでcudaMallocが合計2回呼ばれていることがわかります。
コールスタック（変更後）
1フレーム目で色変換に用いるメモリを確保しておき、それを使い回すことで2フレーム目以降はcudaMallocを呼び出す必要がないので以下のような呼び出しになります。
librealsense::unpack_yuy2
rscuda::unpack_yuy2_cuda
rscuda::unpack_yuy2_cuda_helper
cudaMemcpy(HostToDevice)
kernel_unpack_yuy2_rgb8_cuda ... 色変換処理のCUDAカーネル
cudaGetLastError
cudaDeviceSynchronize
cudaMemcpy(DeviceToHost)
Nsight Systemsで確認したタイムラインは以下の通りです。2フレーム目以降はcudaMallocが呼ばれていないことがわかります。
CPU使用率、処理時間
「cudaMalloc呼び出しを減らす」対応を入れることでCPU使用率、処理時間が改善できていることがわかります。
| ビルドオプション、高速化施策 | CPU使用率[%] |
|---|---|
| WITH_CUDA=OFF | 112 |
| WITH_CUDA=ON | 31 |
| WITH_CUDA=ON、 cudaMalloc呼び出しを減らす |
29 |
| ビルドオプション、高速化施策 | 処理時間[ms]（デバイス、ホスト間のデータコピー時間含む） |
|---|---|
| WITH_CUDA=OFF | 3.644 |
| WITH_CUDA=ON | 1.952 |
| WITH_CUDA=ON、 cudaMalloc呼び出しを減らす |
1.122 |
Jetsonデバイス（iGPU）向けのメモリ確保にする
次にJetsonデバイス（iGPU）向けのメモリ確保にできないかを考えてみます。https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html#memory-managementにあるメモリ構成図を以下に示します。
また、同ページにある
In Tegra, device memory, host memory, and unified memory are allocated on the same physical SoC DRAM.
という記載に着目します。このことから、https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html#memory-selectionにあるようにJetsonデバイス向けの対策としては、Unified memoryやZero Copyが使えそうに思えます。
ただし、librealsense2では画像バッファなどのフレームデータをrs2::frameクラスのメンバとして管理しています。このrs2::frameクラスはCUDA実装以外でも参照されることと、画像データ以外も格納される可能性があることから、このクラスをNVIDIA GPU向けに変更してしまうと影響範囲が大きすぎるという問題があります。すぐに良い手が思い付かなかったのでこの部分は一旦保留としています。
今後の予定
色変換処理
- 詳細設計
- 動作確認（YUV422->RGB以外も含む）
- librealsense2へのMR作成
pixelからpointへのdeproject処理
- 詳細設計
- 動作確認（YUV422->RGB以外も含む）
- librealsense2へのMR作成