Computer Vision Seminar 1/コンピュータビジョンセミナーvol.1 OpenCV活用

Copyright © Fixstars Group
コンピュータビジョンセミナーvol.1
OpenCV活用
OpenCVでCUDAを活用するためのGpuMat解説

View Slide

本日のAgenda
● はじめに
● フィックスターズのご紹介
● OpenCVでCUDAを活用するためのGpuMat解説
○ GpuMatの概要
○ GpuMatと自作CUDAカーネルの連携
○ GpuMatとNPP連携
○ cv::cuda::Streamについて
○ cv::cuda::BufferPoolについて
○ GpuMat Tips
○ GpuMat注意ポイント
● Q&A / 告知
2

View Slide

はじめに

View Slide

● 弊社でサービス展開しているコンピュータビジョン領域において、
複数回に渡る技術セミナーの開催を計画しています
● 今回は、コンピュータビジョン領域で最も利用されているフレームワーク OpenCVを
題材とし、CUDAを活用する上で重要なデータ構造 GpuMat について解説します
● こんな方に向いています
○ OpenCVと連携させるCUDAカーネルを自作したい
○ GPUを活用して処理を高速化したい
本講演の位置づけ
4

View Slide

発表者紹介
冨田明彦
ソリューションカンパニー
執行役員
2008年に入社。金融、医療業界において、ソ
フトウェア高速化業務に携わる。その後、新規
事業企画、半導体業界の事業を担当し、現職。
5
吉村康弘
Fixstars Autonomous Technologies
リードエンジニア
2015年に入社。主に画像処理、コンピュータビ
ジョンのアルゴリズム開発やCUDA高速化業務
を担当。

View Slide

フィックスターズの
ご紹介

View Slide

フィックスターズの強み
コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団
ハードウェアの知見アルゴリズム実装力各産業・研究分野の知見
7
目的の製品に最適なハードウェアを見抜き、
その性能をフル活用するソフトウェアを開
発します。
ハードウェアの特徴と製品要求仕様に合わ
せて、アルゴリズムを改良して高速化を実
現します。
開発したい製品に使える技術を見抜き、実
際に動作する実装までトータルにサポート
します。

View Slide

開発サービス提供分野
8
半導体
自動車
産業機器
生命科学
金融
●NAND型フラッシュメモリ向け
ファームウェア開発
●次世代AIチップの開発環境基盤
●自動運転の高性能化、実用化
●次世代パーソナルモビリティの
研究開発
●Smart Factory実現への支援
●マシンビジョンシステムの高速化
●ゲノム解析の高速化
●医用画像処理の高速化
●AI画像診断システムの研究開発
●デリバティブシステムの高速化
●HFT(アルゴリズムトレード)の高速化

View Slide

サービス領域
様々な領域でソフトウェア開発サービスを提供しています。大量データの高速処理は、
お客様の製品競争力の源泉となっています。
9
組込み高速化
画像処理・アルゴリズム
開発
分散並列システム開発
GPU向け高速化
FPGAを活用した
システム開発
量子コンピューティング
AI・深層学習
自動車向け
ソフトウェア開発
フラッシュメモリ向けフ
ァームウェア開発

View Slide

自動車向けソフトウェア開発
アルゴリズム開発から量産車ターゲット向けの高速化まで、
自動運転の実現に向けた統合的な技術開発を行っています。
ご支援内容

View Slide

画像処理アルゴリズム開発
高速な画像処理需要に対して、経験豊富なエンジニアが
責任を持って製品開発をご支援します。
お客様の課題
高度な画像処理や深層学習等のアルゴリズム
を開発できる人材が社内に限られている
機能要件は満たせそうだが、ターゲット機器
上で性能要件までクリアできるか不安
製品化に結びつくような研究ができていない
ご支援内容
深層学習ネットワーク精度の改善
様々な手法を駆使して深層学習ネットワークの精度を改善
論文調査・改善活動
論文調査から最先端の手法の探索
性能向上に向けた改善活動を継続
アルゴリズム調査・改変
課題に合ったアルゴリズム・実装手法を調査
製品実装に向けて適切な改変を実施

View Slide

