Jetson活用　ROS2自律走行実現に向けて～自律移動ロボット用ー自己位置推定のCUDA高速化～（2021/11/04）

10.7K Views

November 04, 21

jetson ros2 cuda 自律走行 cuda高速化 gpu ros2シリーズ cuda高速化シリーズ高速化シリーズ

スライド概要

Jetsonを用いたAGV/AMR 開発をしているが、ROS2への移行の仕方に悩んでいる方や、Jetsonが十分に生かし切れていないと考えている方向けのオンラインセミナーです。

メカナムローバーを題材にROS2とJetsonを用いての搬送ロボットの組み上げ方法や、自己位置推定アルゴリズムのCUDA高速化方法を紹介します。

＜過去資料＞
・ROS2自律走行実現に向けて 1 ：　https://www.docswell.com/s/fixstars/57VGDN-20221017
・ROS2自律走行実現に向けて 2：　https://www.docswell.com/s/fixstars/K8G1N9-20221130
・Jetson活用　ROS2自律走行実現に向けて～自律移動ロボット用ー自己位置推定のCUDA高速化：　https://www.docswell.com/s/fixstars/KQ814X-20211104

・フィックスターズの自動車向けソフトウェア開発：　https://at.fixstars.com/ja

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

@fixstars

スライド一覧

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

またはPlayer版

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

ダウンロード

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

関連スライド

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

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

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

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

arm neon

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

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

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

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

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

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

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

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

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

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

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

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

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

各ページのテキスト

Fixstars Corporation www.fixstars.com 本日のアジェンダ • フィックスターズ会社紹介 • 会社概要と各種サービス • テクニカルセッション • ros1_brigeを用いた搬送ロボットメカナムローバーのROS2対応 • 自己位置推定アルゴリズムamclの高速化（プロファイリング、CUDA化） • Q＆A Copyright © Fixstars Corporation 2

Fixstars Corporation www.fixstars.com 発表者紹介 • 山田真介 (Shinsuke Yamada) • ディレクター • 博士 (工学) • 主に自動運転関係のプロジェクトマネジメントを担当 • • • 深層学習における量子化車両運動制御プログラムの高速化坂本浩平 (Kohei Sakamoto) • シニアエンジニア • プロトタイプ自律ロボットのアルゴリズム開発や CUDA 高速化、自動運転システムの構築を担当 • 高速化に興味 • アルゴリズム/SIMD/CUDA/高速コード実装 Copyright © Fixstars Corporation 3

Fixstars Corporation www.fixstars.com ソフトウェア高速化サービス概要お客様のソースコードをご提供いただき、ソフトウェアの最適化やアルゴリズムの改良を行い高速化したコードをお返しします。オリジナルソースコードのご提供当社高速化したソースコードコンサルティング高速化お客様サポート性能評価アルゴリズムの改良・開発レポートやコードへのQ&A ボトルネックの特定ハードウェアへの最適化実製品への組込み支援レポート作成 Copyright © Fixstars Corporation 6

Fixstars Corporation www.fixstars.com ソフトウェア高速化サービス領域大量データの高速処理がお客様の製品競争力の源泉となる、様々な領域でソフトウェア開発・高速化サービスを提供しています。 Semiconductor Industrial ・NAND型フラッシュメモリ向けファー・Smart Factory化支援ムウェア開発・マシンビジョンシステムの高速化・次世代AIチップ向け開発環境基盤開発 Mobility Life Science ・自動運転の高性能化、実用化・ゲノム解析の高速化・次世代パーソナルモビリティの研究開発・医用画像処理の高速化・AI画像診断システムの研究開発 Finance ・デリバティブシステムの高速化・HFT(アルゴリズムトレード)の高速化 Copyright © Fixstars Corporation 7

Fixstars Corporation www.fixstars.com Fixstars Autonomous Technologies 高速化技術のフィックスターズと豊田通商グループのネクスティエレクトロニクスによる合弁会社並列化・最適化・機械学習・コンピュータビジョン技術を駆使し、自動運転社会の実現を加速する会社名： Fixstars Autonomous Technologies 設立：2018年2月1日資本金：3,000万円（株）フィックスターズ 66.6% （株）ネクスティエレクトロニクス 33.4% 所在地：東京都港区芝浦3-1-1 代表取締役会長 : 蜂須賀利幸代表取締役社長：羽田哲 Copyright © Fixstars Corporation 8

Fixstars Corporation www.fixstars.com 画像処理・アルゴリズム開発サービス • お客様の課題 • 高度な画像処理、深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件までクリアできるか不安研究開発の成果が製品化にうまく結びつかない • • 弊社の支援内容 • • • 課題に応じたアルゴリズム調査深層学習ネットワーク精度改善、推論高速化手法調査論文調査、実装 https://www.cs.toronto.edu/~frossard/post/vgg16/ Copyright © Fixstars Corporation 9

https://www.cs.toronto.edu/~frossard/post/vgg16/

10.

