本文へ移動

Turing TechTalk#43 エッジでリアルタイムを成立させる最適化 ─ 自動運転パイプラインの性能設計と実践

この記事に登場する人
CTO
山口 祐 Yu Yamaguchi
産業技術総合研究所 研究員/米国NIST客員研究員として研究する傍ら、独自にゲームAIの深層学習の開発を開始。日本の囲碁AIプロジェクトの開発代表として、最大1100GPUの並列分散強化学習を設計・開発し、世界大会準優勝、将棋AIでも世界大会優勝などの実績がある。HEROZ株式会社 執行役員を経て、2022年、チューリングに創業メンバーとして参画。2024年8月、CTO就任。基盤AIチームマネージャー兼任。
Driving System3チーム マネージャー
鈴木 達矢 Tatsuya Suzuki
Canonで画像処理や3次元計測、ロボットビジョンに取り組み、日産で量産に関わるシステム開発を経験。SenseTime japanで、モジュールベースの自動運転システム全体の開発に携わる。チューリングでは、E2E自動運転の開発においてシステム全体を俯瞰してさまざまな技術イシューに取り組んでいる。

はじめに

チューリングが開発する完全自動運転AIは、車載のエッジ計算機上でリアルタイムに動作しています。学習だけでなく、実際に車の中で毎秒10回、100ミリ秒以内に画像取得から推論、車両制御までの一連の処理を終える必要があり、この極めてタイトな制約のもとでいかに高速化・最適化を行うかが大きな技術的課題になっています。

今回のTech Talkでは、CTOの山口祐と、Driving System3チームのチームマネージャーである鈴木達矢が登壇。チューリングの自動運転パイプラインを題材に、ゼロコピー設計、CUDAストリームの活用、TensorRTによるレイヤー・カーネル融合、演算精度の使い分け、そしてNsight Systemsを用いたボトルネック解析まで、エッジ環境で“リアルタイム”を成立させるための技術的な工夫を詳しく紹介しました。

※本記事は「Turing Tech Talk #43」の内容をもとに、一部編集のうえお届けします。

山口:皆さん、こんにちは。Turing Tech Talk 第43回「エッジでリアルタイムを成立させる最適化 ── 自動運転パイプラインの性能設計と実践」を始めていきたいと思います。私はTuringのCTOの山口です。本日はDriving System3チームのチームマネージャーである鈴木達矢さんに来てもらっています。鈴木さん、今日はよろしくお願いします。

鈴木:よろしくお願いいたします。

山口:今日は、自動運転AIをエッジ計算機(車載コンピュータ)上でリアルタイムに動かすための高速化・最適化がテーマです。AIには学習と推論(学習済みモデルを実際に動かす処理)の2つのフェーズがありますが、今回は推論側の話が中心になります。それでは鈴木さん、簡単に自己紹介をお願いします。

鈴木:2009年、新卒でキヤノンに入社し、画像処理やロボットビジョンの研究開発に取り組みました。その後デンソーで自動運転開発に携わり、車両業界に飛び込みました。その後SenseTime japanでモジュールベースの自動運転システム開発に従事し、現在はチューリングで推論周りの開発を担当しています。在籍は1年半になります。

山口:私はチューリングのCTOを務めています。TuringTechTalkは、チューリングの最新の研究開発内容を担当エンジニアが直接解説するオンラインイベントです。今日はエッジ環境での最適化ということで、GPUやCUDA関連の話が中心になるかと思います。チューリング株式会社は2021年8月設立、本社は東京都平和島(羽田空港近くの物流倉庫街の一角)にあります。累計資金調達額は最近更新されまして365億円、社員数はまもなく100人です。事業内容は完全自動運転システムの開発で、どんな場所・シチュエーションでも人間の介入なしに運転できる、いわば究極の自動運転を指しており、こうした開発を進めるディープテックスタートアップとして日々取り組んでいます。それでは、ここから鈴木さんに説明してもらいます。

リアルタイム推論を支えるエッジ計算機

鈴木:今日は、1つのエッジデバイス上でリアルタイムに処理させるためのパイプラインの設計方法と、動かしていく中で生じるボトルネックの解析方法について解説していきます。

