Jetson NanoでGPGPU (CUDA) プログラミング ~ライフゲームの開発~

この記事について

Conway's Game of Life (ライフゲーム) をCUDA実装して高速化します。さらに、Shared MemoryやStream等の最適化手法を試し、PCとJetson Nanoでの挙動の違いを確認します。

最近、Deep Learningが非常に流行っています。その影響もあり、GPUといえば機械学習とセットで語られることが多いです。しかし、もともとはGenral Purpose GPUという言葉の通り、Deep Learningに限らず一般の計算用途に使用できるものです。ひと昔前は仮想通貨のマイニングが流行りましたね。今回は初心に戻って、Jetson NanoでGPGPUしてみようと思います。

実は、Windows PC用にCUDA実装したライフゲームのプロジェクトは以前作ったことがあります。今回は、そのプロジェクトをベースに、CMakeを使いマルチプラットフォーム化して、Jetson Nanoでも動くようにしました。

I created Conway's Game of Life on Jetson Nano.
500 Generations Per Second for 1024 x 1024 thanks to CUDA. https://t.co/nrXj9zMgIi pic.twitter.com/HVITJS7Xa2

— iwatake (@iwatake2222) 2019年5月11日

ソースコードの場所

ソースコードとビルド、実行手順です
https://github.com/iwatake2222/LifeGameFast

昔つくったビデオ

このビデオを見れば、雰囲気がつかめると思います。

環境

• Windows PC
  • Core i7-6400 @3.4GHz x 4 cores (8 logical cores)
  • GeForce GTX 1070 @1.531GHz
    • Pascal
    • 1920 CUDA Cores
    • 256GByte/s (256 bit x 8Gbps)
    • VRAM: 8GByte
  • Windows 10 64-bit
  • CUDA 10.0
  • Visual Studio 2017 64-bit
• Jetson Nano
  • ARM Cortex A57 @1.43GHz x 4 cores
  • GPU @921MHz
    • Maxwell
    • 128 CUDA Cores
    • 25.6GByte/s (64 bit x 1.6GHz x dual rate?)
    • 共有VRAM: 4GByte
  • jetson-nano-sd-r32.1-2019-03-18 (Ubuntu 18.04.2 LTS)
  • CUDA 10.0
  • GCC 7.4.0

今回作るライフゲームについて

基本的なルールは通常のConway's Life Gameと同じです。ただ、単にセルが生きているか死んでいるかだけだと面白くないので、今の年齢も併せて分かるように色分けで表示しました。年齢は、今の世代数を100として正規化した数字です。
各年代のセルがどれくらいいるかといった統計情報を別ウィンドウに表示するようにしました。

また、画面の大きさやパラメータ、アルゴリズム選択を行うためのコントローラUIも表示しています。

下記はWindows上でのCapture画面ですが、左側がコントローラUI、右上がゲーム画面、下が統計情報画面です。統計情報には計算速度が分かるように、Generation Per Second (GPS)と、Frame Per Second (FPS) も表示しています。
このGPSが今回の指標となる数字です。大きいほど計算が速いことになります。

ソフトウェア構造

1年以上前に書いたコードなので、思い出しながら記事を書いています。。。
上の図が全体のクラス図です。

左下にLogicNormalXXXというクラスがあります。Strategyパターンのように、色々なアルゴリズムを切り替えられるようにしています。今回は、通常のCPU実装、OpenMP実装、CUDA実装、とその他アルゴリズム色々(本記事では触れない)を切り替えています。CUDA実装内での最適化手法の切り替えは、格好悪いけどifdefで切り替えています。
LogicNormalCudaクラス(とそこから呼び出しているCudaコード)が今回の主戦場になります。

一番上にあるWorldContextというクラスが全体を管理するクラスです。使用するアルゴリズムと、ライフゲームのゲーム画面(WorldView)と統計情報画面(AnalVew)を持ちます。

画面描画にはOpenGLのWrapperであるFreeGlutを使用しています。また、UIのためにAntTweakBarを外部ライブラリとして使用しています。

まずはざっくり比較

