JetsonとUnifiedMemoryでOpenCVをもっと高速に使う

# はじめに
CUDAにはあまり注目されない(？)機能としてUnified Memoryというものがあります。
メモリ空間を統一して扱うことができ、CPU⇔GPUのデータ転送を簡素化することができます。
必要に応じてデータが自動的にCPU/GPUメモリ領域に転送される機能です。

一方、JetsonはCPU/GPUで一つのDRAMを共有して使用します。
共有DRAM上にデータがあるならデバイス間のデータコピーはもったいないな、と思いませんか？
ここでUnified Memoryを使用すると、なんと「コピーコストなしでCPUでもGPUでも読み書きできるメモリ」が手に入ります。
実際に試してみましょう。

# 環境
* JetsonTX2(Linux nvidia-desktop 4.9.140-tegra)
* OpenCV 4.1.1

# 使い方
Unified Memoryを使うために、専用のmallocがあります。

```
    cudaMallocManaged(&ptr,size)
```
freeは、普通のCUDAメモリと同様に扱えます。

```
    cudaFree(ptr);
```

`cudaMallocManaged`で確保した領域を`cv::Mat`に食わせます。
同様に、確保した領域を`cv::cuda::GpuMat`にも食わすことができます。
自動的にメモリ解放されたりはしないので、使い終わったらfreeしてください。

```cpp
    int width = 1920, height = 1080;

    char* managed_ptr;
    cudaMallocManaged(&managed_ptr, width*height*3);

    cv::Mat cpu_mat(height, width, CV_8UC3, managed_ptr);
    cv::cuda::GpuMat gpu_mat(height, width, CV_8UC3, managed_ptr);
```

これで、`cpu_mat`と`gpu_mat`は共通した中身を持つようになります。
データの中身はGPUドライバが「いい具合」に扱ってくれます。
`gpu_mat`に対して`upload`/`download`のような操作をする必要はありません。便利ですね。

# 使ってみる
`cv::cuda`名前空間に定義されている関数ならば高速化することができます。
試しにリサイズ処理をやってみましょう。

#### CPU
```cpp

#include <opencv2/opencv.hpp>
#include <iostream>
#include <chrono>

int main()
{
	const int loopcount = 10000;

	cv::Mat src = cv::imread("src.png");
	cv::Mat dst(cv::Size(300, 300), CV_8UC3);

	auto st = std::chrono::high_resolution_clock::now();
	for(int i = 0; i < loopcount; ++i)
	{
		cv::resize(src, dst, { 300,300 });
	}
	auto ed = std::chrono::high_resolution_clock::now();

	double dur = std::chrono::duration_cast<std::chrono::microseconds>(ed - st).count() / 1000. / loopcount;
	std::cout << dur << "[ms]" << std::endl;

	cv::imwrite("dst.png", dst);

    return 0;
}
```

#### GPU (Unified Memoryを使わない)
```cpp

#include <opencv2/opencv.hpp>
#include <opencv2/core/cuda.hpp>
#include <iostream>
#include <chrono>

int main()
{
	const int loopcount = 10000;

	cv::Mat src = cv::imread("src.png");
	cv::Mat dst(cv::Size(300, 300), CV_8UC3);
	cv::cuda::GpuMat gpu_src(cv::Size(1920, 1080), CV_8UC3);
	cv::cuda::GpuMat gpu_dst(cv::Size(300, 300), CV_8UC3);

	auto st = std::chrono::high_resolution_clock::now();
	for(int i = 0; i < loopcount; ++i)
	{
		gpu_src.upload(src);
		cv::cuda::resize(gpu_src, gpu_dst, { 300,300 });
		gpu_dst.download(dst);
	}
	auto ed = std::chrono::high_resolution_clock::now();

	double dur = std::chrono::duration_cast<std::chrono::microseconds>(ed - st).count() / 1000. / loopcount;
	std::cout << dur << "[ms]" << std::endl;

	cv::imwrite("dst.png", dst);

	return 0;
}
```
メモリ転送のために`upload`/`download`が必要です。

#### GPU (Unified Memoryを使う)
```cpp
#include <opencv2/opencv.hpp>
#include <opencv2/core/cuda.hpp>
#include <cuda_runtime.h>
#include <iostream>
#include <chrono>
#include <cstring>

int main()
{
	const int loopcount = 10000;

	cv::Mat src = cv::imread("src.png");

	char* managed_src;
	char* managed_dst;

	cudaMallocManaged(&managed_src, 1920 * 1080 * 3);
	cudaMallocManaged(&managed_dst, 300 * 300 * 3);

	cv::Mat cpu_mat_src(cv::Size(1920,1080), CV_8UC3, managed_src);
	cv::Mat cpu_mat_dst(cv::Size(300, 300), CV_8UC3, managed_dst);
	cv::cuda::GpuMat gpu_mat_src(cv::Size(1920, 1080), CV_8UC3, managed_src);
	cv::cuda::GpuMat gpu_mat_dst(cv::Size(300, 300), CV_8UC3, managed_dst);

	// 初回のみコピー
	memcpy(managed_src, src.data, 1920 * 1080 * 3);

	auto st = std::chrono::high_resolution_clock::now();
	for(int i = 0; i < loopcount; ++i)
	{
		cv::cuda::resize(gpu_mat_src, gpu_mat_dst, { 300,300 });
	}
	auto ed = std::chrono::high_resolution_clock::now();

	double dur = std::chrono::duration_cast<std::chrono::microseconds>(ed - st).count() / 1000. / loopcount;
	std::cout << dur << "[ms]" << std::endl;

	cv::imwrite("dst.png", cpu_mat_dst);

	cudaFree(managed_src);
	cudaFree(managed_dst);

	return 0;
}
```
`imread`では直接Unified Memoryに書き込めないので、事前に一回コピーする必要があります。

#### 結果

計測では`sudo nvpmodel -m 0; sudo jetson_clocks`を事前に使用しています。

||JetsonTX2|
|---|---|
|CPU      |2.214ms   |
|GPU (noUM)   | 2.523ms |
|GPU (withUM)  |0.479ms   |

コピーコストは存外に重いことがわかります。
JetsonNanoでも有効なテクニックですので試してはどうでしょうか。

### その他
似たようなことを実現できる機能にZeroCopy Memoryというものがあります。
メモリ領域をページロックしておいて、GPUによる直接メモリアクセスを可能にする技術です。
ただ、GPUドライバによる自動同期があること、キャッシュが効くことからUnified Memoryのほうが勧められているようです。[^1]
[^1]: [Zero-Copy and Managed memory on Jetson](https://devtalk.nvidia.com/default/topic/932957/zero-copy-and-managedmemory-on-jetson)