CUDA対応を有効化したlibrealsense2を使う

## はじめに

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
  - Power mode: MAXN ※詳細は<https://docs.nvidia.com/jetson/archives/r35.1/DeveloperGuide/text/SD/PlatformPowerAndPerformance/JetsonOrinNxSeriesAndJetsonAgxOrinSeries.html#supported-modes-and-power-efficiency>を参照
- librealsense 2.53.1

## RGB-DキャプチャのCPU使用率

今回、効果測定のためにlibrealsense2のサンプルプログラムの1つである[rs-align](https://github.com/IntelRealSense/librealsense/tree/v2.53.1/examples/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など）
  - <https://github.com/IntelRealSense/librealsense/blob/v2.53.1/src/cuda/cuda-conversion.cu>
- pixelからpointへのdeproject処理
  - <https://github.com/IntelRealSense/librealsense/blob/v2.53.1/src/cuda/cuda-pointcloud.cu>
- RGBとdepthのalign処理
  - <https://github.com/IntelRealSense/librealsense/blob/v2.53.1/src/proc/cuda/cuda-align.cu>

### 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）に向いたメモリ確保になっていない
  - <https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html>参照
- 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回呼ばれていることがわかります。

![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/31158/435f1b0b-7e66-dfce-b913-a8c17efeb800.png)

#### コールスタック（変更後）

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`が呼ばれていないことがわかります。

![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/31158/ab6c5fba-c38d-22ab-c395-cb21e9c6a91d.png)

#### CPU使用率、処理時間

「cudaMalloc呼び出しを減らす」対応を入れることでCPU使用率、処理時間が改善できていることがわかります。

|ビルドオプション、高速化施策|CPU使用率\[%\]|
|---|---:|
|WITH_CUDA=OFF|112|
|WITH_CUDA=ON|31|
|WITH_CUDA=ON、<br>cudaMalloc呼び出しを減らす|29|

|ビルドオプション、高速化施策|処理時間\[ms\]（デバイス、ホスト間のデータコピー時間含む）|
|---|---:|
|WITH_CUDA=OFF|3.644|
|WITH_CUDA=ON|1.952|
|WITH_CUDA=ON、<br>cudaMalloc呼び出しを減らす|1.122|

## Jetsonデバイス（iGPU）向けのメモリ確保にする

次にJetsonデバイス（iGPU）向けのメモリ確保にできないかを考えてみます。<https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html#memory-management>にあるメモリ構成図を以下に示します。

![](https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/_images/dGPU-connected-tegra-memory-system.png)

また、同ページにある

> 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](https://fareastprogramming.hatenadiary.jp/entry/2016/10/19/203637)が使えそうに思えます。

ただし、librealsense2では画像バッファなどのフレームデータを[rs2::frame](https://intelrealsense.github.io/librealsense/doxygen/classrs2_1_1frame.html)クラスのメンバとして管理しています。この`rs2::frame`クラスはCUDA実装以外でも参照されることと、画像データ以外も格納される可能性があることから、このクラスをNVIDIA GPU向けに変更してしまうと影響範囲が大きすぎるという問題があります。すぐに良い手が思い付かなかったのでこの部分は一旦保留としています。

## 今後の予定

### 色変換処理

- 詳細設計
- 動作確認（YUV422->RGB以外も含む）
- librealsense2へのMR作成

### pixelからpointへのdeproject処理

- 詳細設計
- 動作確認（YUV422->RGB以外も含む）
- librealsense2へのMR作成