まずは、OpenMPやCUDAといった、大きな粒度で分類した高速化結果を比較します。

高速化手法

1. 境界チェックの削除

通常、ライフゲームではトーラス平面を考えます。つまり、ドラクエなどのRPGのように画面端まで行ったら反対側につながっているとします。そのため、画面端にあるセルを計算するときには特別な対応が必要になります。
しかし、これをif文で書いていると遅くなります。そのため、処理の最初にパディングを付けてしまいます(上図の青い部分)。この領域は、トーラス平面であれば反対側のセルの状態をコピー、非トーラス平面であれば0を設定します。

これによって、各セルの計算をするときには余計なことは考えずに周囲8セルの状態をみるだけでOKになります。
以後、基本的には常にこの手法を使用します。

2. OpenMP

通常のC/C++実装でプログラムを作成したら、1つのCPUしか使われません。
ライフゲームでは多くのforループ処理が出てきます。それらを複数のCPUで並列実装することで高速化します。
OpenMPを使用して、各for文に#pragma omp parallel for を付けるだけで簡単に並列化が可能となります。

3. CUDA

今回の本題です。
OpenMPではCPUを用いた並列化であるため、せいぜい4~8並列にしかなりません。
CUDAではGPUを用いるため、CUDAコアの数だけ並列化されます。GTX 1070なら1920、Jetson Nanoなら128並列になります。理論的には、ですが。CPUに比べて圧倒的な並列度であることが分かります。

4. CUDA + 最適化

詳細は後で出てきますが、CUDA実装の中でも様々な実装方法があります。それによってさらに高速化が可能となります。

結果

GTX 1070 (PC)

Generation Per Second	512x512	1024x1024	2048x2048	4096x4096	8192x8192
C++	500	90	24	5.5	1.3
OpenMP	1000	200	55	13	4.6
CUDA	1000	500	111	27	8
CUDA + Optimization	1000	1000	1000	166	43

Jetson Nano

Generation Per Second	512x512	1024x1024	2048x2048	4096x4096	8192x8192
C++	250	55	12	2.1	0.4
OpenMP	166	76	26	4.9	1
CUDA	250	100	32	8.2	1.9
CUDA + Optimization	1000	500	142	36	8.7

PC, Jetson Nanoどちらも、通常のC++実装 < OpenMP < CUDA < CUDA + 最適化 の順に高速化され、傾向はほぼ同じでした。

PC(GTX1070)とJetson Nanoの性能を比べると、3～5倍程度の差でした。スペック上での性能差に比べると、そこまで大きな差にはなりませんでした。(リソースを限界まで使用した実装になっていない、ということでもありますが。)

※ GPSが500や1000と一気に飛んでいるのは、処理時間が1msec, 2msecなど非常に小さいケースです。
※ JetsonでOpenMPで512x512のケースだけ何故か傾向が違いました。何回か試したのですが毎回同じ結果になりました。謎。

CUDA最適化実装

先ほどまではざっくり「CUDA + 最適化実装」と書いていましたが、実は色々なテクニックがあります。それだけで本が書けてしまうレベルですし、僕も簡単なことしか分かっていません。
ここでは、基本的なレベルの最適化を試してみます。また、一般に遅いといわれているZero Copy Memoryも試してみます。これは、PC(GTX 1070)とJetson Nanoでのアーキテクチャの違いによる挙動差を確認するためです。

データフローの基本

上の図が、CUDAプログラミングする際の最も基本的なデータの流れです。常にこの流れを意識しておく必要があります。

上部の灰色がCPU側(Host側)です。C言語上でallocやnewをしたら確保される領域になります。メモリはPCスペックなどに出てくる所謂「メモリ」です。
下部の緑色がGPU側(Device側)です。こちらのメモリは「一般的には」CPU側とは物理的に異なるものになります。いわゆる「VRAM」です。
ここで、「一般的には」と書きましたが、PCではPCIeバスを介して物理的に離れています。そのため遅いです。一方、、JetsonNanoのような組み込み型では物理的に近い場所だったり、共有しているケースもあります。