Fixstars Corporation www.fixstars.com AI・深層学習関連サービス • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績ネットワーク設計データの前処理、データ拡張精度改善分散処理による学習高速化各種DLフレームワークターゲットデバイスへのポーティング及び推論高速化 ◼ Visconti, ARM, GPU, DSP ◼ SIMD,NEON,CUDA,TensorRT モデル圧縮 - 量子化 - 枝刈り - 蒸留クラウド・サーバ Copyright © Fixstars Corporation エッジ 10

11.

Fixstars Corporation www.fixstars.com GPU向け高速化サービス • お客様の課題 • GPU 高速化の知見がない • 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • GPU 高速化に関するコンサルティング • • ボトルネック調査、GPU プログラムの高速化 CPU/GPU が混在するヘテロジニアス環境での最適化 Copyright © Fixstars Corporation 10～150 倍の高速化事例あり 11

12.

Fixstars Corporation www.fixstars.com AGV / AMR 開発における課題と当社サービス当社サービス開発課題量産機を見越して… 本セミナーで紹介 • 製造原価を抑えたい • 計算機資源の選定 • 組込み規模の限られた計算資源の中で性能を出したい • ロボット試作品の組み上げ • 複数台のロボットにも対応したい • ロボットの動作確認・評価後継機に向けて… • • 実製品向けの開発プラットフォームに変更したい計算資源はそのまま性能を落とさずに機能追加したい • ロボット開発プラットフォームの適用・移植 • アルゴリズム開発・改良 • 処理時間短縮・高速化 Copyright © Fixstars Corporation 12

13.

14.

15.

Fixstars Corporation www.fixstars.com 1. はじめに「ROS2上でナビゲーション処理を行う車両型ロボット開発環境を整備し、ROS2のナビゲーション処理を高速化したい」という動機から、今回、以下のような取り組みを行った。 • メカナムローバーのROS2対応 • ros1_bridgeによってROS1資産を活用しながらのROS2対応 • 自己位置推定パッケージamclのCUDA高速化 • ROS2ノードのプロファイリング（perf、NVIDIA Nsight Systems） • amclのCUDA実装およびプロファイリング（nvprof、NVIDIA Visual Profiler） Copyright © Fixstars Corporation 15

16.

17.

18.

Fixstars Corporation www.fixstars.com 2.1 メカナムローバーとは • ヴイストン株式会社が販売する4輪のメカナムホイールを搭載した全方位移動台車ロボット[1] • 可搬重量約40kg、最高速度1.3m/s • 2021/11/1時点でROS1に対応 • 今回の一連の取り組みではメカナムローバーVer2.0を用いた [1] https://www.vstone.co.jp/products/wheelrobot/ver2.0.html Copyright © Fixstars Corporation 18

https://www.vstone.co.jp/products/wheelrobot/ver2.0.html

19.

Fixstars Corporation www.fixstars.com 2.1 メカナムローバーとは • ハード構成は以下の通り（購入時オプションよって構成が異なる）。メカナムローバー本体 2D LiDAR USB URG-04LX-UG01[2] 測距範囲：0.02～5.6m 走査角度：240° USB PC DCモーター I2C Arduino互換 CPUボードモータードライバエンコーダ [2] https://www.hokuyo-aut.co.jp/search/single.php?serial=17 Copyright © Fixstars Corporation 19

https://www.hokuyo-aut.co.jp/search/single.php?serial=17

20.

Fixstars Corporation www.fixstars.com 2.1 メカナムローバーとは • LiDARをVelodyne VLP-16に変更した • Jetson AGX Xavierを用いた ✓ 測距範囲を広くし、広域のフィールドでも自己位置推定を行いやすくしたい ✓ 今後、3D SLAM検討もできるようにしたいメカナムローバー本体 3D LiDAR Velodyne VLP-16[3] 測距範囲：0.5～100m Ethernet 水平FOV：360° 垂直FOV：30°（±15°） USB PC DCモーター I2C Arduino互換 CPUボードモータードライバエンコーダ Jetson AGX Xavier [3] https://velodynelidar.com/products/puck Copyright © Fixstars Corporation 20

https://velodynelidar.com/products/puck

21.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • 2021/11/1時点でROS1のみサポートされている • 公式で動作検証されているナビゲーション時のROS1ノード構成は以下の通り。 2D LiDAR urg_node キャプチャ scan move_base 地図配信 map_server map オドメトリ pub_odom 変換 odom 自己位置推定 amcl amcl_pose 経路生成、走行制御 dwa_local_ planner rover_twist odom rover_odo Copyright © Fixstars Corporation 21

22.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • 2021/11/1時点でROS1のみサポートされている • 公式で動作検証されているナビゲーション時のROS1ノード構成は以下の通り。 urg_node scan move_base map_server map rover_twist amcl_pose amcl dwa_local_ planner odom pub_odom rover_odo odom geometry_msgs/Twistから nav_msgs/Odometryに変換 Copyright © Fixstars Corporation 22