我々が使用しているのはNVIDIA DRIVE AGX Orin(車載用の計算機。内部にJetson AGX Orinを搭載)です。推論だけでなく、カメラやGNSS(衛星測位システム)、車載のCAN(Controller AreaNetwork、車両内の通信プロトコル)情報の取得もあわせて行っています。

デバイス上では複数のプロセスが協調して動作しています。カメラ画像や車両情報、ナビ情報を取り込み、学習済みモデルがそれを受け取って推論し、車のコントローラーに制御信号として出力します。これを10Hz(1秒間に10回、つまり100ms以内に一連の処理を終える必要がある)というリアルタイム制約のもとで動かしており、制御と並行して各種センサーデータや通信ログをストレージに保存しています。

山口:この計算機はNVIDIA製で、Ampere世代(NVIDIAのGPUアーキテクチャの1つ。現行最新はBlackwell、その前がHopperで、Ampereはさらに前の世代)の車載用に特化したチップです。スペックは約2650TOPS(Tera Operations Per Second、演算性能の指標)で、データセンター向けGPUには劣りますが、身近な例で言うとNintendo Switch 2のSoC(NVIDIA製カスタムチップ)に近い性能とイメージしていただければと思います。このパイプラインで重くなりやすいのは、カメラの画像処理とモデルの推論処理だと聞いています。10Hz・100msという制約は、たとえば言語モデルをデータセンターのサーバー上で推論サービスとして提供するような場合とは性質が異なり、少しでも遅延すればそのまま車の制御が遅れてしまう、非常にタイトなリアルタイム性が求められる領域です。

鈴木:はい、今日の話の中心もそこになります。

ボトルネックの所在とゼロコピー設計

山口:先日、ワールドカップ観戦から帰国されたばかりの本田 圭佑さんに、弊社の自動運転車で空港から送迎させていただいた動画をチューリングのYouTubeチャンネルでご紹介しました。よくできた自動運転というのは、乗っていても「自動運転をしている」ことを意識させないくらい自然になるものです。運転が上手なタクシーの運転手ほど印象に残らないのと同じで、安定して自然であることこそが良い自動運転だと感じています。

鈴木:システム設計で苦労するのは、カメラが5メガピクセルという比較的高画質であるため演算量が大きくなり、大容量画像のリアルタイム処理が課題になる点です。画像をコピーするとCPU負荷や転送遅延が増え、フレームがドロップしてしまうと、モデルは画像が順序通り届くことを前提に動いているため、車両の挙動が乱れる原因になります。それを防ぐには、コピーを発生させず同じメモリ領域からポインタで受け渡す「ゼロコピー」の実装が重要になります。

山口:計算機の内部では、ストレージからメモリへの展開、メモリからCPUレジスタへの転送など、随所でコピーが発生します。特に今回のDRIVE AGX OrinはCPUとGPUが一体化したSoCですが、内部では処理系統が分かれているため、CPUとGPUの間でデータをやり取りするたびにオーバーヘッドが生じます。

8台のカメラの映像はそのままAI推論に使うだけでなく、記録のために動画としてエンコード(圧縮)してストレージに保存する処理も同時に走っています。スマートフォンで長時間動画を撮影しているとカメラ部分が熱くなることがありますが、あれは撮影した信号をリアルタイムに動画へエンコードする処理が走っているためです。カメラ1台でもそれだけの負荷がかかるところを、8台分を同時にこなす必要があり、しかもメモリ・CPU・GPUといったリソースはすべて推論処理と共有しているため、計算機のリソースは非常にカツカツな状態になります。

鈴木:まさにそうですね。だからこそ、できる限り効率的にリソースを使い切る設計が求められます。このボトルネックを解消するため、カメラ画像の前処理から推論までを同一のCUDA(NVIDIAが提供するGPU向け並列計算プラットフォーム・開発キット)プロセスに集約し、GPUから直接参照可能なメモリ上でデータを受け渡すことで、不要なコピーとプロセス間の同期待ちを削減しています。あわせて前処理・推論・記録処理の順序をうまく制御し、リソースの競合を避けることもポイントです。