データの流れを左上から見ていきます。
まずCPUが元となるデータをホストのDDR上に書き込みます。このデータをGPUで処理するためにはGPU側のデバイスメモリに転送する必要があります。しかし、通常のnew/alloc等で確保した領域では論理的には連続していても物理的には連続でない可能性があります。DMA転送などを使う都合だと思うのですが、物理的に連続である必要があります。また、ページアウトされてそもそも物理メモリ上に存在しなくなる可能性もあります。そのため、まずはNon-pageable memory(Pinned memory)として確保した領域にコピーされます。これはプログラマが意識せずとも勝手にやってくれます(やられてしまいます。)

その後、ホストメモリからデバイスメモリにデータを転送します(cudaMemcpyHostToDevice )。GPUはこのデータを用いて何らかの計算をして結果をデバイスメモリ上に書き込みます。

計算が完了したら、デバイスメモリからホストメモリにデータを転送して完了です(cudaMemcpyDeviceToHost )。

高速化手法

高速化として、大きく2つの軸に分けました。
1つは制御変更を伴うもの(以下のALGORITHM_X)、もう1つはメモリ確保方法に関係するものです。結果はこの2つの組み合わせで見てみようと思います。
※Shared Memoryも、もろメモリに関することなのですが、実装変更が大きいので別枠にしました。

ALGORITHM_0: 境界チェック有 + パディング無し

先ほど、「境界チェックがあるから遅くなるのでパディングを付けてif文をなくします」みたいなことを書きましたが、僕の中で、実はそっちの方が遅いのではないか疑惑があり、このケースも試してみました。

ALGORITHM_1: 境界チェック無し + パディング有り

周囲にパディングを付けて、境界チェックロジックを削除したアルゴリズムです。
このアルゴリズムを基準として考えます。先ほどの表・グラフの「CUDA」実装はこれです

ALGORITHM_2: Shared Memory

先ほどのデータフローの説明の図では、GPU(デバイス)側のメモリとして、Global Memoryを記載しました。これは、いわゆるVRAMで8GByteなど大きな容量を持っています。これはオンボード上のDDRであるため、遅いです。GPU内にも当然キャッシュ機構がありますが、さらに低遅延のオンチップメモリであるShared Memoryというものがあります。
Shared Memoryは、同一ブロック内のスレッドで共有される小容量だけど高速なオンチップメモリです。今回、ブロックサイズを32x32としました。そのため、32x32セル単位で共有するメモリを使用します。

Shared Memoryを使用しないと、1つのセルの計算に対して、Global MemoryへのReadが9回、結果のWriteが1回発生します。
一方、Shared Memoryを使用すると、Global MemoryへのReadは1回だけになります。この1回でまずGlobal Memory -> Shared Memoryへコピーして、計算時にはコピーしたShared MemoryのデータをReadします。これによって、Global Memoryへのアクセス回数を減らして、高速化を狙います。

Shared Memoryを使うには、カーネル関数内で__shared__ キーワードを付けた領域(配列)を宣言します。

ALGORITHM_2_STREAM: Shared Memory + Stream

先ほどのデータフローの所で説明した通り、基本的には以下の流れで処理が進みます。

1. Host -> Device(GPU) へのデータコピー
2. GPU側で計算
3. Device(GPU) -> Host へのデータコピー

何も考えずに実装すると、上記のフローはシーケンシャルに実行されます。
Streamでは、これらの操作を非同期に実行し、パイプライン化して処理を並列化させます。

パイプライン化する粒度として、フレーム単位でやるのが一般的だと思います。例えば、n-1 フレーム目のcalculation中に、n フレーム目のデータ転送をする、など。
今回は簡単のために、1枚のフレームを、横長の長方形に8等分して、各エリアのデータ転送、計算、データ転送をStream化しました。

Streamを使うには、cudaStreamCreate で生成したstreamに対して、cudaMemcpyAsync を設定することで非同期に処理を進めてくれます。(同じstream内での操作順序は維持してくれます)

ALGORITHM_3_STREAM: Shared Memory + Stream + host2deviceコピー抑制

この方法はライフゲームだから適用できる手段です。

