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

12.9K Views

August 05, 22

#opencv #cuda #gpumat #コンピュータビジョン #コンピュータビジョンシリーズ #OpenCV #CUDA #GpuMat #Software development #Fixstars

スライド概要

コンピュータビジョンセミナーのシリーズ第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

株式会社フィックスターズ

@fixstars

スライド一覧

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

またはPlayer版

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

ダウンロード

ダウンロード(pdf - 2.3MB)

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

kaggle 画像処理機械学習深層学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 61.3K

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

arm neon

株式会社フィックスターズ 36.8K

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

kaggle 画像処理機械学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 28.5K

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

cuda gpu cuda高速化高速化シリーズ

株式会社フィックスターズ 27.7K

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

ros2 gpu ロボット自律走行 ros2シリーズ高速化シリーズ

株式会社フィックスターズ 26.9K

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

gpu cuda cuda高速化高速化シリーズ cuda高速化シリーズ

株式会社フィックスターズ 25.4K

各ページのテキスト

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

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

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

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

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

10.

11.

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

12.

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

13.

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.

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.

21.

22.

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

https://zenn.dev/onihusube/articles/d5c671870564b2

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

https://github.com/opencv/opencv/blob/4.6.0/modules/core/include/opencv2/core/cuda.inl.hpp#L207-L211

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.

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.

34.

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

https://developer.nvidia.com/npp

35.

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

https://docs.nvidia.com/cuda/npp/modules.html

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.

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

https://docs.opencv.org/4.6.0/db/d8c/group__cudaimgproc__color.html#ga48d0f208181d5ca370d8ff6b62cbe826

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.

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.

45.

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

https://docs.opencv.org/4.6.0/d5/d08/classcv_1_1cuda_1_1BufferPool.html

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.

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.

51.

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.

57.

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

https://github.com/opencv/opencv/blob/4.6.0/cmake/OpenCVDetectCUDA.cmake#L296

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

https://developer.nvidia.com/cuda-gpus

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.

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.

66.

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

https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html

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

https://docs.opencv.org/4.6.0/d9/d41/group__cudacore__struct.html#ga3a55474eb59c884697edf397fe0f871c

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

https://docs.opencv.org/4.6.0/d0/d60/classcv_1_1cuda_1_1GpuMat.html

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.

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

株式会社フィックスターズ

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

いまさら聞けない！CUDA高速化入門 ～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

各ページのテキスト

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

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）