GPU向け高速化
高性能なGPUの本来の性能を十分に引き出し、
ソフトウェアの高速化を実現します。
お客様の課題
GPUで計算してみたが期待した性能が出ない
GPU/CPUを組み合わせた全体として最適な
設計がしたい
ご支援内容
GPU高速化に関するコンサルティング
CPU・GPU混在環境でのシステム設計
アルゴリズムのGPU向け移植
GPUプログラム高速化
継続的な精度向上
原価を維持したまま機能を追加するため、も
う少し処理を速くしたい
品質確保のため、精度を上げたく演算量は増
えるが性能は維持したい

View Slide

GpuMatの概要

View Slide

GpuMatの概要
● OpenCV[1]
は、NVIDIA GPUに処理をオフロードするためのデータ構造として
GpuMatクラス
[2]
を提供している
14
cudaモジュール
coreモジュール
GpuMat
NVIDIA GPU
処理をオフロード
[1] https://opencv.org/
[2] https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html

View Slide

GpuMatの概要
● OpenCVのcudaモジュールの処理フロー概要は以下の通り
○ NPP（NVIDIA Performance Primitives）については後述
15
チャンネル数、depthチェック
NPPを使う？
NPPの関数を呼ぶ
OpenCV同梱の
CUDAカーネルを呼ぶ
OpenCVで
サポート？
Y
Y
N
N
エラー

View Slide

GpuMatの概要
● OpenCVのGpuMatを使うには
○ NVIDIA CUDA ToolKitをインストールする
○ OpenCVをソースコードからインストールする際、以下のCMakeオプションをONにする
16
CMakeオプション意味
WITH_CUDA OpenCVでCUDAを使った実装を有効にする
WITH_CUFFT OpenCVでcuFFTを使った実装を有効にする
WITH_CUBLAS OpenCVでcuBLASを使った実装を有効にする

View Slide

GpuMatの概要
● サンプルコード
17
cv::Mat src(cv::Size(3840, 2160), CV_8UC3, cv::Scalar(0, 0, 0));
cv::cuda::GpuMat d_src;
// GPUに転送
d_src.upload(src);
// GPUで処理
cv::cuda::GpuMat d_gray, d_bin;
cv::cuda::cvtColor(d_src, d_gray, cv::COLOR_BGR2GRAY);
cv::cuda::threshold(d_gray, d_bin, 200, 255, cv::THRESH_BINARY);
// ホストに転送
cv::Mat bin;
d_bin.download(bin);
cv::cuda::GpuMat d_src(src);
と書くのでもよい

View Slide

GpuMatの概要
● cudaモジュール概要
○ 代表的なモジュールは以下の通り
18
モジュール名概要
cudaarithm 行列操作
cudabgsegm Background Segmentation
cudacodec Video Encoding/Decoding
cudafeatures2d Feature Detection and Description
cudafilters フィルタ処理
cudaimgproc 色変換、ヒストグラム、コーナー検出など

View Slide

GpuMatの概要
● cudaモジュール概要
○ 代表的なモジュールは以下の通り
19
モジュール名概要
cudalegacy
レガシーなアルゴリズム
（オプティカルフロー、背景分離など）
cudaobjdetect 物体検出
cudaoptflow オプティカルフロー（Sparse、Dense）
cudastereo ステレオマッチング
cudawarping リサイズ、アフィン変換など
cudev Device layer

View Slide

GpuMatと自作CUDA
カーネルの連携

View Slide

GpuMatと自作CUDAカーネルの連携
● GpuMatと自作CUDAカーネルは簡単に連携できる
21
coreモジュール
GpuMat
NVIDIA GPU
処理をオフロード
自作CUDAカーネル

View Slide

● GpuMatを入力としたCUDAカーネルを作るメリットは以下の通り
○ OpenCVのcudaモジュールと連携しやすい
■ cudaモジュールにあるアルゴリズムと自作CUDAカーネルを組み合わせて実行する等
22
coreモジュール
cudaモジュール
GpuMat

View Slide

● GpuMatを入力としたCUDAカーネルを作るメリットは以下の通り
○ 画像ファイルの入出力、画像表示をOpenCVで行うことができる
23
coreモジュール
imgcodecsモジュール
画像ファイル入出力
highguiモジュール
画像表示
GpuMat

View Slide