23.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • Velodyne VLP-16に変更したことで以下のようなノード構成となった velodyne_driver 3D LiDARキャプチャ velodyne_pointcloud points_raw velodyne_packets pointcloud_to_ 3D点群、2Dスキャン変換 laserscan move_base scan map_server 地図配信 map オドメトリ pub_odom 変換 odom amcl 自己位置推定 amcl_pose 経路生成、走行制御 dwa_local_ planner rover_twist odom rover_odo Copyright © Fixstars Corporation 23

24.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • Velodyne VLP-16に変更したことで以下のようなノード構成となった velodyne_pointcloud velodyne_driver points_raw velodyne_packets pointcloud_to_ laserscan move_base scan map_server map rover_twist amcl_pose amcl dwa_local_ planner odom pub_odom odom rover_odo Copyright © Fixstars Corporation 24

25.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • 速度指令（rover_twistトピック）からモーター制御指令までの処理の流れメカナムローバー本体 ③モーターを回転させる ① rover_twistトピックを配信 rover_twist PC ②（ROS1の） rover_twistトピックを購読し、モータードライバを介して制御指令を出す ※Arduinoのスケッチ実装 Arduino互換 CPUボード DCモーターモータードライバエンコーダ Copyright © Fixstars Corporation 25

26.

Fixstars Corporation www.fixstars.com 2.2 メカナムローバーのROS対応状況 • エンコーダ情報からオドメトリ（rover_odoトピック）を算出する処理の流れメカナムローバー本体 ③rover_odoトピックを購読 ④pub_odomノードでodomトピックに変換 ②モータードライバ経由で得られるモーター回転数からオドメトリを計算し、（ROS1の）rover_odoトピックを配信 ※Arduinoのスケッチ実装 DCモーター rover_odo PC Arduino互換 CPUボードモータードライバエンコーダ ①モーター回転数を取得 Copyright © Fixstars Corporation 26

27.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応 • 今回、以下のような方針でROS2対応を行った • メカナムローバーのROS2対応のファーストステップとして「amclを用いた自己位置推定まで」を行う（＝ナビゲーションまではやらない） • ros1_bridge[4]を用いることで既存のROS1資産を使いつつ、センサ関連や自己位置推定処理をROS2化する • 自己位置推定に関連する主な処理をROS2 Galactic上で動作させる • Galacticより前のROS2ディストリビューションでは使いたい機能が使えないことが多かったため [4] https://github.com/ros2/ros1_bridge Copyright © Fixstars Corporation 27

https://github.com/ros2/ros1_bridge

28.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応各ノードが動作するROS1、ROS2ディストリビューション • A：ROS1 Noetic、B：ROS2 Foxy、C：ROS2 Galactic • B points_raw velodyne_packets velodyne_driver pointcloud_to_ laserscan velodyne_pointcloud scan C A amcl_pose odom odom pub_odom ros1_bridge amcl initial_pose rviz2 map rover_odo メカナムローバー内蔵のArduinoスケッチ実装はROS1のままで動かす Copyright © Fixstars Corporation map_server 28

29.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応各ノードが動作するROS1、ROS2ディストリビューション • A：ROS1 Noetic、B：ROS2 Foxy、C：ROS2 Galactic Velodye関連のROS2パッケージはROS2 • Foxy対応済なのでFoxyで動かす B points_raw velodyne_packets velodyne_driver pointcloud_to_ laserscan velodyne_pointcloud scan C A amcl_pose odom odom pub_odom ros1_bridge amcl initial_pose rviz2 map rover_odo map_server Copyright © Fixstars Corporation 29

30.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応各ノードが動作するROS1、ROS2ディストリビューション • A：ROS1 Noetic、B：ROS2 Foxy、C：ROS2 Galactic • B points_raw velodyne_packets velodyne_driver pointcloud_to_ laserscan velodyne_pointcloud scan C A amcl_pose odom odom pub_odom ros1_bridge amcl initial_pose rviz2 map rover_odo ROS1、ROS2間の通信の相互変換 map_server 自己位置推定関連は ROS2 Galiacticで動かす Copyright © Fixstars Corporation 30

31.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応各ノードが動作するコンピュータ • A：Jetson AGX Xavier、B：ノートPC（Intel CPU） • A points_raw velodyne_packets velodyne_driver pointcloud_to_ laserscan velodyne_pointcloud Jetson AGX Xavier B scan amcl_pose odom odom pub_odom ros1_bridge amcl initial_pose rviz2 map rover_odo map_server 可視化用ノートPC Copyright © Fixstars Corporation 31

32.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応各ノードが動作するコンピュータ • A：Jetson AGX Xavier、B：ノートPC（Intel CPU） • A points_raw velodyne_packets velodyne_driver rviz2による可視化及び初期位置指定を行うために別PCを用意 pointcloud_to_ laserscan velodyne_pointcloud B scan amcl_pose odom odom pub_odom ros1_bridge amcl initial_pose rviz2 map rover_odo map_server 可視化用ノートPC Copyright © Fixstars Corporation 32

