SC13 - CPU-GPU統一メモリを実現したCUDA 6(前編)

今回のSC13にあわせて、NVIDIAは「Tesla K40 GPU」と「CUDA 6」を発表した。既報のように、Tesla K40は、現在のK20xから20～40%性能が向上した製品であり、CUDA 6はCPUとGPU統一メモリ(Unified Memory)を実現したというのがウリである。

しかし、どのようにして統一メモリが実現されているかについてはその詳細は分からなかった。ということで、SC13の会場で、NVIDIAのGPUコンピューティングソフトウェア部門のジェネラルマネージャのIan Buck氏への取材をアレンジして貰った。

しかし、手違いでその日には会えず、翌日、アクセラレーテッドコンピューティング部門のジェネラルマネージャのSumit Gupta氏に時間を取っていただくことになった。そこで、Gupta氏に、K40とCUDA 6の発表のプレスリリースの内容を確認し、統一メモリはどのように実現されているのかを質問すると、「技術的に詳しい内容は承知していない」という回答であった。ハード側の方なので、ソフトの実装まではご存知ないようであった。しかし、NVIDIAはSC13では最大規模の展示ブースを構え、CUDA 6の開発を行った専門家も来ているとのことで、翌日、現役のバリバリのソフトウェアアーキテクトであるStephen Jones氏から話しを聞くことができた。

良く知られているように、CPUはDDR3などのDRAMのメインメモリを使い、GPUは高いメモリバンド幅を実現するため、GDDR5などのグラフィック用のメモリを使っている。そして、CPUとGPUは、PCI Expressで繋がっている。このように、両者のメモリは別物であるので、CPUからGPUに処理を依頼する前に、CPUメモリから必要なデータをGPUメモリにコピーしておく必要がある。また、GPUでの処理が終わると、GPUメモリの中の処理結果をCPUメモリにコピーする必要がある。

そして、CPUメモリとGPUメモリのアロケーションや同期、転送などの指示をCUDAプログラムに書く必要がある。ということで、GPUコンピューティングが取っ付き難いものとなっている。この障害を取り除くのが統一メモリ(Unified Memory)である。

CUDA 6では、右側の図のようにCPUとGPUは共通のメモリをアクセスする「ように見え」、データのコピーの必要が無くなる。しかし、CPUとGPUのメモリは別個であり、「ように見え」るというところが曲者である。

NVIDIAはCUDA 4の時代から、CPUとGPUのメモリを統一した仮想アドレスでアクセスできるUVA(Unified Virtual Address)というメカニズムをサポートしている。この仕掛けで、CPUがGPUメモリをアクセスしたり、GPUがCPUメモリをアクセスしたりすることが出来るのであるが、このUVAはPCI Expressを経由してメモリアクセスを行う。CPUメモリは数10GB/s、GPUメモリは数100GB/sのメモリバンド幅を持っているが、PCI Expressを使うUVAでは、メモリバンド幅はPCI Expressのバンド幅(x16のGen2では双方向合計で16GB/s)で制限されてしまい、GPUネイティブの高いメモリバンド幅を生かせない。

CUDA 6では、次のようにして、この問題を解決している。まず、CPUとGPUで共通にアクセスするメモリ領域は、cudaMallocManaged( )という関数でメモリを確保する。そうすると、CPUメモリとGPUメモリの両方に領域が確保される。

そして、CUDA 6は、CPUからGPUのカーネルを起動する直前にcudaMallocManaged( )で確保されたCPUメモリ領域のすべてのデータを、GPUメモリの対応する領域にDMAでコピーする。これで、GPUで走るカーネルプログラムは同じ仮想アドレスにCPUからコピーされたデータが入っている状態になる。

カーネルの起動時にCPUメモリからGPUメモリへのDMAによるコピーが必要になるが、カーネルが走っている状態では、メモリバンド幅の高いGPUメモリをアクセスして動作するので、GPU本来の高い性能が得られる。