CUDAストリームとTensorRTによる高速化

鈴木:CUDAでは「ストリーム」という単位でGPU上の処理順序を制御します。CPUが1つずつ命令を出してはGPUの処理完了を待つ、という進め方をすると待機時間が無駄になってしまうため、複数の命令をまとめてストリームに投入し、GPUを常に稼働させ続けること、そしてCPU・GPU間の同期はできる限り減らすことがポイントです。

山口:なるほど。これはたとえるなら、指示を出す人と実際に働く人が2人いて、その都度「これ終わりました、次は何をしますか」と確認を挟んでしまうと仕事が進まない、という状況に近いですね。朝のうちに今日やることをまとめて指示しておき、休みなく作業を進めてもらった方が、1日のアウトプットは大きくなります。

鈴木:まさにそのとおりで、GPUを絶えず働かせ続けることが重要になります。学習済みモデルの推論には、TensorRT(NVIDIAが提供する推論高速化ライブラリ・SDK)を使用しています。TensorRTは学習済みモデルの計算グラフを解析し、畳み込みや活性化関数など連続する演算を1つのGPUカーネルにまとめる「レイヤー・カーネル融合」を行うことで、GPU起動のオーバーヘッドやメモリアクセスを削減します。また、入力サイズやGPUの種類に応じて最適なカーネルサイズを自動的に探索するチューニング機能も備えています。

山口:GPUはもともとグラフィックス、つまり映像処理のために作られたもので、大量の演算器を並列に並べ、色データなどを一斉に処理することを得意としています。その並列処理の仕組みをAIの推論にも活用しているのが今のGPUです。ハードウェアの特性を踏まえてソフトウェア側から最適な形に変換してくれるのがTensorRTというわけです。CUDAやTensorRT自体は10年以上の歴史がありますが、当初は数倍程度だった高速化効果が、NVIDIAによるソフトウェアとハードウェア双方の作り込みによって、現在では大幅に向上しています。

演算精度とゼロコピー実装の詳細

鈴木:あわせて重要なのが演算精度です。モデルの学習時にはFP32(32ビット浮動小数点)を使いますが、推論時にはその半分のFP16(16ビット)や、さらに低精度なINT8(8ビット整数)に精度を落とすことで、演算に使うメモリを削減し、データの取り回しを速くできます。低精度化は精度とのトレードオフになりますが、コンピュータは整数演算を得意とするため、INT8化(量子化)は高速化に大きく寄与します。TensorRTはこの量子化のチューニングも自動的に行ってくれます。

鈴木:ゼロコピーの実装面では、通常のcudaMalloc(GPUメモリを確保するCUDAの基本関数)は仮想アドレスの予約と物理メモリの割り当てを同時に行ってしまうため、これらを分離できる低レベルAPI(バーチャルメモリマネジメント)を用いています。これにより、プロセス間でIPC(Inter-Process Communication、プロセス間通信)を介してGPUメモリをそのまま共有する仕組みを実装しています。あわせて、過去の画像や推論結果を再利用するために、タイムスタンプ付きのリングバッファ(あらかじめ確保したバッファ数を使い回す仕組み)で時系列データを蓄積しています。画像処理と推論はCUDAストリームを共有し、同期は推論が完了した1回だけにとどめることを徹底しています。

Nsight Systemsによるボトルネック解析

鈴木:ボトルネックの解析には、Nsight Systems(NVIDIA製のプロファイリングツール)を使用しています。これはCPU、GPU、OS、DLA(Deep Learning Accelerator、GPU以外の演算アクセラレータ)など、すべてのプロセッサの処理タイミングを可視化してくれるツールです。

山口:実際のプロファイルを見ると、画像の前処理とGPUでの推論処理の間に隙間があり、GPUが遊んでしまっている状態が見つかりました。

鈴木:原因は、CPU側が画像処理の完了を待ってから次の命令を出していたことでした。また前処理のブロックが長くなっている箇所では、不要な画像コピーが発生していました。パイプライン全体では104ms程度かかっており、そのうち推論部分(FP32使用時)だけで96ms程度を占めていました。これは100msというリアルタイム制約をわずかに超えてしまう水準です。そこで、CPU・GPU間の同期の削減、不要なコピーの削減、そしてTensorRTでのFP16化を行った結果、104msだった処理時間は60ms未満まで短縮できました。FP16程度の精度低下であれば、運転の制御性能への影響はないこともこれまでの検証で確認しています。