33.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応複数ディストリビューションがあるためDockerコンテナ化した • A：ROS1 Noetic、B：ROS2 Foxy、C：ROS2 Galactic、D：ROS2 Galactic（CUDA） • B points_raw velodyne_packets velodyne_driver pointcloud_to_ laserscan velodyne_pointcloud D C A amcl_pose odom odom pub_odom scan ros1_bridge amcl initial_pose rviz2 map rover_odo map_server Copyright © Fixstars Corporation 33

34.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応このDocker環境を構築するにあたり、以下のようなハマりポイントがあった • 扱うDockerコンテナが多いため手順が煩雑 • 背景：様々なROS1、ROS2ディストリビューションが混在するため • 対応：docker-composeを使うことで簡単に起動できるようした • Jetson AGX Xavier、ノートPC間でROS2の通信ができない • 課題 • • 今回のハードウェア構成では、有線LAN、Wi-Fiといった複数のNICを同時に使用するが、 CycloneDDS側で複数NICを適切にハンドリングしてくれない対応 • CycloneDDS設定で明示的にNICを設定することで対応[5] 設定例（wlan0となっている箇所はネットワークインターフェース名） export CYCLONEDDS_URI="<CycloneDDS><Domain><General><NetworkInterfaceAddress>wlan0</NetworkInterfaceAddress></General></Domain></CycloneDDS>" [5] http://www.robotandchisel.com/2020/08/12/cyclonedds/ Copyright © Fixstars Corporation 34

http://www.robotandchisel.com/2020/08/12/cyclonedds/

35.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応 Jetson AGX Xavier上で動作するCUDA、ROS2 Galactic環境構築 • 課題 • ROS2公式のDockerイメージ (ros:galactic)[6]だとROS2 Galacticは使えるが、 CUDAを使えない • NGC(NVIDIA GPU CLOUD)配布のDockerイメージ(nvcr.io/nvidia/l4t-base)[7] だとCUDAは使えるが、Ubuntu 18.04ベースであるため、ROS2 Galacticをソースコードから入れる必要あり[8] [6] https://hub.docker.com/_/ros [7] https://ngc.nvidia.com/catalog/containers/nvidia:l4t-base [8] https://docs.ros.org/en/galactic/Installation/Ubuntu-Development-Setup.html Copyright © Fixstars Corporation 35

36.

Fixstars Corporation www.fixstars.com 2.3 メカナムローバーのROS2対応 Jetson AGX Xavier上で動作するCUDA、ROS2 Galactic環境構築 • 今回の対応 • NGC配布Dockerイメージ（nvcr.io/nvidia/l4t-base）をベースにROS2 GalacticをインストールしたDockerイメージ（dustynv/ros:galactic-ros-base-l4t-r32.6.1） [9]を活用 • amclおよび依存ROS2パッケージをソースコードからビルド • ROS2 Galacticをソースコードからインストールしている都合上、aptでROS2 Galacticパッケージのインストールができないため [9] https://github.com/dusty-nv/jetson-containers Copyright © Fixstars Corporation 36

https://github.com/dusty-nv/jetson-containers

37.

38.

39.

40.

Fixstars Corporation www.fixstars.com 3.1 背景 • ナビゲーション処理に関わるROS2パッケージをCUDA実装することで、処理の高速化およびCPU負荷低減を実現したい • 以下の理由から今回はamclパッケージ（自己位置推定処理）[10]に着目した • パーティクル毎に処理が独立しており、並列度が高く、GPUと相性がよい（amclのアルゴリズム詳細は後述） • GPUにオフロードできれば自己位置推定精度を改善するためにパーティクル数を増やしやすい • navigation2でもGPU、CPU高速化対象アルゴリズムの候補として挙がっている • • • https://github.com/ros-planning/navigation2/issues/1391 https://github.com/ros-planning/navigation2/issues/1781 https://navigation.ros.org/2021summerOfCode/projects/multithreading.html [10] https://github.com/ros-planning/navigation2/tree/galactic/nav2_amcl Copyright © Fixstars Corporation 40

41.

Fixstars Corporation www.fixstars.com 3.2 amcl とは • Adaptive Monte Carlo Localization • • 自己位置推定アルゴリズム • 内界センサ、外界センサの情報からロボットの現在位置を推定する赤い点群: パーティクル外周の点群: LiDAR によるスキャン特徴 • パーティクルフィルタをベースとしたアルゴリズム • • • 位置姿勢の候補とその尤度をパーティクルで表現するパーティクルの散布範囲と数を動的に変更する自己位置が不確かな場合大域的リセットを行う Copyright © Fixstars Corporation 41

42.

Fixstars Corporation www.fixstars.com 3.2 amcl とは • amcl のおおまかな処理の流れパーティクルの初期化予測 (初期の位置姿勢が取り得る範囲にパーティクルをランダム散布) (ロボットの移動分を内界センサで取得し、ノイズを混ぜてパーティクルを移動させる) 更新 (各パーティクルについて、外界センサ情報との合致度に基づき尤度を更新) 推定 (各パーティクルの情報から位置姿勢を一意に推定) リサンプリング (尤度の低いパーティクルを除去し、尤度の大きいパーティクル付近にランダムに散布) Copyright © Fixstars Corporation 42