● cv::cuda::PtrStepSz
○ CUDAカーネル（デバイスコード）で画像データにアクセスが容易となるデータ型
■ 日本語情報だと https://zenn.dev/onihusube/articles/d5c671870564b2 の解説がとてもわ
かりやすい
○ cv::cuda::PtrStepSzは名前から類推できるようにstep、サイズ情報（rows、cols）にアクセス
できる
○ step情報にアクセスできるcv::cuda::PtrStepも用意されている
24

View Slide

● GpuMatにある画像バッファのアドレス参照
○ 代表的な方法は以下の通り
■ GpuMatクラスのptrメソッド
■ cv::cuda::PtrStepSz
25

View Slide

○ GpuMatクラスのptrメソッド
■ cv::cuda::GpuMat.ptr(y)
● T：データ型
● y：参照する列（デフォルトだとy=0）
26
cv::cuda::GpuMat src(cv::Size(320, 240), CV_8UC1);
// 画像バッファの先頭アドレスを取得
// uchar* psrc = src.data でも等価
uchar* psrc = src.ptr(0);
int width = src.cols;
int height = src.rows;
int step = src.step;
https://github.com/opencv/opencv/blob/4.6.0/modules/core/include/opencv
2/core/cuda.inl.hpp#L207-L211
uchar* GpuMat::ptr(int y)
{
CV_DbgAssert( (unsigned)y < (unsigned)rows );
return data + step * y;
}
ptrメソッドの内部実装
サンプルコード

View Slide

○ cv::cuda::PtrStepSz
27
__global__ void inversionGpu(const cv::cuda::PtrStepSz src,
cv::cuda::PtrStepSz dst)
{
const int x = blockDim.x * blockIdx.x + threadIdx.x;
const int y = blockDim.y * blockIdx.y + threadIdx.y;
if ((y >= 0) && (y < src.rows))
{
if ((x >= 0) && (x < src.cols))
{
uchar value = src(y, x);
dst(y, x) = 255 - value;
}
}
} CUDAカーネル
uchar value = src.ptr(y)[x];
dst.ptr(y)[x] = 255 - value;
と書くのでもよい

View Slide

● GpuMatを入力としたCUDAカーネルの作り方
○ 方法１：画像バッファおよびサイズ情報を引数にしたCUDAカーネルを実装する
○ 方法２：cv::cuda::PtrStepSzを引数にしたCUDAカーネルを実装する
28

View Slide

● サンプルコード（CUDAカーネル）
29
__global__ void inversionGpu
(
uchar* src, uchar* dst,
int width, int height, int step
)
{
if ((y >= 0) && (y < height))
{
if ((x >= 0) && (x < width))
{
uchar val = src[y*step + x];
dst[y*step + x] = 255 - val;
}
}
} CUDAカーネル
画像バッファのポインタ、
width、height、ステップを引数で渡す
自前でアドレスを計算して読み書きする

View Slide

● サンプルコード（CUDAカーネル呼び出し）
30
void launchInversionGpu
(
cv::cuda::GpuMat& src,
cv::cuda::GpuMat& dst
)
{
const dim3 block(32, 32);
const dim3 grid(cv::cudev::divUp(dst.cols, block.x), cv::cudev::divUp(dst.rows, block.y));
// CUDAカーネル呼び出し
inversionGpu<<>>(src.ptr(0), dst.ptr(0), src.cols, src.rows, src.step);
// エラーチェック
CV_CUDEV_SAFE_CALL(cudaGetLastError());
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());
}
src.data、src.datastartと書くのでもよい

View Slide

● サンプルコード（CUDAカーネル）
31
(
const cv::cuda::PtrStepSz src,
cv::cuda::PtrStepSz dst
)
{
if((y >= 0) && (y < src.rows))
{
if((x >= 0) && (x < src.cols))
{
dst.ptr(y)[x] = (255 - src.ptr(y)[x]);
}
}
} CUDAカーネル
cv::cuda::PtrStepSzを引数で渡す
ptrメソッドで座標(x, y)のピクセルのアド
レスを参照し、読み書きする

View Slide

