はじめに
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のほうが勧められているようです。1