CUDAを一枚の絵にまとめてみた

CUDA
このエントリーをはてなブックマークに追加 follow us in feedly

まえがき

初めまして、R&Dチームの宮﨑です。趣味はFPSをやってます。150時間やってますが未だにドンかつを食べられていません。普段は深層学習のベンチマークを取ったりしています。

実はひと昔前からあった深層学習。近年実用性が増して一大ブームとなっています。実用的になった背景の一つとしてGPUを用いた高速化により深層学習の学習及び推論が実用的な時間で実行できるようになったことがあります。NVIDIAから提供されるCUDAを用いるとC言語に拡張を加えた形式でCPU+GPUのヘテロジニアスコンピューティングを記述できます。 CUDAを使用するにはCPU+GPUの環境がどのようなハードウェア構成をしており、その上でCUDAがどのようなシステムを構築しているのかを理解する必要があります。この理解なしに漫然とサンプルコードを真似するだけでは意図通りのパフォーマンスが出なかったり、そもそもAPIの意味が理解できずに書けない・・・なんてことになります。そこでCUDAを中心とした環境を一枚の絵にまとめてみました！

本記事では以下のような方々を対象としております。

  • CUDAにちょっと興味をもって概要をぱっと見てみたい方 ⇒ 各章の概要を読んでいくことをお勧めします♪
  • CUDAを勉強しているがいまいち全体像がつかめず、どう書けばよいかよくわからない方 ⇒ 全体像の一枚絵を見てイメージを掴んでください!(^^)!
  • CUDAを使って高速化したはずなのに思うような速度が出ずに困っている方 ⇒ 詳細まで読めばヒントを掴めるかもしれません

※参考にした「CUDA C プロフェッショナル プログラミング」はCUDA Toolkit 6.0とFermi・Kepler世代のGPUを基本として解説されていますので、本記事も同様です。本記事で得た理解を起点として差分を見ることでお使いになるバージョン・GPUでの動作を理解できると思います。CUDA Toolkitの各バージョンの情報は公式が出しているリリースノート、GPUはCC(コンピュートケイパビリティ)が参考になると思います。

全体像

青四角がハードウェア、黄色四角がCUDAで用いられる概念を表しています。 青矢印はCUDAがハードウェアのメモリにどのソフトウェアのメモリを割り当てるか、黄矢印はCUDAで扱う概念の流れを表しています。 赤線はデータ通信を表しています。 各所の説明を述べていきます。 f:id:optim-tech:20190814205720p:plain

基礎編

①カーネル/メモリ転送

概要

GPUに計算させるにはホストメモリ(CPUのメモリ)からデバイスメモリ(GPUのメモリ)へ計算に必要なデータを転送(Host to Device, H to D)、GPUで計算(カーネル、メソッドの形式で記述)、計算結果をGPUからCPUに転送(Device to Host, D to H)が基本的な流れになります。またCPUのプロセスからカーネルをコールしGPUで計算を行いますが、CPUのプロセスとカーネルは基本的に非同期です。すなわちCPUはカーネルをコールした後カーネルの終了を待たずに次の処理へ移ります。

詳細

CPU-GPU間は物理的にはPCIe バスで結ばれておりデータ転送で使います。 メモリ転送(H to D, H to H, D to H, D to D)をするにはcudaMemcpy、デバイスメモリの領域を確保/解放をするにはcudaMalloc/cudaFree、デバイスメモリに値をセットするにはcudaMemsetを使います。カーネル(GPUの計算処理)は処理を記述したメソッドに__global__などの修飾子をつけ、通常のメソッドと区別します。CPUーGPU間で同期を取りたい場合はcudaDeviceSynchronizeというバリアを用いて足並みを揃えます。Kepler世代以降ではダイナミックパラレリズムという機能がついており、カーネルからさらにカーネルを呼び出せるようになりました。

②ストリーム ③ストリームスケジューリング/ハードウェアキュー

概要

カーネル・メモリ転送はストリームというキューに順に追加し、複数のストリームを並行して走らせることができます。 作成されたストリームはスケジューリング処理に渡され、CPU-GPU間のタスクの受け渡しをするハードウェアキューに押し込まれます。 またストリームにおける処理の各地点にマーカーを設けることができ、イベントと言います。イベントを使うことでストリーム同士で同期を図ったり、進行状況の監視ができます。

詳細