● サンプルコード（CUDAカーネル呼び出し）
32
void launchInversionGpu
(
cv::cuda::GpuMat& src,
cv::cuda::GpuMat& dst
)
{
const dim3 block(32, 32);
const dim3 grid(cv::cudev::divUp(dst.cols, block.x), cv::cudev::divUp(dst.rows, block.y));
// CUDAカーネル呼び出し
inversionGpu<<>>(src, dst);
// エラーチェック
CV_CUDEV_SAFE_CALL(cudaGetLastError());
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());
}
GpuMatクラスのインスタンスを渡す

View Slide

GpuMatとNPPの連携

View Slide

● NPP[3]
とは
○ NVIDIA Performance Primitivesのこと
○ 画像処理、信号処理等の各種アルゴリズムのCUDA実装ライブラリ
● NPPとOpenCVの実装は連携することができる
○ 以降でNPPの基本的な情報と連携方法を紹介します
34
[3] https://developer.nvidia.com/npp

View Slide

● NPPの機能概要（画像処理）
○ 色変換
○ フィルタ処理
○ Geometry Transforms（回転、反転、リサイズなど）
○ モルフォロジー変換（Erode、Dilate）
○ Statistical Operations（総和、最大値・最小値計算、ヒストグラム計算、インテグラルイメー
ジ作成など）
○ etc...
35
詳細は https://docs.nvidia.com/cuda/npp/modules.html 参照のこと

View Slide

● NPPで用いる基本的なデータ型
○ Nppという命名規則となっている
36
NPPで定義されるデータ型実際に用いられるデータ型
Npp8u 8-bit unsigned char
Npp8s 8-bit signed char
Npp16u 16-bit unsigned integer
Npp16s 16-bit signed integer
Npp32f
32-bit (IEEE) floating-point
number
Npp64f 64-bit floating-point number

View Slide

● サンプルコード（GpuMatとNPPの連携）
○ GpuMatクラスのインスタンスを入出力として、NPPのメディアンフィルタを実行する
37
// GpuMatクラスのインスタンス生成
cv::cuda::GpuMat d_src(src);
cv::cuda::GpuMat d_dst(dst);
// 中略
// stepの参照
Npp32s nSrcStep = d_src.step;
Npp32s nDstStep = d_dst.step;
// NPP API呼び出し
NppStatus status = nppiFilterMedian_8u_C1R(d_src.datastart, nSrcStep, d_dst.datastart, nDstStep,
roi, mask, anchor, d_median_filter_buffer);
8u：8-bit unsigned char
C1：1channel
R：Region-of-Interest(ROI)

View Slide

cv::cuda::Stream

View Slide

cv::cuda::Stream
● Streamとは
○ GPU処理のスケジュール管理の単位
39

View Slide

cv::cuda::Stream
● cv::cuda::Streamとは
○ CUDAのStreamをOpenCV内で実装されているCUDA実装で使うためにラップしたもの
○ OpenCVに実装されているCUDA実装でStreamを使ってスケジューリングする場合に用いる
○ OpenCV APIで明示的に指定しない場合、default stream(cv::cuda::Stream::Null())が用いられ
る
40
https://docs.opencv.org/4.6.0/db/d8c/group__cudaimgproc__color.html#ga48d0f208181d5ca370d8ff6b62cbe826

View Slide

cv::cuda::Stream
● サンプルコード（cv::cuda::Streamを明示的に指定しない）
41
cv::cuda::HostMem gray[2];
cv::cuda::GpuMat d_src[2], d_resize[2], d_gray[2];
// HtoD転送（default stream）
d_src[0].upload(src);
cv::cuda::resize(d_src[0], d_resize[0], cv::Size(), 2.0, 2.0, cv::INTER_LINEAR);
cv::cuda::cvtColor(d_resize[0], d_gray[0], cv::COLOR_BGR2GRAY, 0);
// DtoH転送（default stream）
d_gray[0].download(gray[0]);
// HtoD転送（default stream）
d_src[1].upload(src);
cv::cuda::resize(d_src[1], d_resize[1], cv::Size(), 2.0, 2.0, cv::INTER_LINEAR);
cv::cuda::cvtColor(d_resize[1], d_gray[1], cv::COLOR_BGR2GRAY, 0);
// DtoH転送（default stream）
d_gray[1].download(gray[1]);

View Slide

cv::cuda::Stream
● タイムライン（cv::cuda::Streamを明示的に指定しない）
42
Default streamのみが使われているデータ転送でブロッキングされている

View Slide