43.

Fixstars Corporation www.fixstars.com 3.3 amcl のプロファイリング • プロファイリング • 目的のプログラムを分析し、計算時間が大きい箇所を特定する ⇒ amcl の処理のうち、どこを高速化すべきか目星をつける amcl amcl プロファイリング〇〇[%] 重い処理〇〇[%] 〇〇[%] 〇〇[%] 〇〇[%] Copyright © Fixstars Corporation 43

44.

Fixstars Corporation www.fixstars.com 3.3 amcl のプロファイリング • 実際に amcl のプロファイリングを実施 • 説明内容 • • • • 3.3.1 3.3.2 3.3.3 3.3.4 プロファイリングの実行環境プロファイリング時のノード構成プロファイリング方法プロファイリング結果 Copyright © Fixstars Corporation 44

45.

Fixstars Corporation www.fixstars.com 3.3.1 プロファイリングの実行環境 • ハードウェア • NVIDIA® Jetson AGX Xavier™ • • • • CPU: ARMv8 Processor rev 0 (v8l)、64bit 8コア GPU: NVIDIA Volta™, CUDAコア512基 RAM: 32GB 動作モード: 30W 8コア • • 動作周波数(Max) … CPU: 1.2GHz、GPU: 905 MHz ソフトウェア • • NVIDIA® JetPack SDK 4.6 (L4T 32.6.1) (*) • OS: Ubuntu18.04 • CUDA: 10.2 コンパイラ: gcc version 7.5.0 (Ubuntu/Linaro 7.5.0-3ubuntu1~18.04) (*) JetPack : Jetson用の開発キット．L4T 32.6.1を含む L4T (NVIDIA® Jetson™ Linux Driver Package) : Jetson向けのサポートパッケージ Copyright © Fixstars Corporation 45

46.

Fixstars Corporation www.fixstars.com 3.3.1 プロファイリングの実行環境 • ROS2の動作環境 • ROSディストリビューション : ROS2 Galactic • 2021年11月現在の最新バージョン • Galactic Geochelone リリース: 2021/5/23 Dockerコンテナ内にROS2環境を構築 • Linux 18.04 (L4T 32.6.1) 上にROS2 Galacticをビルドしたコンテナ[11]を使用アプリケーションゲストOS Docker コンテナ (L4T 32.6.1) ホストOS Linux 18.04 (L4T 32.6.1) ハードウェア [11] https://github.com/dusty-nv/jetson-containers ROS2 Galactic Copyright © Fixstars Corporation Jetson AGX Xavier 46

https://github.com/dusty-nv/jetson-containers

47.

Fixstars Corporation www.fixstars.com 3.3.2 プロファイリング時のノード構成 • 以下のノード構成でamclを実行 ROS2 Galactic rosbag2 各種センサ情報、オドメトリ情報など amclの処理結果の再現性を確保するため、 rosbag2 で記録したデータを入力に使用する nav2_amcl 地図情報ステート管理 nav2_lifecycle _manager nav2_map_ server Copyright © Fixstars Corporation 47

48.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法今回実施したプロファイリングの流れ ROS2 Galactic 実行 (1) amclノード (ROS2)のビルド・実行ビルド amcl ノード分析 perf amcl 実行ファイルソースコード (2) perf[12]、FlameGraph[13]を使用して分析 FlameGraph (3) 分析結果から計算負荷の高い関数を特定する amcl::funcA1_1 amcl::funcA1 [12] perf … Linuxカーネルで使用可能な性能分析ツール [13] FrameGraph … 性能解析結果を階層図で可視化できるツール https://www.brendangregg.com/flamegraphs.html A2 B2 amcl::funcA amcl::funcA1_1 (…) { amcl ……… ; ...... ; cpu utilization [%] ……… ; …………… ; ⇒高速化の余地あり ………… ; } Copyright © Fixstars Corporation funcB 48

https://www.brendangregg.com/flamegraphs.html

49.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法 (1) amcl ノード (ROS2) のビルド・実行 ROS2 Galactic amcl ノード • 実行 Dockerコンテナを起動する • Dockerコンテナ内でのプロファイリングを有効にするため、ビルド (colcon) 実行ファイル amcl ソースコード --privileded オプションを付与して起動する $ docker run --gpus all --rm -it –-privileged container:latest • amclを含む ROS2パッケージをビルドする • プロファイリングに必要な情報を実行ファイルに付与するため、以下のオプションを指定 -DCMAKE_CXX_FLAGS="-fno-omit-frame-pointer -funwind-tables -g3" $ colcon build --cmake-args -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_CXX_FLAGS="-fno-omit-frame-pointer -funwind-tables -g3" Copyright © Fixstars Corporation 49