cudaStreamCreateでストリームを作成し、非同期で実行する用のメソッド(例えばcudaMemcpyAsync)に渡すことでその操作をストリームのキューに追加します。ストリーム同士で足並みをそろえるためにはバリアを設けるcudaStreamSynchronizeやストリームの進行状況を問い合わせるcudaStreamQueryを使います。使い終わったストリームを破棄するためにはcudaStreamDestroyを使います。他にもcudaStreamCreateWithPriorityを使って各ストリームに優先度を設定したり、cudaStreamAddCallbackを使ってコールバックを設定できます。PCIeバスの仕様上、異なるストリームに属していても方向が異なるデータ転送(H to D & D to H)はオーバーラップできますが、同方向のデータ転送はできないことに注意してください。CUDAではいかに処理をオーバーラップさせるかが並列性の高さへのカギです。ハードウェアキューを複数にしてより高いパフォーマンスを目指したものがHyper-Qと呼ばれる技術です。

イベントはcudaEvent_t eventと宣言しcudaEventCreateメソッドを使って作成します。このイベントをcudaEventRecordを使ってストリームに登録し、cudaSynchronizeでスレッドのバリア作成・cudaEventQueryでイベントが完了したかの確認・cudaEventElapsedTimeで時間計測などを行い、不要になったイベント削除する際はcudaEventDestroyを使います。

注意点としてCUDAには暗黙的な同期があります。デバイスメモリの確保・設定、メモリコピー、ピンメモリの確保、L1キャッシュとシェアードメモリの設定の変更などです。CUDAは通常CPU-GPUが非同期で動作してくれますが、暗黙的な同期によってそれがブロックされ、パフォーマンスの低下を招くことがあります。

④スレッド ⑤ブロック ⑥グリッド

概要

GPUにおいて計算処理を行うカーネルは処理を行う１つ以上のスレッド、スレッドを束にした一つ以上のブロック、ブロックを束にしたグリッドの三層構造となっています。 各スレッドは異なるデータに対し同一の命令を実行するため、SIMDのような形式です。各スレッドにおいて処理を行う際はSM上のユニットを使います。

詳細

起動するスレッド・ブロックの数はGPUが許す範囲で自由に指定できます。各スレッド/ブロックはblockIdx・threadIdxを用いて割り振られたインデックスを取得し、それらを用いて処理の差別化をします。スレッドを同期させる機能としては各ブロック内のスレッド同士を同期させるバリアを設ける__syncthreadsや書き込まれたデータへの参照を保証するメモリフェンスを設ける__threadfenceがあります。

SM上のユニットは以下のものがあります。

  • 整数または浮動小数点数の命令を実行するCUDAコア
  • 倍精度の命令を実行するDP
  • 正弦、余弦、平方根といった組み込み命令を実行するSFU
  • ロード/ストアを実行するロード/ストアユニット

渡されたグリッドに含まれるブロックは実行する際GigaThreadエンジンによって各SMのワープスケジューラに分配されます。さらにブロックはワープという32個の連続したスレッドによって構成されるものに分解されます。ワープ内のスレッドはすべて同じ命令を実行するため、条件分岐でワープ内のスレッド間で異なる命令を実行させようとすると、全てのスレッドが同じ命令を実行した後に該当しないスレッドが無効になり非効率的です。このことをワープダイバージェンスといいます。ワープごとに同一の命令を実行するように条件分岐を切り分けるなどして回避することが望ましいです。並列に動作する各ワープに必要な情報は常にチップに乗っているため、実行するワープは高速に切り替えることができます。従って多くのワープを起動できるようにすることがリソースの高い使用率に繋がります。またワープ内のスレッド同士はレジスタを直接読み取ることができ、この機能をワープシャッフル命令といいます。

ワープは32個のスレッドを含むことからブロックあたりのスレッド数を32の倍数にすると任意のワープ内で遊ぶスレッドができず効率的です。また起動するブロック数やスレッド数が少なすぎるとリソースがフルに活用されない、多すぎると必要なリソースをスムーズに確保できない可能性があることも考慮すべきです。CUDAではリソースをフルに稼働させること、実行モデルに即した実装をすることが重要になってきます。リソースの使用状況及び実行モデルに即した実装をする上で必要となる情報・指標を取得するには後述するプロファイリングツール等を使います。

メモリ編

CUDAにはCPUのメモリモデルと同様に用途に応じた複数種類のメモリが用意されます。各メモリの利用率及び動作の並列性を高めることでパフォーマンスを上げることができます。 CPU-GPU間はPCIe バスを通じて通信しますがCPU内・GPU内の通信と比較しかなり低速であるため、速度を確保するにはCPU-GPU間の通信を最小限に抑えたり、カーネルとオーバーラップさせ隠蔽する構成を考える必要があります。CPU側とGPU側双方とも基本的にもう片方のメモリへのアクセスはできないことに注意して下さい。どちらのメモリも同ファイルに記述するため混同しやすいです。変数名の命名規則などで区別しやすいようにしておくことをお勧めします。