cv::cuda::Stream
● サンプルコード（cv::cuda::Streamを明示的に指定する）
43
cv::cuda::HostMem gray[2];
cv::cuda::GpuMat d_src[2], d_resize[2], d_gray[2];
cv::cuda::Stream stream[2];
// HtoD転送（stream0）
d_src[0].upload(src, stream[0]);
cv::cuda::resize(d_src[0], d_resize[0], cv::Size(), 2.0, 2.0, cv::INTER_LINEAR, stream[0]);
// HtoD転送（stream1）
d_src[1].upload(src, stream[1]);
cv::cuda::cvtColor(d_resize[0], d_gray[0], cv::COLOR_BGR2GRAY, 0, stream[0]);
// DtoH転送（stream0）
d_gray[0].download(gray[0], stream[0]);
cv::cuda::resize(d_src[1], d_resize[1], cv::Size(), 2.0, 2.0, cv::INTER_LINEAR, stream[1]);
cv::cuda::cvtColor(d_resize[1], d_gray[1], cv::COLOR_BGR2GRAY, 0, stream[1]);
// DtoH転送（stream1）
d_gray[1].download(gray[1], stream[1]);

View Slide

cv::cuda::Stream
● タイムライン（cv::cuda::Streamを明示的に指定する）
44
作成したstreamが使われているデータ転送とCUDAカーネル実行がオーバーラップしている

View Slide

cv::cuda::BufferPool

View Slide

● cv::cuda::BufferPool[4]
とは
○ あらかじめ確保したGPUのデバイスメモリ領域領域からGpuMatのメモリを割り当てることが
できるGpuMat専用のメモリプール機能
○ cudaMalloc、cudaMallocPitch、cudaFreeなどのCUDA API呼び出しを減らし、メモリ確保、
解放のオーバーヘッドを減らすことができる
○ メモリ確保サイズがあらかじめ決まっていて、規模の小さいプログラムで活用できる
■ 詳細は後述
46
[4] https://docs.opencv.org/4.6.0/d5/d08/classcv_1_1cuda_1_1BufferPool.html

View Slide

● サンプルコード（cv::cuda::BufferPool未使用）
47
cv::cuda::GpuMat d_src = cv::cuda::GpuMat d_src(cv::Size(1024, 1024), CV_8UC3);
cv::cuda::GpuMat d_dst = cv::cuda::GpuMat d_src(cv::Size(1024, 1024), CV_8UC1);
// GpuMatを使った処理
cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY, 0);

View Slide

● タイムライン（cv::cuda::BufferPool未使用）
48
GpuMatクラスのインスタンス生成の度に
cudaMallocPitchが呼ばれている

View Slide

● サンプルコード（cv::cuda::BufferPool使用）
49
//メモリプール機能の有効化
cv::cuda::setBufferPoolUsage(true);
// メモリプールのサイズ変更
cv::cuda::setBufferPoolConfig(cv::cuda::getDevice(), 1024 * 1024 * 64, 2);
cv::cuda::Stream stream;
// メモリプール生成
cv::cuda::BufferPool pool(stream);
// メモリプールから確保
cv::cuda::GpuMat d_src = pool.getBuffer(1024, 1024, CV_8UC3);
cv::cuda::GpuMat d_dst = pool.getBuffer(1024, 1024, CV_8UC1);
cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY, 0, stream);

View Slide

● タイムライン（cv::cuda::BufferPool使用）
50
メモリプールから確保すると
cudaMallocPitchが呼ばれない

View Slide

● タイムライン（cv::cuda::BufferPool使用）
51
メモリプールを新規確保
setBufferPoolConfig（メモリプール設定変更）を呼ぶと
デフォルトで確保していた領域を解放

View Slide

● cv::cuda::BufferPool使用時の注意点
○ BufferPoolクラスのインスタンス生成前にsetBufferPoolUsageをコールする必要がある
○ 解放順に気を付ける
○ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく
52

View Slide