50.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法 (1) amcl ノード (ROS2) のビルド・実行 ROS2 Galactic amcl ノード実行ビルド (colcon) 実行ファイル • amcl ソースコード amclノードの実行方法について • 後述のperfによる分析において分析対象を絞り込むため、amclノード単体で起動するビルドされたamclのバイナリを直接実行することで可能 (ros2 run や launch を使わない) $ install/nav2_amcl/lib/nav2_amcl/amcl --ros-args [options] ビルドされたamclノードのバイナリを実行することでamclノードのみ起動できる Copyright © Fixstars Corporation 50

51.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法 (2) perf、FlameGraphを使用して分析実行手順 1. Dockerコンテナにperfをインストール $ apt-get install linux-tools-generic 2. perfによる分析を有効にする sudo sysctl -w kernel.perf_event_paranoid=-1 3. perfの引数にamclの実行コマンドを指定して実行 perf record –g: perf によるデータ採取コマンド (-g: 関数の階層構造を読み取る) $ /usr/lib/linux-tools-4.15.0-161/perf record –g \ install/nav2_amcl/lib/nav2_amcl/amcl --ros-args [options] Copyright © Fixstars Corporation 51

52.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法 (2) perf、FlameGraphを使用して分析実行手順 4. FlameGraphのスクリプト[14]を使用し、 perfの分析結果を可視化 $ perf script | ./stackcollapse-perf.pl \ | ./flamegraph.pl > output.svg [14] https://github.com/brendangregg/FlameGraph Copyright © Fixstars Corporation FrameGraph 52

https://github.com/brendangregg/FlameGraph

53.

Fixstars Corporation www.fixstars.com 3.3.3 プロファイリング方法 (3) 分析結果から計算負荷の高い関数を特定する • FlameGraphから計算負荷が高い関数を特定する • FlameGraphにより、関数構造と処理負荷の比率が分かる • タイマの埋め込みにより処理時間を計測する • FlameGraphで判明した処理負荷の大きい関数周辺にタイマを埋め込み、具体的な処理時間を計測する関数の呼び出し順序処理負荷 Copyright © Fixstars Corporation 53

54.

Fixstars Corporation www.fixstars.com 3.3.4 プロファイリング結果 • FlameGraphから計算負荷が高い関数を特定する • amclの主処理の中で、sensorFunction関数の計算負荷が最も高い sensorFunction関数 : amclノードによる自己位置推定処理の中で、最も計算負荷が高い関数 laserReceived関数 : amclノードによる自己位置推定処理の本体となる関数ロボットの動作中、VLP-16からのセンサ情報を繰り返し受け取り、処理し続ける Copyright © Fixstars Corporation 54

55.

Fixstars Corporation www.fixstars.com 3.3.4 プロファイリング結果 • タイマの埋め込みにより処理時間を計測する laserReceived: 41.761 msec getOdomPose: 0.083 msec shouldUpdateFilter: 0.016 msec laserReceived関数処理時間内訳 pf_update_resample 12% odometryUpdate 6% getOdomPose odometryUpdate: 2.504 msec shouldUpdateFilter updateFilter (≒ sensorFunction): 33.179 msec odometryUpdate pf_update_resample: 4.875 msec updateFilter publishParticleCloud: 0.472 msec pf_update_resample updateFilter 79% getMaxWeightHyp: 0.526 msec publishAmclPose: 0.128 msec publishParticleCloud getMaxWeightHyp publishAmclPose calculateMaptoOdomTransform: 0.194 msec calculateMaptoOdomTransform sendMapToOdomTransform: 0.175 msec sendMapToOdomTransform Copyright © Fixstars Corporation 55

56.

Fixstars Corporation www.fixstars.com 3.4 amclのCUDA化による高速化アウトライン • 3.4.1 NaiveなCUDA化処理の実装 • 3.4.2 CUDA実行モデル • 3.4.2 CUDA化処理プロファイリング • 3.4.3 CUDA化処理最適化 • 3.4.4 Jetson上における性能測定の留意点 Copyright © Fixstars Corporation 56

57.

Fixstars Corporation www.fixstars.com 3.4 amclのCUDA化による高速化(sensorFunction CPU処理時間) • CPU処理での処理時間を計測 • amclでは適応的にパーティクル数を変化させる • rosbagファイルを流して一連の自己位置推定を行った時のパーティクル数と処理時間の関係を測定 • 結果 • 処理時間はパーティクル数に比例 • 最大パーティクル数で約40msec • LiDARの入力周期が10Hzなので最大レイテンシが約半周期 Copyright © Fixstars Corporation 57

58.

Fixstars Corporation www.fixstars.com 3.4.1 NaiveなCUDA化処理の実装 – 検討 • sensorFunctionの処理 • 各パーティクルに対して LiDAR観測の尤度を求めて重みを更新する • amclデフォルトではパーティクル数 (500-2000)、走査方向 (60) の計算が必要 • 計算はパーティクルで並列となる各パーティクル処理 : 走査線各走査方向に対する尤度の和を観測の尤度とする Copyright © Fixstars Corporation 58

59.