⑦レジスタ ⑧シェアードメモリ

概要

スレッドごとに割り当てられるレジスタと同ブロック内のスレッド間で共有されるシェアードメモリがあります。

詳細

各SMのスレッド間で分割されるレジスタはレジスタファイルに格納され、GPUにおいて最も高速なメモリです。修飾子をつけない変数は基本レジスタに格納され、スレッドのライフタイムに渡って維持されます。使用するレジスタの数がハードウェアの上限を超えた場合、超えた分はローカルメモリに退避されます。この現象をレジスタスピルと言い性能低下を招く恐れがあります。

オンチップはシェアードメモリとL1キャッシュに分割され、その割合はcudaFuncSetCacheConfigメソッドを使用して指定できます。シェアードメモリはブロックのライフタイムに渡って維持され、オンチップなメモリなのでオンボードなメモリであるグローバルメモリに比べて高速に動作してくれます。修飾子__shared__をつけることで1・2・3次元の配列として宣言でき、extern __shared__を修飾としてつけることでサイズが動的な一次配列として宣言できます。また必要となるレジスタとシェアードメモリの領域の大きさによってSMに割り当てられるブロックとワープの数が決まります。CUDAではこのようなある要素の要求する内容によって他の要素の割り当てが決定される関係が多くあるため、それらの関係を把握し最適化の際に調整できるようにする必要があります。またシェアードメモリはプログラマブルなことからフル活用すべきリソースであり、具体的な用途としてはグローバルメモリへのキャッシュ、中間に挟むことで高速なアクセスパターンを構成するなどがあります。シェアードメモリは32個の同じサイズのメモリモジュール(バンク)で構成され、複数のスレッドが同バンクの異なるアドレスにアクセスした場合その数だけリプレイされ性能が低下する恐れがあります。この現象をバンクコンフリクトと言います。同一のアドレスに複数のスレッドがアクセスした場合は一度だけアクセスが実行され、スレッドへブロードキャストします。この場合もアクセスする領域は少ないため帯域幅は低いです。シェアードメモリのどのアドレスがどのバンクに属するかは算出できることから、きめ細かい配置をプログラマでデザインできるため、バンクコンフリクトなどのパフォーマンスを下げる現象を回避できるかもしれません。具体的なテクニックとしてはアルゴリズムを変更したり、パディングを入れることでメモリ配置を変更することが考えられます。

⑨L1 ⑩L2 ⑪ローカル

概要

ローカルメモリは先述したようにレジスタの退避用などに使用されます。 ローカルメモリとグローバルメモリのキャッシュとしてSMごとのL1キャッシュとデバイスごとのL2キャッシュがあります。

詳細

ローカルメモリが配置される場所はグローバルメモリと同じです。データはL1,L2キャッシュにも格納されることがあり、ライフタイムはスレッドと同じです。

L1キャッシュは有効無効を切り替えることができます。また__ldgメソッドを使用するか__restrict__修飾子を使用することでL1キャッシュの代わりにリードオンリーキャッシュを使うことができます。ばらつきのあるアクセスではL1よりもリードオンリーキャッシュの方が高い性能を発揮することが多いという性質があります。この性質は空間的局所性のあるアクセスで高い性能を発揮するコンスタントキャッシュとは対称的です。またCPUと違ってGPUのL1キャッシュは空間的局所性を考慮していますが時間的局所性は考慮されていないことに注意が必要です。

⑫グローバルメモリ ⑬コンスタントメモリ ⑭テクスチャメモリ

概要

全てのスレッドから参照可能であり、アプリケーションのライフタイムに渡って維持されるメモリとして、グローバルメモリ・コンスタントメモリ・テクスチャメモリがあります。 グローバルメモリはGPU/CPUの両方から、コンスタントメモリ・テクスチャメモリはCPUからのみ書き込みができます。コンスタントメモリは同一のアドレスにアクセスするスレッドが多い場合、テクスチャメモリは二次元の空間的局所性が高い場合などで高い性能を発揮します。

詳細

グローバルメモリはGPUで最大の容量と遅延を持つオンボードなメモリです。動的に確保するには上述のcudaMallocを使い、静的に確保するには__device__修飾子をつけます。グローバルメモリへのワープのアクセスに関するパフォーマンスはアクセスするアドレスの分散具合・アラインメントに大きく影響されます。基本的にコアレスであり、アラインされたメモリアクセスが最も高いパフォーマンスを発揮します。シェアードメモリやグローバルメモリの宣言時にvolatile修飾子をつけるとキャッシュする最適化が行われなくなり、コヒーレンシが考えられなくなります。

