コンピュータビジョンセミナーvol.1 OpenCV活用 ~OpenCVでCUDAを活用するためのGpuMat解説~(2022/08/05)

5.8K Views

August 05, 22

スライド概要

コンピュータビジョンセミナーのシリーズ第1回として、GpuMatの活用方法を解説します。GpuMatは、OpenCVでCUDAによる高速化の恩恵を受ける手段の一つです。
コンピュータビジョンに関連する開発業務を行っているエンジニアに最適なセミナーです。

<講演内容>
・GpuMatの概要
・GpuMatと自作CUDAカーネルの連携
・GpuMatとNPP連携
・cv::cuda::Streamについて
・cv::cuda::BufferPoolについて
・GpuMat Tips
・GpuMat注意ポイント
・Q&A

<過去資料>
・vol.1 OpenCV活用(OpenCVでCUDAを活用するためのGpuMat解説): https://www.docswell.com/s/fixstars/ZRXQ72-20220805
・vol.2 視差計算ライブラリ libSGM のアルゴリズム解説と CUDA高速化: https://www.docswell.com/s/fixstars/5ENE7J-20221220
・vol.3 視差計算を利用した障害物検出: https://www.docswell.com/s/fixstars/ZQ8447-20231027

・フィックスターズの画像処理アルゴリズム開発支援: https://www.fixstars.com/ja/services/image-processing

profile-image

フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。 並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。       ・開催セミナー一覧:https://www.fixstars.com/ja/seminar   ・技術ブログ :https://proc-cpuinfo.fixstars.com/

シェア

埋め込む »CMSなどでJSが使えない場合

関連スライド

各ページのテキスト
1.

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

2.

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

3.

はじめに Copyright © Fixstars Group

4.

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

5.

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

6.

フィックスターズの ご紹介 Copyright © Fixstars Group

7.

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

8.

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

9.

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

10.

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

11.

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

12.

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

13.

GpuMatの概要 Copyright © Fixstars Group

14.

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

15.

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

16.

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

17.

GpuMatの概要 ● サンプルコード cv::Mat src(cv::Size(3840, 2160), CV_8UC3, cv::Scalar(0, 0, 0)); cv::cuda::GpuMat d_src; // GPUに転送 d_src.upload(src); cv::cuda::GpuMat d_src(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); Copyright © Fixstars Group 17

18.

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

19.

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

20.

GpuMatと自作CUDA カーネルの連携 Copyright © Fixstars Group

21.

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

22.

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

23.

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

24.

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

25.

GpuMatと自作CUDAカーネルの連携 ● GpuMatにある画像バッファのアドレス参照 ○ 代表的な方法は以下の通り ■ GpuMatクラスのptrメソッド ■ cv::cuda::PtrStepSz Copyright © Fixstars Group 25

26.
[beta]
GpuMatと自作CUDAカーネルの連携
● GpuMatにある画像バッファのアドレス参照
○ GpuMatクラスのptrメソッド
■ cv::cuda::GpuMat.ptr<typename T>(y)
● T:データ型
● y:参照する列(デフォルトだとy=0)
サンプルコード

ptrメソッドの内部実装

cv::cuda::GpuMat src(cv::Size(320, 240), CV_8UC1);

uchar* GpuMat::ptr(int y)
{
CV_DbgAssert( (unsigned)y < (unsigned)rows );
return data + step * y;
}

// 画像バッファの先頭アドレスを取得
// uchar* psrc = src.data でも等価
uchar* psrc = src.ptr<uchar>(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

Copyright © Fixstars Group

26

27.
[beta]
GpuMatと自作CUDAカーネルの連携
● GpuMatにある画像バッファのアドレス参照
○ cv::cuda::PtrStepSz
__global__ void inversionGpu(const cv::cuda::PtrStepSz<uchar> src,
cv::cuda::PtrStepSz<uchar> 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.ptr(y)[x];
{
dst.ptr(y)[x] = 255 - value;
uchar value = src(y, x);
と書くのでもよい
dst(y, x) = 255 - value;
}
}
}
CUDAカーネル
Copyright © Fixstars Group

27

28.

GpuMatと自作CUDAカーネルの連携 ● GpuMatを入力としたCUDAカーネルの作り方 ○ 方法1:画像バッファおよびサイズ情報を引数にしたCUDAカーネルを実装する ○ 方法2:cv::cuda::PtrStepSzを引数にしたCUDAカーネルを実装する Copyright © Fixstars Group 28