Fixstars Corporation www.fixstars.com 3.4.1 NaiveなCUDA化処理の実装 – 結果 • まずはCPU処理をそのままCUDA化 • CUDAの各スレッドで各パーティクルを処理 • CUDAで必要なメモリ領域は cudaMallocで確保 • CUDA処理前にデータを cudaMallocで確保した領域にコピー • 高速化の効果が見える • パーティクル数、走査方向数を増やして自己位置推定精度を確保することが可能 • レイテンシが低くなる • CPU負荷の削減 • 削減分を他処理にまわせる Copyright © Fixstars Corporation cpu : CPU処理時の処理時間 gpu-naive : NaiveなCUDA実装の処理時間 59

60.

Fixstars Corporation www.fixstars.com 3.4.2 CUDA実行モデル GPU 命令ユニット • CUDAの内部には複数のSMがある • SM : Streaming Multiprocessor • 独立して実行できるプロセッサー • Jetson AGX Xavierでは20個 • カーネルの実行スレッドはBlockで区切られる • Block 内部は32スレッド単位のWarp で区切られる • Block単位でSMに割り当て • 同じタスクが違うSMで処理される • Warpごとに処理される • Warpに所属する32スレッドが同時に処理を行うコア … レジスタキャッシュ SM Block を割り当て、Warpごとに処理タスク Block (～1024Threds) Warp (32Threds) … … … Copyright © Fixstars Corporation 60

61.

Fixstars Corporation www.fixstars.com 3.4.3 CUDAプログラムのプロファイリング • プロファイリング実行環境 ROS2 Galactic (Docker コンテナ) amcl ノードは Docker コンテナ内で実行 amcl ノード分析プロファイリングツール (Nsight Systems、nvvp) 分析データデスクトップPC（Ubuntu 18.04） SSH Jetson AGX Xavier Copyright © Fixstars Corporation 61

62.

Fixstars Corporation www.fixstars.com 3.4.3 CUDAプログラムのプロファイリング • NVIDIA® Nsight™ Systems • NVIDIA社製のプロファイリングツール • プロファイリング結果をGUIで可視化 • GPU、CPUの処理どちらも分析可能 • GPUはCUDA API/カーネル実行時間をプロファイリング可能 ROS2 Galactic amcl ノード Nsight Systems (可視化) 分析 Nsight Systems デスクトップPC (分析) Jetson AGX Xavier Copyright © Fixstars Corporation 62

63.

Fixstars Corporation www.fixstars.com 3.4.3 CUDAプログラムのプロファイリング • NVIDIA® Nsight™ Systems 実行手順 ① Jetson上でDockerコンテナを起動 Nsight Systems プロファイリングで使う一時ディレクトリをマウント $ docker run --gpus all --rm -it --privileged --shm-size=1024m \ -v /opt/nvidia:/opt/nvidia -v /tmp/nvidia:/tmp/nvidia cuda-galactic:root ② Dockerコンテナ上で環境変数を設定 export CUDA_INJECTION64_PATH="/opt/nvidia/nsight_systems/libToolsInjection64.so“ export QUADD_INJECTION_PROXY="OSRT, $QUADD_INJECTION_PROXY" ③ Dockerコンテナ上で実行 LD_PRELOAD=“/opt/nvidia/nsight_systems/libToolsInjectionProxy64.so” \ install/nav2_amcl/lib/nav2_amcl/amcl --ros-args [options] ④ デスクトップPC上で Nsight Systems を立ち上げ Jetson に SSH 接続し docker 内で実行したプロセスを指定してプロファイリング Copyright © Fixstars Corporation 63

64.

Fixstars Corporation www.fixstars.com 3.4.3 CUDAプログラムのプロファイリング • NVIDIA Visual Profiler (nvvp) • NVIDIA社製のプロファイリングツール • CUDA カーネル内部のプロファイリング結果をGUIで可視化 • nvprof • NVIDIA社製のプロファイリングツール • プロファイリングのメトリクス取得 ROS2 Galactic amcl ノード Linux 18.04 nvvp 分析 nvprof Jetson AGX Xavier デスクトップPC Copyright © Fixstars Corporation 64

65.

Fixstars Corporation www.fixstars.com 3.4.3 CUDAプログラムのプロファイリング • nvvp 実行手順 ① Jetson上でDockerコンテナ起動 $ docker run --gpus all --rm -it --privileged --shm-size=1024m cuda-galactic:root ② Dockerコンテナ上でnvprof実行 nvprof --analysis-metrics -o amcl_cuda.nvvp \ ../install/nav2_amcl/lib/nav2_amcl/amcl --ros-args [params] ③ 生成したプロファイル結果をscpコマンド等でデスクトップPCにコピー ④ デスクトップPCでnvvpを起動してプロファイリング結果を可視化 sudo /usr/local/cuda/bin/nvvp \ -vm /usr/lib/jvm/java-8-openjdk-amd64/bin/java amcl_cuda.nvvp Copyright © Fixstars Corporation 65

66.

67.

Fixstars Corporation www.fixstars.com 3.4.3 CUDA化処理プロファイリング nvvp – kernel analysis 計算資源をフルに使えていない自動変数の扱いに問題？ Occupancy(*) が低い (*) Occupancy: CUDA へのWarp充填度 GPU 処理はレイテンシが高いのでそれを隠すためにSMに多数のWarp を投入する必要がある Copyright © Fixstars Corporation 67