繰り返しになりますが、基本的には以下の流れで処理が進みます。

1. Host -> Device(GPU) へのデータコピー
2. GPU側で計算
3. Device(GPU) -> Host へのデータコピー

ここで、1番の「Host -> Device(GPU) へのデータコピー」は不要なのではないかと考えました。
n フレーム目の入力画像は、n-1 フレーム目の出力であり、そのデータは既にデバイス側メモリにあるのだから、それをそのまま使います。その結果、1番の「Host -> Device(GPU) へのデータコピー」は不要になります。

1番の「Host -> Device(GPU) へのデータコピー」は、一番最初と、ユーザが操作してセルの配置状態を変更した時だけ行います。

ALGORITHM_3_REPEAT: Shared Memory + Stream + host2deviceコピー抑制 + device2hostコピースキップ

先ほどの続きで、そもそも結果の出力(3番の「Device(GPU) -> Host へのデータコピー」)も毎回やらないでいいんじゃない? という考えです。
例えば、ライフゲームのロジック計算で100GPS出たとしても、描画処理やディスプレイ出力部分で結局60fps程度に抑えられます。であれば、結果出力用のデータ転送は少しくらい間引いても問題ないはずです。
こうすると、Host<->Device間のデータ転送がほとんどなくなります。

今回は、10フレーム間隔で間引くことにしました。

※
今回はライフゲームなので結果をそこそこの頻度で取得して表示する必要がありました。
Tensorflow等を用いたDeep Learningの学習は、一度データをDevice側に転送したら学習中はずっとDevice側だけで行い、本当にすべての処理が完了したタイミングでHost側に結果を返すことで、高速実行されているのだと思います。

メモリ確保: Pinned Memory

メモリの確保方法でも最適化が出来ます。
左上の図は最初に説明したものと同じです。Host側のCアプリケーションでnew/allocすると、その領域はDDR上のPageable memoryとして確保されます。しかし、これだとDeviceへの転送に使えないため、暗黙的にNon-pageable memoryにコピーされます。

これが無駄です。

コピーされるくらいなら、最初からNon-pageable memoryとして確保すればいいじゃない、という考えです。これを、Pinned Memoryというようです。

Pinned Memoryとして領域を確保するには、cudaMallocHost を使用します。

デメリットとしては、一般に論理アドレス導入によってもたらされる恩恵が無くなる点があります。Pinned Memoryとして領域確保すると、実アドレス空間がそれだけ占有されます。とはいえ、これは今の時代そんなに大きなデメリットではないかと思います。

メモリ確保: Zero Copy Memory

これまで教科書通りに、Host->Deviceへデータ転送して、GPU側はDeviceメモリに対してアクセスをするようにしていました。しかし、実はGPUはHostメモリに対してもアクセス可能です。
こうすることで、Host<->Deviceのデータ転送が不要になります(Zero Copy)。
一方、GPUがHostメモリにアクセスするコストが大きい場合、計算途中のレイテンシも非常に大きくなってしまい、結果としてパフォーマンスが低下します。

一般に、PCの場合は、GPUとHostメモリはPCIeでつながっているため遅いです。ハイスペックなPCIe3.0 x 16でも16GByte/sであり、GTX 1070のVRAM帯域256GByte/sに比べるとはるかに遅いです。
一方で、Jetson Nanoの場合は、共有メモリなので、実はデメリットはないのではないかな? と思っています。これについてはこれから結果を見ていきたいと思います。

Zero Copy Memoryを使うには、ホストメモリをcudaHostAllocMapped フラグを付けて取得して、cudaHostGetDevicePointer でデバイス側からアクセスできるポインタを取得します。

結果

GTX 1070 (PC)

画面サイズは8192x8192

	Generation Per Second	Normal	Pinned Memory	Zero Copy Memory
ALGORITHM_0	Normal without padding	8.5	20	2.8
ALGORITHM_1	Normal	8.4	20.8	2.3
ALGORITHM_2	Shared Memory	8.5	21.7	7.3
ALGORITHM_2_STREAM	Shared Memory + Stream	8.5	29.4	7.3
ALGORITHM_3_STREAM	Shared Memory+ Stream+ No Copy(h2d) unless update	16.3	43.5	7.3
ALGORITHM_3_REPEAT	Shared Memory+ No Copy(h2d) unless update+ No Copy(d2h) while skipping (10 times)	100	150	9.4