29.
[beta]
GpuMatと自作CUDAカーネルの連携
● サンプルコード(CUDAカーネル)
__global__ void inversionGpu
(
uchar* src, uchar* dst,
int width, int height, int step
)
{
const int x = blockDim.x * blockIdx.x + threadIdx.x;
const int y = blockDim.y * blockIdx.y + threadIdx.y;
if ((y >= 0) && (y < height))
{
if ((x >= 0) && (x < width))
{
uchar val = src[y*step + x];
dst[y*step + x] = 255 - val;
}
}
}

画像バッファのポインタ、
width、height、ステップを引数で渡す

自前でアドレスを計算して読み書きする

Copyright © Fixstars Group

CUDAカーネル

29

30.
[beta]
GpuMatと自作CUDAカーネルの連携
● サンプルコード(CUDAカーネル呼び出し)
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<<<grid, block>>>(src.ptr<uchar>(0), dst.ptr<uchar>(0), src.cols, src.rows, src.step);

// エラーチェック
CV_CUDEV_SAFE_CALL(cudaGetLastError());
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());
}

Copyright © Fixstars Group

src.data、src.datastartと書くのでもよい

30

31.
[beta]
GpuMatと自作CUDAカーネルの連携
● サンプルコード(CUDAカーネル)
__global__ void inversionGpu
(
const cv::cuda::PtrStepSz<uchar> src,
cv::cuda::PtrStepSz<uchar> 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))
{
dst.ptr(y)[x] = (255 - src.ptr(y)[x]);
}
}
}

cv::cuda::PtrStepSzを引数で渡す

ptrメソッドで座標(x, y)のピクセルのアド
レスを参照し、読み書きする

Copyright © Fixstars Group

CUDAカーネル
31

32.
[beta]
GpuMatと自作CUDAカーネルの連携
● サンプルコード(CUDAカーネル呼び出し)
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<<<grid, block>>>(src, dst);

// エラーチェック
CV_CUDEV_SAFE_CALL(cudaGetLastError());
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());

GpuMatクラスのインスタンスを渡す

}

Copyright © Fixstars Group

32

33.

GpuMatとNPPの連携 Copyright © Fixstars Group

34.

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

35.

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

36.

GpuMatとNPPの連携 ● NPPで用いる基本的なデータ型 ○ Npp<ビット数><データ型>という命名規則となっている NPPで定義されるデータ型 実際に用いられるデータ型 Npp8u 8-bit unsigned char Npp8s 8-bit signed char Npp16u 16-bit unsigned integer Npp16s 16-bit signed integer Npp32u 32-bit unsigned integer Npp32s 32-bit signed integer Npp64u 64-bit unsigned integer Npp64s 64-bit signed integer Npp32f 32-bit (IEEE) floating-point number Npp64f 64-bit floating-point number Copyright © Fixstars Group 36

37.

GpuMatとNPPの連携 ● サンプルコード(GpuMatとNPPの連携) ○ GpuMatクラスのインスタンスを入出力として、NPPのメディアンフィルタを実行する // GpuMatクラスのインスタンス生成 cv::cuda::GpuMat d_src(src); cv::cuda::GpuMat d_dst(dst); // 中略 8u:8-bit unsigned char C1:1channel R:Region-of-Interest(ROI) // 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); Copyright © Fixstars Group 37

38.

cv::cuda::Stream Copyright © Fixstars Group

39.

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

40.

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

41.

cv::cuda::Stream ● サンプルコード(cv::cuda::Streamを明示的に指定しない) 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]); Copyright © Fixstars Group 41

42.

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

43.

cv::cuda::Stream ● サンプルコード(cv::cuda::Streamを明示的に指定する) 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]); Copyright © Fixstars Group 43

44.

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

45.

cv::cuda::BufferPool Copyright © Fixstars Group

46.

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

47.

cv::cuda::BufferPool ● サンプルコード(cv::cuda::BufferPool未使用) 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); Copyright © Fixstars Group 47

48.

cv::cuda::BufferPool ● タイムライン(cv::cuda::BufferPool未使用) GpuMatクラスのインスタンス生成の度に cudaMallocPitchが呼ばれている Copyright © Fixstars Group 48

49.

cv::cuda::BufferPool ● サンプルコード(cv::cuda::BufferPool使用) //メモリプール機能の有効化 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); // GpuMatを使った処理 cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY, 0, stream); Copyright © Fixstars Group 49