山口:この手のプロファイル分析は、まさにエンジニアの間でよく言われる「推測するな、計測せよ」という格言そのものですね。専門知識があるとつい「こうだろう」と感覚で議論しがちですが、結局コンピュータの速度や性能を決めるのは実際に計測した結果でしかありません。自動運転そのものについても同じで、こういうアーキテクチャなら運転が良くなるはずだ、というのはあくまで仮説であり、実際に車の上、計算機の上で動かして検証して初めて確かめられます。

山口:モデルは大きいほど賢くなる一方で、パラメータ数が多いほど演算コストも増えるというトレードオフがあります。パイプラインを高速化して生まれた余裕を、より大きく賢いモデルの搭載に充てるという積み重ねを重ねてきた結果、これまで全体として2~3倍程度の高速化を実現してきました。かつてはギリギリでしか動かなかった処理も、今では余裕を持って動かせるようになっています。

チーム体制と今後

鈴木:このチームは現在、私を含めて3人体制です。もう1人は監視カメラなど別のデバイスで同様のノウハウを培ってきたメンバー、もう1人は新卒からチューリングに入りずっとこのシステムに関わってきたメンバーです。チューリングとしてもまだ5年ほどのスタートアップですが、若いメンバーと経験のあるメンバーがともに1つの課題に取り組んでいます。

山口:鈴木さん、今日はありがとうございました。非常に勉強になりました。次回のテックトーク第44回は、1ヶ月後の8月12日(水)18時から、「車両の制約下で成立させる自動運転 ── 機構と統合設計」というテーマでお届けします。次回は、いま我々が開発している自動運転車、つまり市販車にセンサーや計算機を追加改造して作っている車両について、その改造・設計・メンテナンスを支えるハードウェア技術を扱います。Driving System2チームのシニアエンジニアである山岡さんに登壇いただき、走行系・電装系をどのように設計し施工しているかを解説する予定です。

山口:鈴木さん、今日はいかがでしたか?

鈴木:今日は自分なりにわかりやすく説明したつもりでしたが、振り返るとやや難しくなってしまった部分もあり、反省点があるかなと思います。また機会があれば、皆さんにもっとうまく解説できるようブラッシュアップしていきたいと思います。

山口:今日は、ありがとうございました。

Q&A一覧(一部抜粋)

以降は視聴者のQAに回答していきました。詳しくは動画をご覧ください。

  • この車載計算機(DRIVE AGX Orin)はどれくらいの金額なのでしょうか?
  • ログを撮りながら10Hzで動くのがすごいです。ログはどこまで撮っているのですか?
  • システム設計とアルゴリズムは何が異なるのでしょうか?リアルタイムにするためにどういったアルゴリズムを使っているのか気になりました。
  • CPU・GPU間の同期はプロセス間の依存関係を解決するために必要だと思います。仮想メモリ部分から実装し直す必要があったと思いますが、難しかったことはありましたか?
  • 通常、学習モデルにはPyTorchなどを使うと思いますが、推論部分はどの言語を使われていますか?CUDAライブラリを直接コールしているのでしょうか?
  • 初心者でもこういった高速化のためのシステム設計を行いたいと思ったら、とっつきやすい材料などはありますか?
  • INT8を使うことでモデルを高速化するとのことですが、運転挙動やモデル精度にはどれほど影響があるのでしょうか?

チューリングでは、完全自動運転の技術を共に創る仲間を募集しています。今日お話ししたDriving Systemチームはもちろんのこと、機械学習エンジニア、リサーチャー、ソフトウェアエンジニア、組み込みエンジニア、インフラエンジニアなど、非常に幅広いエンジニア職種で仲間を募集しています。ご興味のある方は、ぜひ採用ページをご確認ください。多様な職種がありますので、ご自身がどれに当てはまるか、ぜひチェックしてみてください。