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

8月5日に開催した「コンピュータビジョンセミナーvol.1 OpenCV活用～OpenCVでCUDAを活用するためのGpuMat解説～」の当日資料です。

Transcript

1. Copyright © Fixstars Group コンピュータビジョンセミナーvol.1 OpenCV活用 OpenCVでCUDAを活用するためのGpuMat解説
2. Copyright © Fixstars Group 本日のAgenda • はじめに • フィックスターズのご紹介 • OpenCVでCUDAを活用するためのGpuMat解説 ◦ GpuMatの概要 ◦ GpuMatと自作CUDAカーネルの連携 ◦ GpuMatとNPP連携 ◦ cv::cuda::Streamについて ◦ cv::cuda::BufferPoolについて ◦ GpuMat Tips ◦ GpuMat注意ポイント • Q&A / 告知 2
3. Copyright © Fixstars Group はじめに
4. Copyright © Fixstars Group • 弊社でサービス展開しているコンピュータビジョン領域において、 複数回に渡る技術セミナーの開催を計画しています • 今回は、コンピュータビジョン領域で最も利用されているフレームワーク OpenCVを 題材とし、CUDAを活用する上で重要なデータ構造 GpuMat について解説します • こんな方に向いています ◦ OpenCVと連携させるCUDAカーネルを自作したい ◦ GPUを活用して処理を高速化したい 本講演の位置づけ 4
5. Copyright © Fixstars Group 発表者紹介 冨田 明彦 ソリューションカンパニー 執行役員 2008年に入社。金融、医療業界において、ソ フトウェア高速化業務に携わる。その後、新規 事業企画、半導体業界の事業を担当し、現職。 5 吉村 康弘 Fixstars Autonomous Technologies リードエンジニア 2015年に入社。主に画像処理、コンピュータビ ジョンのアルゴリズム開発やCUDA高速化業務 を担当。
6. Copyright © Fixstars Group フィックスターズの ご紹介
7. Copyright © Fixstars Group フィックスターズの強み コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団 ハードウェアの知見 アルゴリズム実装力 各産業・研究分野の知見 7 目的の製品に最適なハードウェアを見抜き、 その性能をフル活用するソフトウェアを開 発します。 ハードウェアの特徴と製品要求仕様に合わ せて、アルゴリズムを改良して高速化を実 現します。 開発したい製品に使える技術を見抜き、実 際に動作する実装までトータルにサポート します。
8. Copyright © Fixstars Group 開発サービス提供分野 8 半導体 自動車 産業機器 生命科学 金融 •NAND型フラッシュメモリ向け ファームウェア開発 •次世代AIチップの開発環境基盤 •自動運転の高性能化、実用化 •次世代パーソナルモビリティの 研究開発 •Smart Factory実現への支援 •マシンビジョンシステムの高速化 •ゲノム解析の高速化 •医用画像処理の高速化 •AI画像診断システムの研究開発 •デリバティブシステムの高速化 •HFT(アルゴリズムトレード)の高速化
9. Copyright © Fixstars Group サービス領域 様々な領域でソフトウェア開発サービスを提供しています。大量データの高速処理は、 お客様の製品競争力の源泉となっています。 9 組込み高速化 画像処理・アルゴリズム 開発 分散並列システム開発 GPU向け高速化 FPGAを活用した システム開発 量子コンピューティング AI・深層学習 自動車向け ソフトウェア開発 フラッシュメモリ向けフ ァームウェア開発
10. Copyright © Fixstars Group 自動車向けソフトウェア開発 アルゴリズム開発から量産車ターゲット向けの高速化まで、 自動運転の実現に向けた統合的な技術開発を行っています。 ご支援内容
11. Copyright © Fixstars Group 画像処理アルゴリズム開発 高速な画像処理需要に対して、経験豊富なエンジニアが 責任を持って製品開発をご支援します。 お客様の課題 高度な画像処理や深層学習等のアルゴリズム を開発できる人材が社内に限られている 機能要件は満たせそうだが、ターゲット機器 上で性能要件までクリアできるか不安 製品化に結びつくような研究ができていない ご支援内容 深層学習ネットワーク精度の改善 様々な手法を駆使して深層学習ネットワークの精度を改善 論文調査・改善活動 論文調査から最先端の手法の探索 性能向上に向けた改善活動を継続 アルゴリズム調査・改変 課題に合ったアルゴリズム・実装手法を調査 製品実装に向けて適切な改変を実施
12. Copyright © Fixstars Group GPU向け高速化 高性能なGPUの本来の性能を十分に引き出し、 ソフトウェアの高速化を実現します。 お客様の課題 GPUで計算してみたが期待した性能が出ない GPU/CPUを組み合わせた全体として最適な 設計がしたい ご支援内容 GPU高速化に関するコンサルティング CPU・GPU混在環境でのシステム設計 アルゴリズムのGPU向け移植 GPUプログラム高速化 継続的な精度向上 原価を維持したまま機能を追加するため、も う少し処理を速くしたい 品質確保のため、精度を上げたく演算量は増 えるが性能は維持したい
13. Copyright © Fixstars Group GpuMatの概要
14. Copyright © Fixstars Group 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
15. Copyright © Fixstars Group GpuMatの概要 • OpenCVのcudaモジュールの処理フロー概要は以下の通り ◦ NPP（NVIDIA Performance Primitives）については後述 15 チャンネル数、depthチェック NPPを使う？ NPPの関数を呼ぶ OpenCV同梱の CUDAカーネルを呼ぶ OpenCVで サポート？ Y Y N N エラー
16. Copyright © Fixstars Group GpuMatの概要 • OpenCVのGpuMatを使うには ◦ NVIDIA CUDA ToolKitをインストールする ◦ OpenCVをソースコードからインストールする際、以下のCMakeオプションをONにする 16 CMakeオプション 意味 WITH_CUDA OpenCVでCUDAを使った実装を有効にする WITH_CUFFT OpenCVでcuFFTを使った実装を有効にする WITH_CUBLAS OpenCVでcuBLASを使った実装を有効にする
17. Copyright © Fixstars Group 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); と書くのでもよい
18. Copyright © Fixstars Group GpuMatの概要 • cudaモジュール概要 ◦ 代表的なモジュールは以下の通り 18 モジュール名 概要 cudaarithm 行列操作 cudabgsegm Background Segmentation cudacodec Video Encoding/Decoding cudafeatures2d Feature Detection and Description cudafilters フィルタ処理 cudaimgproc 色変換、ヒストグラム、コーナー検出など
19. Copyright © Fixstars Group GpuMatの概要 • cudaモジュール概要 ◦ 代表的なモジュールは以下の通り 19 モジュール名 概要 cudalegacy レガシーなアルゴリズム （オプティカルフロー、背景分離など） cudaobjdetect 物体検出 cudaoptflow オプティカルフロー（Sparse、Dense） cudastereo ステレオマッチング cudawarping リサイズ、アフィン変換など cudev Device layer
20. Copyright © Fixstars Group GpuMatと自作CUDA カーネルの連携
21. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatと自作CUDAカーネルは簡単に連携できる 21 coreモジュール GpuMat NVIDIA GPU 処理をオフロード 自作CUDAカーネル
22. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatを入力としたCUDAカーネルを作るメリットは以下の通り ◦ OpenCVのcudaモジュールと連携しやすい ▪ cudaモジュールにあるアルゴリズムと自作CUDAカーネルを組み合わせて実行する等 22 coreモジュール 自作CUDAカーネル cudaモジュール GpuMat
23. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatを入力としたCUDAカーネルを作るメリットは以下の通り ◦ 画像ファイルの入出力、画像表示をOpenCVで行うことができる 23 coreモジュール 自作CUDAカーネル imgcodecsモジュール 画像ファイル入出力 highguiモジュール 画像表示 GpuMat
24. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • cv::cuda::PtrStepSz ◦ CUDAカーネル（デバイスコード）で画像データにアクセスが容易となるデータ型 ▪ 日本語情報だと https://zenn.dev/onihusube/articles/d5c671870564b2 の解説がとてもわ かりやすい ◦ cv::cuda::PtrStepSzは名前から類推できるようにstep、サイズ情報（rows、cols）にアクセス できる ◦ step情報にアクセスできるcv::cuda::PtrStepも用意されている 24
25. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatにある画像バッファのアドレス参照 ◦ 代表的な方法は以下の通り ▪ GpuMatクラスのptrメソッド ▪ cv::cuda::PtrStepSz 25
26. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatにある画像バッファのアドレス参照 ◦ GpuMatクラスのptrメソッド ▪ cv::cuda::GpuMat.ptr<typename T>(y) • T：データ型 • y：参照する列（デフォルトだとy=0） 26 cv::cuda::GpuMat src(cv::Size(320, 240), CV_8UC1); // 画像バッファの先頭アドレスを取得 // 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 uchar* GpuMat::ptr(int y) { CV_DbgAssert( (unsigned)y < (unsigned)rows ); return data + step * y; } ptrメソッドの内部実装 サンプルコード
27. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatにある画像バッファのアドレス参照 ◦ cv::cuda::PtrStepSz 27 __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(y, x); dst(y, x) = 255 - value; } } } CUDAカーネル uchar value = src.ptr(y)[x]; dst.ptr(y)[x] = 255 - value; と書くのでもよい
28. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • GpuMatを入力としたCUDAカーネルの作り方 ◦ 方法１：画像バッファおよびサイズ情報を引数にしたCUDAカーネルを実装する ◦ 方法２：cv::cuda::PtrStepSzを引数にしたCUDAカーネルを実装する 28
29. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • サンプルコード（CUDAカーネル） 29 __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カーネル 画像バッファのポインタ、 width、height、ステップを引数で渡す 自前でアドレスを計算して読み書きする
30. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • サンプルコード（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<<<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()); } src.data、src.datastartと書くのでもよい
31. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • サンプルコード（CUDAカーネル） 31 __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]); } } } CUDAカーネル cv::cuda::PtrStepSzを引数で渡す ptrメソッドで座標(x, y)のピクセルのアド レスを参照し、読み書きする
32. Copyright © Fixstars Group GpuMatと自作CUDAカーネルの連携 • サンプルコード（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<<<grid, block>>>(src, dst); // エラーチェック CV_CUDEV_SAFE_CALL(cudaGetLastError()); CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize()); } GpuMatクラスのインスタンスを渡す
33. Copyright © Fixstars Group GpuMatとNPPの連携
34. Copyright © Fixstars Group GpuMatとNPPの連携 • NPP[3] とは ◦ NVIDIA Performance Primitivesのこと ◦ 画像処理、信号処理等の各種アルゴリズムのCUDA実装ライブラリ • NPPとOpenCVの実装は連携することができる ◦ 以降でNPPの基本的な情報と連携方法を紹介します 34 [3] https://developer.nvidia.com/npp
35. Copyright © Fixstars Group GpuMatとNPPの連携 • NPPの機能概要（画像処理） ◦ 色変換 ◦ フィルタ処理 ◦ Geometry Transforms（回転、反転、リサイズなど） ◦ モルフォロジー変換（Erode、Dilate） ◦ Statistical Operations（総和、最大値・最小値計算、ヒストグラム計算、インテグラルイメー ジ作成など） ◦ etc... 35 詳細は https://docs.nvidia.com/cuda/npp/modules.html 参照のこと
36. Copyright © Fixstars Group GpuMatとNPPの連携 • NPPで用いる基本的なデータ型 ◦ Npp<ビット数><データ型>という命名規則となっている 36 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
37. Copyright © Fixstars Group GpuMatとNPPの連携 • サンプルコード（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)
38. Copyright © Fixstars Group cv::cuda::Stream
39. Copyright © Fixstars Group cv::cuda::Stream • Streamとは ◦ GPU処理のスケジュール管理の単位 39
40. Copyright © Fixstars Group 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
41. Copyright © Fixstars Group 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]);
42. Copyright © Fixstars Group cv::cuda::Stream • タイムライン（cv::cuda::Streamを明示的に指定しない） 42 Default streamのみが使われている データ転送でブロッキングされている
43. Copyright © Fixstars Group 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]);
44. Copyright © Fixstars Group cv::cuda::Stream • タイムライン（cv::cuda::Streamを明示的に指定する） 44 作成したstreamが使われている データ転送とCUDAカーネル実行がオーバーラップしている
45. Copyright © Fixstars Group cv::cuda::BufferPool
46. Copyright © Fixstars Group cv::cuda::BufferPool • 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
47. Copyright © Fixstars Group cv::cuda::BufferPool • サンプルコード（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);
48. Copyright © Fixstars Group cv::cuda::BufferPool • タイムライン（cv::cuda::BufferPool未使用） 48 GpuMatクラスのインスタンス生成の度に cudaMallocPitchが呼ばれている
49. Copyright © Fixstars Group cv::cuda::BufferPool • サンプルコード（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); // GpuMatを使った処理 cv::cuda::cvtColor(d_src, d_dst, cv::COLOR_BGR2GRAY, 0, stream);
50. Copyright © Fixstars Group cv::cuda::BufferPool • タイムライン（cv::cuda::BufferPool使用） 50 メモリプールから確保すると cudaMallocPitchが呼ばれない
51. Copyright © Fixstars Group cv::cuda::BufferPool • タイムライン（cv::cuda::BufferPool使用） 51 メモリプールを新規確保 setBufferPoolConfig（メモリプール設定変更）を呼ぶと デフォルトで確保していた領域を解放
52. Copyright © Fixstars Group cv::cuda::BufferPool • cv::cuda::BufferPool使用時の注意点 ◦ BufferPoolクラスのインスタンス生成前にsetBufferPoolUsageをコールする必要がある ◦ 解放順に気を付ける ◦ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく 52
53. Copyright © Fixstars Group cv::cuda::BufferPool • cv::cuda::BufferPool使用時の注意点 ◦ 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);
54. Copyright © Fixstars Group cv::cuda::BufferPool • cv::cuda::BufferPool使用時の注意点 ◦ 解放順に気を付ける ▪ cv::cuda::BufferPoolで確保したメモリはLIFO順に解放する必要がある ▪ 解放順を間違うとランタイムエラーが起きる 54 cv::cuda::setBufferPoolUsage(true); 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); d_src2.release(); d_src1.release(); cv::cuda::setBufferPoolUsage(true); 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); d_src1.release(); d_src2.release();
55. Copyright © Fixstars Group cv::cuda::BufferPool • cv::cuda::BufferPool使用時の注意点 ◦ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく ▪ メモリプールからではなくDefaultAllocatorから確保される 55 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
56. Copyright © Fixstars Group cv::cuda::BufferPool • cv::cuda::BufferPool使用時の注意点 ◦ メモリプールの容量以上のメモリを確保した場合の挙動を理解しておく 56 超過分はメモリプールからではなくDefaultAllocatorから確保される （＝cudaMallocPitchが呼ばれる）
57. Copyright © Fixstars Group GpuMat Tips
58. Copyright © Fixstars Group 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")
59. Copyright © Fixstars Group 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
60. Copyright © Fixstars Group 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向けのみでビルドされる
61. Copyright © Fixstars Group GpuMat Tips • GpuMatの画像データのウィンドウ表示 ◦ ホストメモリに転送して表示 ▪ メリット：特別なセットアップ手順が不要 ▪ デメリット：画像表示のためにホストにデータ転送する必要がある ◦ highgui(OpenGL)を使って表示 ▪ メリット：画像表示でホストにデータ転送しなくてよい ▪ デメリット：OpenGLを有効化した設定でOpenCVセットアップが必要 61
62. Copyright © Fixstars Group GpuMat Tips • GpuMatの画像データのウィンドウ表示 ◦ ホストメモリに転送して表示 1. downloadメソッドを使ってホストに転送 2. imshowを使ってウィンドウ表示 62 // 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);
63. Copyright © Fixstars Group GpuMat Tips • GpuMatの画像データのウィンドウ表示 ◦ ホストメモリに転送して表示 ▪ タイムライン（Nsight Systems） 63 CUDAカーネル （cv::cuda::cvtColor） cudaMemCpy2D（デバイス→ホスト）
64. Copyright © Fixstars Group GpuMat Tips • GpuMatの画像データのウィンドウ表示 ◦ highgui(OpenGL有効版)を使って表示 1. OpenCVビルド時にWITH_OPENGL=ONとする 2. namedWindowでcv::WINDOW_OPENGLのフラグを立てる 3. imshowにGpuMatクラスのインスタンスを渡す 64 // 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);
65. Copyright © Fixstars Group GpuMat Tips • GpuMatの画像データのウィンドウ表示 ◦ highgui(OpenGL有効版)を使って表示 ▪ タイムライン（Nsight Systems） 65 cudaMemCpy2Dが呼ばれていない CUDAカーネル （cv::cuda::cvtColor）
66. Copyright © Fixstars Group GpuMat注意ポイント
67. Copyright © Fixstars Group GpuMat注意ポイント • 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.
68. Copyright © Fixstars Group GpuMat注意ポイント • 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::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; 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になっている
69. Copyright © Fixstars Group GpuMat注意ポイント • step、isContinuousの扱いに注意 ◦ width=100、height=100 69 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; 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; 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 d_img.isContinuous(): 0 ソースコード（Mat） ソースコード（GpuMat） 標準出力（Mat） 標準出力（GpuMat） isContinuous()=falseになっている
70. Copyright © Fixstars Group GpuMat注意ポイント • step、isContinuousの扱いに注意 ◦ cv::cuda::createContinuousメソッドを使うことで連続したメモリ確保にすることができる 70 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） 標準出力（GpuMat） isContinuous()=trueになっている https://docs.opencv.org/4.6.0/d9/d41/group__cudacore__struct.html#ga3a55474eb59c884697edf397fe0f871c
71. Copyright © Fixstars Group GpuMat注意ポイント • cudevモジュールが大量のconstant memoryを消費する ◦ opencv2/cudev.hppをインクルードするとconstant memoryを確保してしまう ▪ インクルードしない場合はconstant memoryを64KB確保できるが、インクルードすると ビルドエラーになる 71 #include <opencv2/cudev/common.hpp> #include <iostream> __constant__ float buffer[16384]; // 64KB((64*1024)/4) int main(int argc, char *argv[]) { std::exit(EXIT_SUCCESS); } #include <opencv2/cudev.hpp> #include <iostream> __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)
72. Copyright © Fixstars Group GpuMat注意ポイント • cv::cuda::PtrStepSzのオーバーヘッドに注意 ◦ 以下のCUDAカーネルの処理時間を比較 72 __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; } } } __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; } } } PtrStepSz使用 PtrStepSz未使用 CUDAカーネル CUDAカーネル
73. Copyright © Fixstars Group GpuMat注意ポイント • cv::cuda::PtrStepSzのオーバーヘッドに注意 ◦ 入力は画像サイズ4096x4096のグレースケール8bit画像 ◦ 前頁の白黒反転するCUDAカーネルを100回実行して、平均時間を計算 ◦ cv::cuda::PtrStepSzは便利だが、cv::cuda::GpuMatからcv::cuda::PtrStepSzへの暗黙キャスト の分、処理時間がわずかに遅くなる 73 条件 処理時間 [ms] PtrStepSz使用 0.83 PtrStepSz未使用 0.78
74. Copyright © Fixstars Group GpuMat注意ポイント • 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.
75. Copyright © Fixstars Group 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
76. Copyright © Fixstars Group OpenCV書籍執筆 • 「OpenCVではじめよう ディープラーニングによる画像認識」という書籍を 書きました。 ◦ 出版社：技術評論社 ◦ 著者：吉村康弘、五木田和也、杉浦司 ◦ 書籍URL：https://gihyo.jp/book/2022/978-4-297-12775-6 ◦ サンプルコード：https://github.com/ghmagazine/opencv_dl_book 76 OpenCVの基礎的な解説、dnnモジュールを用いたDNN推論処理を取り扱った書籍です。
77. Copyright © Fixstars Group Thank you! お問い合わせ窓口 : hr-seminar@fixstars.com