50.

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

51.

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

52.

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

53.

cv::cuda::BufferPool ● cv::cuda::BufferPool使用時の注意点 ○ BufferPoolクラスのインスタンス生成前にsetBufferPoolUsageをコールする必要がある // 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); Copyright © Fixstars Group 53

54.

cv::cuda::BufferPool ● cv::cuda::BufferPool使用時の注意点 ○ 解放順に気を付ける ■ cv::cuda::BufferPoolで確保したメモリはLIFO順に解放する必要がある ■ 解放順を間違うとランタイムエラーが起きる cv::cuda::setBufferPoolUsage(true); cv::cuda::setBufferPoolUsage(true); cv::cuda::Stream stream; cv::cuda::BufferPool pool(stream); cv::cuda::Stream stream; cv::cuda::BufferPool pool(stream); cv::cuda::GpuMat d_src1 = pool.getBuffer(512, 512, CV_8UC1); cv::cuda::GpuMat d_src2 = pool.getBuffer(512, 512, CV_8UC1); cv::cuda::GpuMat d_src1 = pool.getBuffer(512, 512, CV_8UC1); cv::cuda::GpuMat d_src2 = pool.getBuffer(512, 512, CV_8UC1); d_src2.release(); d_src1.release(); d_src1.release(); d_src2.release(); Copyright © Fixstars Group 54

55.

cv::cuda::BufferPool ● cv::cuda::BufferPool使用時の注意点 ○ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく ■ メモリプールからではなくDefaultAllocatorから確保される size_t stack_size = 1024 * 1024 * 64; cv::cuda::setBufferPoolConfig(cv::cuda::getDevice(), stack_size, 1); cv::cuda::Stream stream; cv::cuda::BufferPool pool(stream); cv::cuda::GpuMat d_img1 = pool.getBuffer(cv::Size(4096, 4096), CV_8UC3); // 48MB cv::cuda::GpuMat d_img2 = pool.getBuffer(cv::Size(4096, 4096), CV_8UC1); // 16MB cv::cuda::GpuMat d_img3 = pool.getBuffer(cv::Size(4096, 4096), CV_8UC1); // 16MB Copyright © Fixstars Group 55

56.

cv::cuda::BufferPool ● cv::cuda::BufferPool使用時の注意点 ○ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく 超過分はメモリプールからではなくDefaultAllocatorから確保される (=cudaMallocPitchが呼ばれる) Copyright © Fixstars Group 56

57.

GpuMat Tips Copyright © Fixstars Group

58.

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

59.

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="" Copyright © Fixstars Group 59

60.

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

61.

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

62.

GpuMat Tips ● GpuMatの画像データのウィンドウ表示 ○ ホストメモリに転送して表示 1. downloadメソッドを使ってホストに転送 2. imshowを使ってウィンドウ表示 // GpuMatを使った処理 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); Copyright © Fixstars Group 62

63.

GpuMat Tips ● GpuMatの画像データのウィンドウ表示 ○ ホストメモリに転送して表示 ■ タイムライン(Nsight Systems) cudaMemCpy2D(デバイス→ホスト) CUDAカーネル (cv::cuda::cvtColor) Copyright © Fixstars Group 63

64.

GpuMat Tips ● GpuMatの画像データのウィンドウ表示 ○ highgui(OpenGL有効版)を使って表示 1. OpenCVビルド時にWITH_OPENGL=ONとする 2. namedWindowでcv::WINDOW_OPENGLのフラグを立てる 3. imshowにGpuMatクラスのインスタンスを渡す // GpuMatを使った処理 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); Copyright © Fixstars Group 64

65.

GpuMat Tips ● GpuMatの画像データのウィンドウ表示 ○ highgui(OpenGL有効版)を使って表示 ■ タイムライン(Nsight Systems) CUDAカーネル (cv::cuda::cvtColor) cudaMemCpy2Dが呼ばれていない Copyright © Fixstars Group 65

66.

GpuMat注意ポイント Copyright © Fixstars Group

67.

GpuMat注意ポイント ● step、isContinuous ○ https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html に以下の記載がある 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. ○ GpuMatの画像バッファのメモリは、ハードウェアに依存してアラインメントされるため、多 くのケースではisContinuous()==falseとなることに気を付ける ■ 例外として行数が1のGpuMatクラスのインスタンスはisContinuous()==trueとなる Copyright © Fixstars Group 67