68.

Fixstars Corporation www.fixstars.com 3.4.3 CUDA化処理プロファイリング nvvp – latency analysis パーティクル数が1280以下だと SM を全部(20 個 )使い切らない nvvp – compute analysis Divergence に問題を抱えたコードが存在 Divergence(*) に問題を抱えた箇所がある (*) Divergence: 命令を実行したスレッドの割合。条件分岐があると一部のスレッドしか実行を行わないことがあり、実行を行わないスレッドが多いほどこの値が低くなる Copyright © Fixstars Corporation 68

69.

Fixstars Corporation www.fixstars.com 3.4.4 CUDA化処理最適化 – 方針 (1) • cudaMemcpyで時間がかかる • Integrated GPU なのでコピーは本来不要 • ⇒ Unified Memory を使用してコピーを削除する ※ Unified Memory: ホスト, GPU から同一のメモリアドレスでアクセスできるマネージドメモリ cudaMemcpy での明示的なメモリ転送を行わなくてよい Discrete GPU では DRAM, GDRAM 間で暗黙的にメモリ転送される Jetson では同一の SoC DRAM にアクセスするため転送が不要 CPU (Host) CPU (Host) GPU iGPU 同一メモリアドレスでアクセス同一メモリアドレスでアクセス暗黙的に転送 DRAM PCIe GDRAM SoC DRAM Discrete GPU での Unified Memory Jetson での Unified Memory Copyright © Fixstars Corporation 69

70.

Fixstars Corporation www.fixstars.com 3.4.4 CUDA化処理最適化 – 方針 (2) • Divergence • 全パーティクルの重み合計をatomicAddでReduce計算している箇所が原因 • • ⇒ ワープ内シャッフル命令によりReduce ⇒ ワープ間の Reduce は CPU で計算 samples[global_thread_num].weight *= p; // atomicAdd の実装は NVIDIA ドキュメント[5]参照 atomicAdd(total_weight, samples[global_thread_num].weight); for (int i = 16; i >= xdim; i >>= 1) { current_weight += __shfl_down_sync(0xffffffff, current_weight, i); } • Occupancy • Blockのスレッド数は32まで引き下げられる • ⇒ 64 から 32 に変更 Copyright © Fixstars Corporation 70

71.

72.

73.

Fixstars Corporation www.fixstars.com 3.4.4 CUDA化処理最適化 – プロファイリング nvvp – kernel analysis メモリの使用効率はよくないメモリ帯域はまだまだ空きがある演算ユニットはほぼ使い切っている Copyright © Fixstars Corporation Occupancy は低いが演算ユニットを使い切っているので影響は低い 73

74.

Fixstars Corporation www.fixstars.com 3.4.4 CUDA化処理最適化 – プロファイリング nvvp – latency analysis Occupancy は改善していないがこれ以上 Warp を詰め込んでも演算ユニットが詰まるだけ nvvp – compute analysis 強烈な Divergence はなくなった Copyright © Fixstars Corporation 74

75.

Fixstars Corporation www.fixstars.com 3.4.4 CUDA処理最適化 – さらなる検討 • 演算ユニットの負荷削減 • ⇒ 演算強度を軽減する (double -> float) • Occupancy改善 • ⇒ 内側のループ (各パーティクル内の処理) も並列処理する • メモリ使用効率、Divergence軽減 • ランダムアクセスが必要なアルゴリズムのため難度が高いただし他の処理と比較して十分高速なため、これ以上の最適化は効果が薄い Copyright © Fixstars Corporation 75

76.

Fixstars Corporation www.fixstars.com 3.4.5 Jetson上における性能測定の留意点 • Jetsonにはいくつかの性能モードがあるので測定時には性能モードをあわせる • 下記例ではモード3 (30W、8 CPU) $ sudo /usr/sbin/nvpmodel -m 3 • CPU/GPU動作周波数を固定する • デフォルトだと、処理負荷、CPU温度に応じて動的に動作周波数が変わり、処理時間がばらつく • jetson_clocksコマンドで動作周波数を固定する • 同時にファンも最大まで回しておく $ sudo /usr/bin/jetson_clocks --fan 参考 : NVIDIA 公式ドキュメント Copyright © Fixstars Corporation 76

https://docs.nvidia.com/jetson/l4t/index.html

77.

Fixstars Corporation www.fixstars.com 4 今後の課題 • メカナムローバーでの ROS2 navigationによる自律走行 • amcl高速化によるnavigationへの効果測定 • navigation2内の他パッケージのCUDA高速化 • amclのさらなるCUDA高速化 • オドメトリ更新 • リサンプリング (KD木構築のCUDA化が必要) Copyright © Fixstars Corporation 77

78.

79.

https://www.fixstars.com/ja/recruit/

80.

Jetson活用 ROS2自律走行実現に向けて ～自律移動ロボット用ー自己位置推定のCUDA高速化～（2021/11/04）