○ BufferPoolクラスのインスタンス生成前にsetBufferPoolUsageをコールする必要がある
53
// BufferPoolクラスのインスタンス生成前に
// setBufferPoolUsageをコールしている
cv::cuda:: setBufferPoolUsage(true);
cv::cuda:: Stream stream;
cv::cuda:: BufferPool pool(stream);
cv::cuda:: GpuMat mat = pool.getBuffer(1024, 1024,
CV_8UC1);
cv::cuda:: Stream stream;
cv::cuda:: BufferPool pool(stream);
// BufferPoolクラスのインスタンス生成後に
// setBufferPoolUsageをコールしている
cv::cuda:: setBufferPoolUsage(true);
cv::cuda:: GpuMat mat = pool.getBuffer(1024, 1024,
CV_8UC1);

View Slide

○ 解放順に気を付ける
■ cv::cuda::BufferPoolで確保したメモリはLIFO順に解放する必要がある
■ 解放順を間違うとランタイムエラーが起きる
54
cv::cuda::GpuMat d_src1 = pool.getBuffer(512, 512, CV_8UC1);
d_src2.release();
d_src1.release();
d_src1.release();
d_src2.release();

View Slide

■ メモリプールからではなくDefaultAllocatorから確保される
55
size_t stack_size = 1024 * 1024 * 64;
cv::cuda::setBufferPoolConfig(cv::cuda::getDevice(), stack_size, 1);
cv::cuda::GpuMat d_img1 = pool.getBuffer(cv::Size(4096, 4096), CV_8UC3); // 48MB

View Slide

56
超過分はメモリプールからではなくDefaultAllocatorから確保される
（＝cudaMallocPitchが呼ばれる）

View Slide

GpuMat Tips

View Slide

GpuMat Tips
● OpenCVのCUDAカーネル内をデバッグしたい
○ WITH_CUDA=ONでOpenCVをビルドした場合、デフォルトだとデバッグ情報が付与されな
いため、OpenCVのCUDAカーネル内をデバッグできない。
○ opencv/cmake/OpenCVDetectCUDA.cmake[5]
を以下のように書き換えてデバック情報を付与
するようにしてOpenCVをビルドするのが簡単。
58
[5] https://github.com/opencv/opencv/blob/4.6.0/cmake/OpenCVDetectCUDA.cmake#L296
# NVCC flags to be set
set(NVCC_FLAGS_EXTRA "")
# NVCC flags to be set
set(NVCC_FLAGS_EXTRA "-G -g")

View Slide

GpuMat Tips
● cudaモジュールのビルド時間を短縮
○ デフォルトだと複数のCompute Capabilityをターゲットとしてビルドするため、ビルドに時間
が掛かってしまう
○ ターゲットのCompute Capabilityを絞ることでビルド時間を短縮できる
■ Compute Capabilityは https://developer.nvidia.com/cuda-gpus で調べられる
■ GeForce GTX 1650の場合のCMakeオプション指定例
● CUDA_ARCH_BIN="7.5"
● CUDA_ARCH_PTX=""
59

View Slide

GpuMat Tips
● cudaモジュールのビルド時間を短縮
○ CMake出力メッセージの違い
○ ビルド時間の違い
60
パターンビルド時間
Compute Capability指定なし 53min 8sec
Compute Capability指定あり 12min 57sec
-- NVIDIA CUDA: YES (ver 11.7, CUFFT CUBLAS)
-- NVIDIA GPU arch: 35 37 50 52 60 61 70 75 80 86
-- NVIDIA PTX archs:
-- NVIDIA CUDA: YES (ver 11.7, CUFFT CUBLAS)
-- NVIDIA GPU arch: 75
-- NVIDIA PTX archs:
Compute Capability指定なし Compute Capability指定あり
Compute Capability 7.5向けのみでビルドされる

View Slide

GpuMat Tips
● GpuMatの画像データのウィンドウ表示
○ ホストメモリに転送して表示
■ メリット：特別なセットアップ手順が不要
■ デメリット：画像表示のためにホストにデータ転送する必要がある
○ highgui(OpenGL)を使って表示
■ メリット：画像表示でホストにデータ転送しなくてよい
■ デメリット：OpenGLを有効化した設定でOpenCVセットアップが必要
61

View Slide

GpuMat Tips
1. downloadメソッドを使ってホストに転送
2. imshowを使ってウィンドウ表示
62
cv::cuda::GpuMat d_src(src), d_dst;
cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY);
// ホストメモリに転送する
cv::Mat dst;
d_dst.download(dst);
// ウィンドウ表示する
cv::namedWindow("dst", cv::WINDOW_AUTOSIZE);
cv::imshow("dst", dst);
cv::waitKey(0);

