💭

OpenCVをNPPにした結果→10倍高速に！

chizuchizu

2023/02/03に公開

この記事は「自動運転システムをエッジデバイスに組み込むための技術」を3回に分けて紹介するTURINGのテックブログ連載の第2回の記事「OpenCVをNPPにした結果→10倍高速に！」です。

第1回の「C++でOpenCV完全入門！」、第3回の「詳解V4L2 (video for linux 2)」もぜひご覧ください！

はじめに

TURINGで働いている木更津高専の越智です。TURINGでは「We Overtake Tesla」を目標に掲げて、完全自動運転EVの開発・製造を行っています。

TURINGでは、社内で使っている自動運転ソフトウェアにおいて、画像処理部分のライブラリをOpenCVからNVIDIA Performance Primitives(NPP)に変更するプロジェクトに取り組んでいました。これによって、CPUで動かしていた画像処理をGPUバックエンドで動かすことができるようになりました。結果として、画像処理の10倍以上の大幅な高速化を実現しながら、同時にCPUリソースの多くを解放することに成功しています。本記事では、この改善をどうやって実現したかについて解説します。

自動運転技術の裏側

本記事にて取り扱う話は、TURINGが販売する車に搭載される自動運転ソフトウェアの裏側に関する話です！

TURINGでは「THE FIRST TURING CAR」という車を販売しています。最近プレスリリースが出ました！

「THE 1st TURING CAR」はレクサスRX450hをベース車両として、チューリングが自社開発したAI自動運転システムと、オリジナルエンブレムが搭載されており、チューリングとして初めてエンドユーザー向けに販売する製品です。

TURINGが販売する「THE FIRST TURING CAR」に載せるコンピュータには、車載の諸要件（耐振動性、コネクタの抜けにくさなど）を満たす必要があります。このような要件を満たすデバイスを選定しましたが、NVIDIAのGPUを搭載しているものの、一般的なPCと比較して性能が貧弱でした。

パソコンに詳しい人のために説明すると、CPUはラズパイや最近のスマホより弱いし、メモリは4GBで動いています。さらに、GPUが好きな人のために説明すると、GPUはGTX 1080世代と同等のコアを1/10しか積んでいません。

今までリッチなPCで開発してきたソフトウェアをエッジデバイスに落とし込んだために、CPU使用率は全コア100%、4GBあるメモリ使用率は95%以上と、リソース的に非常に厳しい状態でした。更には処理が重すぎてFPSが低いといった問題もあり、当初はこのデバイスで本当に動くのかと心配になっていました。

この課題を解決する方法として次のものが挙げられます。

処理を高速化する
- アルゴリズムの改善
- GPUの活用（大規模並列処理）
CPU以外で処理する
- GPUやその他ハードウェアを活用

このうち、「GPUの活用」のため、NPPを選択しました。

NPPとは

NVIDIA Performance Primitivesの略です。NVIDIAによる、GPUで高速化された画像や信号処理関数を提供するライブラリです。

https://developer.nvidia.com/npp

GPUが乗っていて、CUDA Toolkitが入っているデバイスなら動きます。

なぜNPPが必要なのか

NPPを使うと、GPUの上で画像処理が可能になります。GPUの特徴として次のようなものがあります。

GPUはCPUに比べ大規模な並列処理に適しており、100倍以上並列して計算を行うことができる
GPUはCPUと比べ、複雑な処理ができない（画像処理など単純なものに限られる）

自動運転ソフトウェアで扱う画像は100万ピクセル以上あるため、CPUよりも大規模並列処理ができるGPUのほうが適しているといえます。したがって、GPU上で画像処理を行うことでパフォーマンスの向上を期待できます。

誰にオススメできるか

NVIDIAのGPUを持っている人
リアルタイム画像処理をしたい人
大きなデータの画像処理をしたい人
OpenCVに飽き飽きした人

知っておくべきCUDAとCPUのメモリのやりとり

メモリ確保
- cudaMallc
  - GPU上で確保
- cudaMallocHost
  - CPU上で確保
  - ページロックしてくれるので画像のような大規模データとの相性が良い
メモリコピー
- cudaMemcpy
  - 最後の引数には cudaMemcpyHostToDeviceか cudaMemcpyDeviceToHost (もしくは cudaMemcpyDeviceToDevice) を渡す