コンスタントメモリはSMごとのコンスタントキャッシュにキャッシュされ、グローバルメモリと同じ場所に格納されます。__constant__の修飾子をつけることで宣言できます。GPU側では読み取り専用であるため、CPU側でcudaMemcpyToSymbolメソッドを用いて初期化させておく必要があります。コンスタントメモリという独自のキャッシュを使うことでL1キャッシュを使うグローバルメモリへのアクセスと並行して実行できるため、上手く使えば上述した並列性の高いメモリ動作を実現できます。

テクスチャメモリはSMごとのテクスチャキャッシュにキャッシュされ、グローバルメモリと同じ場所に格納されます。テクスチャキャッシュもコンスタントキャッシュと同様にグローバルメモリへのアクセスとは別ルートなので活用できれば並列性の高いメモリ動作を実現できるかもしれません。

⑮ユニファイドメモリ ⑯ピンメモリ/ゼロコピーメモリ

概要

ユニファイドメモリはCPU・GPUの両方から同じアドレスでアクセスできます。データ移行が必要になった際にはCUDA側で自動的に対処してくれるユニファイドメモリをマネージドメモリと言います。 CPUとGPUではメモリは基本的に切り分けられますが、GPUからCPUのデータにアクセスできる仕組みとしてピンメモリ/ゼロコピーメモリがあります。

詳細

マネージドメモリは動的にはcudaMallocManagedメソッドを使うことで、静的には__device__ __managed__を修飾子として使うことで宣言できます。ユニファイドメモリを用いることでプログラムの記述をシンプルにすることができます。また似た技術としてCPU側とGPU側で仮想アドレス空間を共有するUVAという技術があります。UVAではデータが物理的に移動することがありません。

CPU側のデータは基本的にページング可能ですが、GPU側から確実にアクセスするため一時的にピンメモリというページングできない領域をCPU側に設け、そこにデータを移すことでGPUからアクセスできるようにします。ピンメモリを確保する際にはcudaMallocHostメソッドを用います。当然多用するとCPUでページングされるデータが減少するためパフォーマンスが低下する可能性があります。またピンメモリを確保した際はcudaFreeHostメソッドを用いて解放する必要があります。デバイスのアドレス空間にマッピングされたピンメモリをゼロコピーメモリといいます。

調査編

⑰各種情報を取得する機能

nvidis-smiというコマンドを用いればOSがGPUを認識しているのか、メモリの使用率は今どれぐらいかなどGPUの動作状況を取得することができます。また、表示される情報の中にGPUに割り振られたIDがありますが、CUDA_VISIBLE_DEVICESという環境変数にそのIDをセットすれば使用するGPUを制限することができます。

Nsight/nvvp/nvprofといったプロファイリングツールを用いることでCUDA上で実行された処理のパフォーマンスを計測することができます。各命令の所要時間、メモリの効率性、コンフリクトの発生状況など最適化を行う際にほしくなる情報を豊富に提供してくれます。実行時の詳細な情報を提供するツールがあり、メモリへのデータ配置・移動の仕方などハードウェアレベルで細かく動作をカスタマイズできるようにすることでCUDAは最適化を追求できるようになっています。プロファイリングツールを用いてボトルネックを発見し、理論値との差やハードウェアのリソース量・特性を考慮した実装に修正し問題個所を最適化することが基本戦略となります。

cudaDevicePropを用いるとGPUのプロパティを取得できます。起動できるスレッド数・ブロック数の上限やマルチプロセッサ数などCUDAを用いたプログラミングの際に必要となる情報を提供してくれます。

nvccはCUDAのコンパイラですが、オプションを指定することでレジスタ・シェアードメモリなど各種使用されるリソースを実行時に出力させることができます。またレジスタの最大数などリソースの使用に関する指定もできます。

cudaGetErrorStringを使ってCUDAのメソッドから発生するエラーコードを人の読める文字列に変換してくれます。このメソッドを使って例外処理を記述できます。

cuda-gdbを使うとgdb的な機能に加えて、CUDA環境の表示や実行されるスレッドのうち一つだけに着目してprintしてくれるなど嬉しい機能があります。またCUDA4.0以降ではデバイスでのprintfがサポートされており、おなじみのprintfデバッグができます。ただし並列に実行される処理にprintfを入れると大量に表示されるので注意が必要です(汗) CUDA特有のメモリ関連を調べるツールとしてcuda-memcheckがあります。

最後に

今回はCUDAを一枚の絵にまとめ、各所の簡単な解説を行いました。ヘテロジニアスコンピューティングの面白さを少しでも感じていただけたのなら幸いです。ヘテロジニアスコンピューティングは今後ますます拡大して面白くなっていくでしょう。 オプティムは出身や専門分野を問いません。そんなヘテロジニアスな会社で働いてみませんか。応募待ってます！