View Slide

GpuMat Tips
■ タイムライン（Nsight Systems）
63
CUDAカーネル
（cv::cuda::cvtColor）
cudaMemCpy2D（デバイス→ホスト）

View Slide

GpuMat Tips
○ highgui(OpenGL有効版)を使って表示
1. OpenCVビルド時にWITH_OPENGL=ONとする
2. namedWindowでcv::WINDOW_OPENGLのフラグを立てる
3. imshowにGpuMatクラスのインスタンスを渡す
64
cv::cuda::GpuMat d_src(src), d_dst;
cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY);
// namedWindowでcv::WINDOW_OPENGLのフラグを立てる
cv::namedWindow("d_dst", cv::WINDOW_AUTOSIZE | cv::WINDOW_OPENGL);
// imshowにGpuMatクラスのインスタンスを渡す
cv::imshow("d_dst", d_dst);
cv::waitKey(0);

View Slide

GpuMat Tips
○ highgui(OpenGL有効版)を使って表示
■ タイムライン（Nsight Systems）
65
cudaMemCpy2Dが呼ばれていない
CUDAカーネル
（cv::cuda::cvtColor）

View Slide

GpuMat注意ポイント

View Slide

● step、isContinuous
○ https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html に以下の記載がある
○ GpuMatの画像バッファのメモリは、ハードウェアに依存してアラインメントされるため、多
くのケースではisContinuous()==falseとなることに気を付ける
■ 例外として行数が1のGpuMatクラスのインスタンスはisContinuous()==trueとなる
67
In contrast with Mat, in most cases GpuMat::isContinuous() == false . This means
that rows are aligned to a size depending on the hardware. Single-row GpuMat is
always a continuous matrix.

View Slide

● step、isContinuousの扱いに注意
○ width=512、height=512
68
cv::Mat img(cv::Size(512, 512), CV_8UC1);
std::cout << "img.cols: " << img.cols << std::endl;
std::cout << "img.rows: " << img.rows << std::endl;
std::cout << "img.size: " << img.size() << std::endl;
std::cout << "img.step: " << img.step << std::endl;
std::cout << "img.isContinuous(): " << img.isContinuous() << std::endl <<
std::endl;
cv::cuda::GpuMat d_img(img);
std::cout << "d_img.cols: " << d_img.cols << std::endl;
std::cout << "d_img.rows: " << d_img.rows << std::endl;
std::cout << "d_img.size: " << d_img.size() << std::endl;
std::cout << "d_img.step: " << d_img.step << std::endl;
std::cout << "d_img.isContinuous(): " << d_img.isContinuous() <<
std::endl;
img.cols: 512
img.rows: 512
img.size: [512 x 512]
img.step: 512
img.isContinuous(): 1
d_img.cols: 512
d_img.rows: 512
d_img.size: [512 x 512]
d_img.step: 512
d_img.isContinuous(): 1
ソースコード（Mat）
ソースコード（GpuMat）
標準出力（Mat）標準出力（GpuMat）
isContinuous()=trueになっている

View Slide

○ width=100、height=100
69
std::cout << "img.cols: " << img.cols << std::endl;
std::cout << "img.rows: " << img.rows << std::endl;
std::cout << "img.size: " << img.size() << std::endl;
std::cout << "img.step: " << img.step << std::endl;
std::cout << "img.isContinuous(): " << img.isContinuous() << std::endl <<
std::endl;
cv::cuda::GpuMat d_img(img);
std::endl;
img.cols: 100
img.rows: 100
img.size: [100 x 100]
img.step: 100
img.isContinuous(): 1
d_img.cols: 100
d_img.rows: 100
d_img.size: [100 x 100]
d_img.step: 512
ソースコード（Mat）
ソースコード（GpuMat）
標準出力（Mat）標準出力（GpuMat）
isContinuous()=falseになっている

View Slide

○ cv::cuda::createContinuousメソッドを使うことで連続したメモリ確保にすることができる
70
cv::cuda::GpuMat d_img = cv::cuda::createContinuous(100, 100,
CV_8UC1);
d_img.upload(img);
std::endl;
d_img.cols: 100
d_img.rows: 100
d_img.size: [100 x 100]
d_img.step: 100
ソースコード（GpuMat）標準出力（GpuMat）
isContinuous()=trueになっている
https://docs.opencv.org/4.6.0/d9/d41/group__cudacore__struct.html#ga3a55474eb59c884697edf397fe0f871c