まずは列方向に結果を見てみます。Pinnec Memoryを使うことで、全体的にパフォーマンスが向上しています。一方、Zero Copy Memoryを使うと、非常に遅くなっています。使用するアルゴリズムに依らず遅くなっていることから、GPU-Host間のMemory転送がボトルネックになっているのだと思われます。

次に、アルゴリズムの違いによる結果を見てみます。
まず、境界チェックではあまり違いはありませんでした。むしろ少し遅くなってる???。パディング領域を作るためのデータ転送が遅いのが原因かもしれません。
Shared Memoryも、高速化にはあまり貢献していないように見えます。ただ、Zero Copy Memoryのケースでは明らかに改善しているので、実装を間違えているとかではないようです。GTX1070ではGlobal Memoryがそんなに遅くないのかもしれません。そのため、Shared Memoryにしてもあまり変化がなかったのではないかと考えています。これなら、Zero Copy Memory(ホストメモリ)のケースで速くなっていることも納得です。イメージ的には、Shared Memory >= Global Memory >>> Host Memory といった感じでしょうか? あくまで今回の測定結果からですが。
Stream化することで、さらに高速化することが出来ています。効果があるのはPinned Memoryを使ったケースだけなのですが、おそらく他のケースではメモリ転送時間が支配的で効果が実感できないのだと思われます。
ALGORITHM_3_STREAM(ホストメモリからデバイスメモリへの転送をしない)では、明らかに高速化しています。Normal(Pinned Memoryを使わない)のケースで約2倍になっていることから、メモリ転送時間がほぼ全体の処理時間を占めていることが分かります。(h2dで50%、d2hで50%)。Pinned Memoryを使用するケースでもかなりの改善が見られました。一方、Zero Copy Memoryのケースでは変化有りませんでした。これは、Zero Copy Memoryではそもそも転送がされないためです。
最後のALGORITHM_3_REPEATは、デバイスメモリからホストメモリへの転送間隔が1/10に間引かれたのでさらに高速化しています。

Jetson Nano

画面サイズは4096x4096

	Generation Per Second	Normal	Pinned Memory	Zero Copy Memory
ALGORITHM_0	Normal without padding	7.6	13	8.4
ALGORITHM_1	Normal	8.2	17	4.0
ALGORITHM_2	Shared Memory	9.6	27	33
ALGORITHM_2_STREAM	Shared Memory + Stream	6.4	22	31
ALGORITHM_3_STREAM	Shared Memory+ Stream+ No Copy(h2d) unless update	12	36	32
ALGORITHM_3_REPEAT	Shared Memory+ No Copy(h2d) unless update+ No Copy(d2h) while skipping (10 times)	38	45	33

Jetson Nanoの結果はPCと異なるところが多々あります。

ALGORITHM_0/1 境界チェック処理の有無ですが、結局は以下2つの処理のどちらが遅いか、ということだと思います。

• 全セルの処理時に境界チェック(if文)をする
• 処理の最初にパディング領域を作る

これを考えると、メモリアクセスが遅いZero Copy MemoryのケースでALGORITHM_1(パディング領域作成)が大幅に遅くなっているのは納得がいきます。
NormalとPinned Memoryのケースでは、PCでは差はなかったのですが、Jetson Nanoでは少し高速化しています。これは、Jetson NanoのGPU計算能力がGTXに比べて低いため、if文ロジックを大量に行うよりも、一度パディング領域を作る方が速度的に有利になったのだと思われます。

