はじめに

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してください。

    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


#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を使わない)


#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を使う)

#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のほうが勧められているようです。¹

Zero-Copy and Managed memory on Jetson ↩