メモリ解放
- cudaFree

CPU上のメモリに値を書き込み、GPUに渡してからCPUに戻すサンプルコードは次のとおりです。

#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <npp.h>
using namespace std;


int main(void) {
    int N = 1024;
    int *a, *b, *c, *a_d, *b_d, *c_d;
    // ホスト側
    cudaMallocHost(&a, N * sizeof(int));
	  cudaMallocHost(&b, N * sizeof(int));
    // デバイス側
    cudaMalloc(&a_d, N * sizeof(int));

    for (int i = 0; i < N; i++){
        a[i] = i;
    }

    cudaMemcpy(a_d, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(b, a_d, N * sizeof(int), cudaMemcpyDeviceToHost);

    for (int i = 0; i < N; i++) cout << b[i] << ", ";
    cout << endl;

    cudaFree(a);
    cudaFree(a_d);
    cudaFree(b);

    return 0;
}

NPPのインターフェースについて

NPPには同じ処理でも、データの種類によって様々な関数が存在します。

uint8型で、3 channel (RGB)のデータであれば、接尾語が _8u_C3Rになります。1 channelであれば、 _8u_C1Rになります。多くのケースでは、この2つのどちらかになると思います。関数名の違いについては、後半の「NPPを使う上でのTips」にまとめました。

動かしてみる

TURINGの自動運転ソフトウェアでも使っている「warpPerspective」変換を使った例を紹介します。「warpPerspective」変換とは、画角を補正する変換です。位置を変えても同じような画角になるように変換することで、カメラの位置にかかわらず推論ができるようにしています。この変換は一般に「ホモグラフィー変換」と呼ばれています。（ホモグラフィー変換については、連載の第1回でも解説しています！気になる方はぜひ）

warpPerspective関数で画像をホモグラフィー変換するコードです。それぞれopencv.pngとnpp.pngという画像ファイルを出力します。

本記事で動かすコードはここのリポジトリにまとめています。

サンプルコード

OpenCV

#include <opencv2/imgcodecs.hpp>
#include <opencv2/opencv.hpp>

int main(int argc, char *argv[]) {
  // GpuMat
  cv::Mat src = cv::imread("turing_white.png");
  cv::Mat dst;

  float matrix[9] = {9.45135927e-01, -4.92482404e-02, -9.16291224e+01,
                     1.86556287e-02, 9.08238651e-01,  1.29333648e+01,
                     1.78247084e-05, -4.62799593e-05, 9.97536602e-01};

  cv::Mat M = cv::Mat(3, 3, CV_32F, matrix);
  cv::warpPerspective(src, dst, M, cv::Size(src.cols, src.rows),
                      cv::INTER_LINEAR, cv::BORDER_CONSTANT, 1);

  cv::imwrite("opencv.png", dst);
  return 0;
}

NPP

#include <npp.h>
#include <opencv2/imgcodecs.hpp>

int main() {
  // 変換する画像
  cv::Mat src = cv::imread("turing_white.png");
  // 変換された画像が入る変数
  cv::Mat dst;

  // GPU上のメモリ確保
  uint8_t *cudaSrc, *cudaDst;
  cudaMalloc((void **)&cudaSrc, src.rows * src.cols * 3);
  cudaMalloc((void **)&cudaDst, src.rows * src.cols * 3);

  // 画像をGPUにメモリコピー
  cudaMemcpy(cudaSrc, src.datastart, src.rows * src.step,
             cudaMemcpyHostToDevice);

  // 関数の引数の宣言
  NppiSize imSize; // 画像のサイズ
  imSize.width = src.cols;
  imSize.height = src.rows;
  // Roi: Reagion of Interestの略。画像サイズと同じでよい。
  NppiRect SrcRoi = {0, 0, src.cols, src.rows};
  NppiRect DstRoi = {0, 0, src.cols, src.rows};

  // 変換行列の宣言
  double matrix[3][3] = {{9.45135927e-01, -4.92482404e-02, -9.16291224e+01},
                         {1.86556287e-02, 9.08238651e-01, 1.29333648e+01},
                         {1.78247084e-05, -4.62799593e-05, 9.97536602e-01}};

  nppiWarpPerspective_8u_C3R(cudaSrc, imSize, src.step, SrcRoi, cudaDst,
                             src.step, DstRoi, matrix, NPPI_INTER_LINEAR);

  // GPU上の画像をcpuに戻す
  uint8_t *cpuDst = (uint8_t *)malloc(src.rows * src.step);
  cudaMemcpy(cpuDst, cudaDst, src.rows * src.step, cudaMemcpyDeviceToHost);
  cv::Mat nppDst = cv::Mat(src.rows, src.cols, src.type(), cpuDst, src.step);

  cv::imwrite("npp.png", nppDst);

  return 0;
}

動かすと、画像の画角が変換されて画像が出力されます。下図のとおり、NPPでもOpenCVと同様の出力が得られることが確認できました。

OpenCVとNPPの比較実験

1920x1080の画像に対してホモグラフィー変換を行う処理を100ループ行い、その処理にかかった時間を計測します。

実行環境

Jetson TX2
コンパイラ
- clang++

実行コード

OpenCV

#include <opencv2/imgcodecs.hpp>
#include <opencv2/opencv.hpp>
#include <stdio.h>

int main(int argc, char *argv[]) {
  // GpuMat
  cv::Mat src = cv::imread("turing_white.png");
  cv::Mat dst;

  std::chrono::system_clock::time_point start, end;
  double elapsed;

  float matrix[9] = {9.45135927e-01, -4.92482404e-02, -9.16291224e+01,
                     1.86556287e-02, 9.08238651e-01,  1.29333648e+01,
                     1.78247084e-05, -4.62799593e-05, 9.97536602e-01};

  cv::Mat M = cv::Mat(3, 3, CV_32F, matrix);
  start = std::chrono::system_clock::now();
  for (int i = 0; i < 100; i++) {
    cv::warpPerspective(src, dst, M, cv::Size(src.cols, src.rows),
                        cv::INTER_LINEAR, cv::BORDER_CONSTANT, 1);
  }
  end = std::chrono::system_clock::now();
  elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(end - start)
                .count();
  printf("opencv 100 loop: %lf \n", elapsed);

  return 0;
}

NPP

#include <npp.h>
#include <opencv2/imgcodecs.hpp>

#include <stdio.h>

int main() {
  // 変換する画像
  cv::Mat src = cv::imread("turing_white.png");
  // 変換された画像が入る変数
  cv::Mat dst;

  std::chrono::system_clock::time_point start, end;
  double elapsed;

  // GPU上のメモリ確保
  uint8_t *cudaSrc, *cudaDst;
  cudaMalloc((void **)&cudaSrc, src.rows * src.cols * 3);
  cudaMalloc((void **)&cudaDst, src.rows * src.cols * 3);

  // 画像をGPUにメモリコピー
  cudaMemcpy(cudaSrc, src.datastart, src.rows * src.step,
             cudaMemcpyHostToDevice);

  // 関数の引数の宣言
  NppiSize imSize; // 画像のサイズ
  imSize.width = src.cols;
  imSize.height = src.rows;
  // Roi: Reagion of Interestの略。画像サイズと同じでよい。
  NppiRect SrcRoi = {0, 0, src.cols, src.rows};
  NppiRect DstRoi = {0, 0, src.cols, src.rows};

  // 変換行列の宣言
  double matrix[3][3] = {{9.45135927e-01, -4.92482404e-02, -9.16291224e+01},
                         {1.86556287e-02, 9.08238651e-01, 1.29333648e+01},
                         {1.78247084e-05, -4.62799593e-05, 9.97536602e-01}};

  start = std::chrono::system_clock::now();
  for (int i = 0; i < 100; i++) {
    nppiWarpPerspective_8u_C3R(cudaSrc, imSize, src.step, SrcRoi, cudaDst,
                               src.step, DstRoi, matrix, NPPI_INTER_LINEAR);
  }
  end = std::chrono::system_clock::now();
  elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(end - start)
                .count();
  printf("npp 100 loop: %lf \n", elapsed);

  return 0;
}

実験結果

NPPがOpenCVよりも10倍以上速く動作しました！

OpenCVでは1処理あたり28ms程度かかっていたので、たかだか35FPSしか出ないことがわかります。一方で、NPPにすることで2.5msに抑えられることがわかりました。実用的に大きな進化をしたことがわかりました。

NPPを使う上でのTips

ここからは余談です。NPPについては、日本語はおろか、英語ですら解説記事がほぼないので参考になるものを共有します。

CUDAの基本について知りたい

NPPを使う上で、CUDAの基本的知識は絶対に必要になるので、CUDA programmingについて知りたい人は次のNVIDIAがGTCで発表したPDFを読むといいです。

このPDFでは

GPUのアーキテクチャ
CPUとGPUのメモリのやり取り
プログラムの実行方法

が学べます。これだけ知れば、NPPは扱えます。

NPPドキュメントをみても、何の関数を使ったらいいかわからない

試しにwarpPerspective関数のドキュメントを開いてみると、同じ処理だけで30個以上関数があります。

しかし、これらは次のような接尾語で区別されています。

データ型(8u, 16u, 32s, 16f, 32f)
チャネル数(C1R, C3R, C4R, P3R, P4R, AC4R)
- RGBならC3R
- グレースケールならC1R
非同期処理(Ctx)
- streamを渡す

OpenCVはMatと呼ばれる画像データ以外にも型やチャネルの情報が付与されたデータを関数に渡すので、1つの関数でも対応できます。一方NPPでは、画像データはただのポインタなので、型やチャネルの違いを柔軟に受け入れられません。

デバッグをしたい

関数の返り値が NppStatus 型のエラーコードなので、これを表示させることでデバッグが可能になります。

NppStatus status = nppiWarpPerspective_8u_C3R(cudaSrc, imSize, src.step, SrcRoi, cudaDst, src.step, DstRoi, cudaM, NPPI_INTER_LINEAR);
printf("status code = %d\n", status);

エラーコードはNPPドキュメントの次のページに記載されています。0が正常です。

OpenCVのCUDAバージョンはあるの？

あります。

OpenCV Contribとは、OpenCVの拡張版で、GPUを使った画像処理が可能になるモジュールです。こちらは、OpenCVとの互換性があるので導入して実装まで手軽にできます。

しかしながら、我々はOpenCV Contribを採用していません。
理由は2つあります。

OpenCVとの互換が不完全
- 対応していない関数が散見される
CUDAで独自関数を扱いたい

OpenCVには実装されていないような特殊な処理をしたいとき、CUDA Kernelを自前で用意してあげればボトルネックなしで高速動作に期待できる点が嬉しいです。

余談ですが、NPPによって弊社で開発しているソフトウェアからOpenCVごと消えてしまいました。

おわりに

完全自動運転車のためには、機械、ロボット、AIなど多くの分野のエンジニアが必要です。TURINGは自動運転の実現とそれによる社会の未来を信じています。システム開発だけでなく車両開発も進めており、低レイヤーやハード領域の開発も急ピッチで進んでいます。「現代最高のエンジニアリング課題の一つ完全自動運転を一緒に解きたい」「車をゼロから作り、公道を走る感動をチームで味わいたい」そんな方はぜひご応募ください。

著者（私）のTwitterはこちらから〜 DM開放してるのでいつでもご相談ください！

問い合わせ先：弊社求人一覧およびWantedlyをご覧ください。また、 info@turing-motors.com 宛にお問い合わせいただけます。その他ご質問や気になる点がありましたら、お気軽にTwitterのDMをお送りください。共同代表山本・青木どちらもDMを開放しております。→ @issei_y, @aoshun7

Tech Blog - TuringPublication

LV5完全自動運転EVを開発するAIテクノロジースタートアップ・Turing株式会社の公式テックブログです。

はじめに

自動運転技術の裏側

NPPとは

なぜNPPが必要なのか

誰にオススメできるか

知っておくべきCUDAとCPUのメモリのやりとり

NPPのインターフェースについて

動かしてみる

OpenCVとNPPの比較実験

実行環境

実験結果

NPPを使う上でのTips

CUDAの基本について知りたい

NPPドキュメントをみても、何の関数を使ったらいいかわからない

デバッグをしたい

OpenCVのCUDAバージョンはあるの？

おわりに

Discussion