次にShared Memoryです。Normal、Pinned Memoryどちらも高速化しています。また、改善幅はPCよりも大きいです。これは、Jetson NanoではShared MemoryはGlobal Memoryに比べてだいぶ速いといことなんだと思います。というよりも、Global Memoryが遅いという方が正しいかもしれません。これは、共有VRAMだからではないかと思います。この共有VRAMは、次のZero Copy Memoryの結果にも影響します。Zero Copy Memoryでは33GPSにまで高速化され、この3つの中では一番速いです。
この結果の受け止め方ですが、Jetson NanoではVRAMが共有ということを思い出す必要があります。そのため、GPUのGlobal Memory = ホストメモリとなります。Pinned Memoryを使い、いつも通りcudaMemcpyでh2d,d2hとコピーしても物理的に同じメモリ上でコピーしているだけなので、メリットはありません。むしろ無駄です。Zero Copy Memory使用時にはこのコピー処理が不要になるため、最も速くなったのだと思われます。

次のStreamの結果はいずれも遅くなってしまっています。なぜだか理由がさっぱり分かりません。
想像ですが、バスが非常に混んでしまうせいではないかと疑っています。ホストとデバイスでメモリが分かれているPCでは転送とGPU計算がStreamによってパイプライン化され並列実行しても、リソースの競合は起きずらいと思います。しかし、共有VRAMであるJetson Nanoでは、メモリ転送でも計算時でも、物理的に同じメモリアクセスが起きます。そこが混雑しているのではないかな～と思っています。

ALGORITHM_3_STREAM, REPEATを用いることで、Normal、Pinned Memoryどちらもさらに高速化されます。これは、データ転送回数が減るので当然の結果ですね。一方、Zero Copy Memoryのケースでは変化有りませんでした。これは、Zero Copy Memoryではそもそも転送がされないためです。
ただ、ここで気になったのが、NormalとPinned Memoryの速度がZero Copyの速度を上回ることです。結局データ転送時間が支配的で、メモリアクセス時間は変わらないのだから、データ転送を行わないZero Copy Memoryの31GPS付近が上限になると思っていました。しかし、実際にはZeroCopyMemoryの速度を超えています。Zero Copy Memoryだと何らかのオーバーヘッドがついてくるのかもしれません。
(キャッシュのコヒーレンシ保つために毎回キャッシュクリアしてるとか? でもそれならPinned-memoryも同じか。。。というか、CPU側ってどうやってキャッシュコヒーレンシ保ってるんだろう。dirty bitが立つのかな。。。ということは、Zero CopyだとGPUが処理するたびにDirty bitが立つからキャッシュヒット率が落ちるということ? 一方、Pinned memoryだとmemcpyしたときだけDirby bitが立つからそこまでヒット率は落ちないとか。 であればこの結果も納得。)

PCとJetson Nanoを比較すると、色々と異なる部分はありましたが、どちらもALGORITHM_3_STREAM + Pinned Memory が最速でした。先ほどの表・グラフ内の「CUDA + Optimization」はこの結果です。
(ALGORITHM_3_REPEATは、反則的なので除く)

おわりに

残った謎

• Jetson Nanoだとプログラム全体の動きが遅い。例えば8192x8192だとライフゲームの計算が36GPS、描画が20fps程度だったので1 / (1/36+1/20) = 12.8fps程度は全体で出るはず。なのに、実際の動作だと明らかに遅くカクついているように見える。どこかでストールしているのか、freeglutのメインループ処理が回っていないっぽい
• Pinned MemoryやZero Copy Memoryで、GPUが書き込んだ場合、キャッシュコヒーレンシはどうやって保たれるのだろうか。
• PC（GTX1070) でShared Memoryを使用した実装で、あまり高速化されなかった。バンク衝突とかが起きているのかな? それとも、単にGlobal Memoryが十分速いからShared Memoryの効果を感じられないだけ?
• Jetson NanoでStreamを使用した実装で、速度が低下した

まとめ

• マルチプラットフォームで動くライフゲームプロジェクトを作りました
• Open MPやCUDAを使って高速化しました
• 動作をPC, Jetson Nanoで比較しました
• CUDAの実装最適化の結果は、PCとJetsonNanoで異なることが分かりました。また、別のシステム/アーキテクチャでは、別の結果になりそうです
• 最適化実装には一つの正解があるものではないので、システム構成や状況を見て、都度検討が必要そうです