View Slide

● cudevモジュールが大量のconstant memoryを消費する
○ opencv2/cudev.hppをインクルードするとconstant memoryを確保してしまう
■ インクルードしない場合はconstant memoryを64KB確保できるが、インクルードすると
ビルドエラーになる
71
#include
#include
__constant__ float buffer[16384]; // 64KB((64*1024)/4)
int main(int argc, char *argv[])
{
std::exit(EXIT_SUCCESS);
}
#include
#include
__constant__ float buffer[16384]; // 64KB((64*1024)/4)
int main(int argc, char *argv[])
{
std::exit(EXIT_SUCCESS);
}
ptxas error : File uses too much global
constant data (0x1cfc0 bytes, 0x10000 max)

View Slide

● cv::cuda::PtrStepSzのオーバーヘッドに注意
○ 以下のCUDAカーネルの処理時間を比較
72
(
uchar* src, uchar* dst,
int width, int height, int step
)
{
if ((y >= 0) && (y < height))
{
if ((x >= 0) && (x < width))
{
uchar val = src[y*step + x];
dst[y*step + x] = 255 - val;
}
}
}
__global__ void inversionGpu(const cv::cuda::PtrStepSz src,
cv::cuda::PtrStepSz dst)
{
if ((y >= 0) && (y < src.rows))
{
if ((x >= 0) && (x < src.cols))
{
uchar val = src.ptr(y)[x];
dst.ptr(y)[x] = 255 - val;
}
}
}
PtrStepSz使用
PtrStepSz未使用
CUDAカーネル
CUDAカーネル

View Slide

● cv::cuda::PtrStepSzのオーバーヘッドに注意
○ 入力は画像サイズ4096x4096のグレースケール8bit画像
○ 前頁の白黒反転するCUDAカーネルを100回実行して、平均時間を計算
○ cv::cuda::PtrStepSzは便利だが、cv::cuda::GpuMatからcv::cuda::PtrStepSzへの暗黙キャスト
の分、処理時間がわずかに遅くなる
73
条件処理時間 [ms]
PtrStepSz使用 0.83
PtrStepSz未使用 0.78

View Slide

● static、グローバル変数はNG
○ https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html に以下の記載があり、
GpuMatクラスのインスタンスをstatic、グローバル変数として確保することは非推奨となっ
ている
74
You are not recommended to leave static or global GpuMat variables allocated, that
is, to rely on its destructor. The destruction order of such variables and CUDA
context is undefined. GPU memory release function returns error if the CUDA
context has been destroyed before.

View Slide

OpenCVコントリビューション活動事例
● 弊社で開発したlibSGM[6]
がOpenCVのcudastereoモジュールにマージされて
います
○ 弊社メンバおよびアルバイトの大塚さんによる成果で、OpenCV開発メンバーとのやり取り
は https://github.com/opencv/opencv_contrib/pull/2772 にあります
○ ニュース：https://news.fixstars.com/2151/
○ 技術ブログ：https://proc-cpuinfo.fixstars.com/2021/02/libsgmがopencvにマージされました/
75
[6] https://github.com/fixstars/libSGM

View Slide

OpenCV書籍執筆
● 「OpenCVではじめようディープラーニングによる画像認識」という書籍を
書きました。
○ 出版社：技術評論社
○ 著者：吉村康弘、五木田和也、杉浦司
○ 書籍URL：https://gihyo.jp/book/2022/978-4-297-12775-6
○ サンプルコード：https://github.com/ghmagazine/opencv_dl_book
76
OpenCVの基礎的な解説、dnnモジュールを用いたDNN推論処理を取り扱った書籍です。

View Slide

Thank you!
お問い合わせ窓口 : [email protected]

View Slide

Computer Vision Seminar 1/コンピュータビジョンセミナーvol.1 OpenCV活用

Computer Vision Seminar 1/コンピュータビジョンセミナーvol.1 OpenCV活用

fixstars

More Decks by fixstars

Other Decks in Programming

Featured

Transcript