68.

GpuMat注意ポイント ● step、isContinuousの扱いに注意 ○ width=512、height=512 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::Mat img(cv::Size(512, 512), CV_8UC1); 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; ソースコード(Mat) img.cols: 512 img.rows: 512 img.size: [512 x 512] img.step: 512 img.isContinuous(): 1 ソースコード(GpuMat) d_img.cols: 512 d_img.rows: 512 d_img.size: [512 x 512] d_img.step: 512 d_img.isContinuous(): 1 標準出力(Mat) isContinuous()=trueになっている 標準出力(GpuMat) Copyright © Fixstars Group 68

69.

GpuMat注意ポイント ● step、isContinuousの扱いに注意 ○ width=100、height=100 cv::Mat img(cv::Size(100, 100), 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; ソースコード(Mat) img.cols: 100 img.rows: 100 img.size: [100 x 100] img.step: 100 img.isContinuous(): 1 cv::Mat img(cv::Size(100, 100), CV_8UC1); 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; ソースコード(GpuMat) d_img.cols: 100 d_img.rows: 100 d_img.size: [100 x 100] d_img.step: 512 d_img.isContinuous(): 0 標準出力(Mat) isContinuous()=falseになっている 標準出力(GpuMat) Copyright © Fixstars Group 69

70.

GpuMat注意ポイント ● step、isContinuousの扱いに注意 ○ cv::cuda::createContinuousメソッドを使うことで連続したメモリ確保にすることができる cv::cuda::GpuMat d_img = cv::cuda::createContinuous(100, 100, CV_8UC1); d_img.upload(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; d_img.cols: 100 d_img.rows: 100 d_img.size: [100 x 100] d_img.step: 100 d_img.isContinuous(): 1 ソースコード(GpuMat) isContinuous()=trueになっている 標準出力(GpuMat) https://docs.opencv.org/4.6.0/d9/d41/group__cudacore__struct.html#ga3a55474eb59c884697edf397fe0f871c Copyright © Fixstars Group 70

71.
[beta]
GpuMat注意ポイント
● cudevモジュールが大量のconstant memoryを消費する
○ opencv2/cudev.hppをインクルードするとconstant memoryを確保してしまう
■ インクルードしない場合はconstant memoryを64KB確保できるが、インクルードすると
ビルドエラーになる
#include <opencv2/cudev/common.hpp>
#include <iostream>

#include <opencv2/cudev.hpp>
#include <iostream>

__constant__ float buffer[16384]; // 64KB((64*1024)/4)

__constant__ float buffer[16384]; // 64KB((64*1024)/4)

int main(int argc, char *argv[])
{
std::exit(EXIT_SUCCESS);
}

int main(int argc, char *argv[])
{
std::exit(EXIT_SUCCESS);
}

ptxas error : File uses too much global
constant data (0x1cfc0 bytes, 0x10000 max)
Copyright © Fixstars Group

71

72.
[beta]
GpuMat注意ポイント
● cv::cuda::PtrStepSzのオーバーヘッドに注意
○ 以下のCUDAカーネルの処理時間を比較
__global__ void inversionGpu(const cv::cuda::PtrStepSz<uchar> src,
cv::cuda::PtrStepSz<uchar> 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 val = src.ptr(y)[x];
dst.ptr(y)[x] = 255 - val;
}
}
}
CUDAカーネル

__global__ void inversionGpu
(
uchar* src, uchar* dst,
int width, int height, int step
)
{
const int x = blockDim.x * blockIdx.x + threadIdx.x;
const int y = blockDim.y * blockIdx.y + threadIdx.y;
if ((y >= 0) && (y < height))
{
if ((x >= 0) && (x < width))
{
uchar val = src[y*step + x];
dst[y*step + x] = 255 - val;
}
}
}
CUDAカーネル

PtrStepSz使用

PtrStepSz未使用
Copyright © Fixstars Group

72

73.

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

74.

GpuMat注意ポイント ● static、グローバル変数はNG ○ https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html に以下の記載があり、 GpuMatクラスのインスタンスをstatic、グローバル変数として確保することは非推奨となっ ている 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. Copyright © Fixstars Group 74

75.

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にマージされました/ Copyright © Fixstars Group [6] https://github.com/fixstars/libSGM 75

76.

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

77.

Thank you! お問い合わせ窓口 : hr-seminar@fixstars.com Copyright © Fixstars Group