CPUとGPUの連携を簡素化するUnified Memory

CPUの場合は、2つのCPUチップが使われている場合でも、これらのCPUは共通のメモリをアクセスし、一方のCPUが書き込んだ結果を他方のCPUが読み出すことができる。これに対して、一方がCPUで、もう一方がGPUの場合は、次の図の左側の絵のように、CPUはシステムメモリ、GPUはGPUメモリという独立したメモリを持つという構造が使われてきた。

GPUが描画だけに使われ、CPUがシステムメモリに格納した描画データをGPUが読み出して、GPUメモリを使って描画処理を行い、その結果をGPUメモリの中のフレームバッファに書き出し、フレームメモリの内容をGPUに内蔵された専用のハードウェアでディスプレイに転送するという一方向に流れる処理を行っているときは、この独立メモリの構造でも、大きな不便はなかった。

しかし、科学技術計算やゲームの場合でもPhysXなどでシミュレーションを行う場合には、もっと複雑にCPUとGPUの間でデータのやり取りを行うことが必要になる。そうなると、システムメモリとGPUメモリの間のデータのコピーをいちいちプログラムで記述してやるのは面倒だし、間違いも多くなる。

CPUとGPU間のデータ転送を簡単化するUnified Memory

このために考えられたのが、「Unified Memory」である。Unified Memoryは物理的には別物の2つのメモリをあたかも1つのメモリのように扱えるようにするという技術である。

次の図の左側はCPUでSortを行うプログラムで、右側がUnified Memoryを使うGPUシステム用のコードである。違いは、CPU側ではmalloc( )でデータを格納するメモリを確保している部分が、cudaMallocManaged( )という関数の呼び出しになっている点と、qsortの処理がGPUのスレッドグリッドで行われるように書かれ、全スレッドの処理を同期するためのcudaDeviceSynchronize( )の呼び出しが追加されている点である。

NVIDIAは、2013年にCUDA6+Kepler GPUでUnified Memoryをサポートした。これはシステムメモリとGPUメモリにそれぞれUnified Memory用の領域を確保し、CPUがデータを使う時には、システム(CUDAドライバとランタイム)が自動的にCPU側のUnified Memory領域にデータを持ってきて、GPUがデータを使う場合は、これもシステムが自動的にデータを持ってくるというシステムである。

Unified Memoryはどのように実現されているのか

その仕掛けは、多少、筆者の推測を含んでいるが、次の図のようになっている。CPU、GPUともに、ページテーブルを使うメモリ管理機構を持っている。アプリケーションが使う仮想アドレス(VA)でページテーブルを引くと、そのVAを含む物理ページのアドレスが得られるので、ハードウェアは、この物理アドレスを使ってシステムメモリ、あるいはGPUメモリをアクセスする。

システム全体で1つのVA体系をなすUnified Virtual Address Spaceの場合は、同じVAがCPU側のページテーブルとGPU側のページテーブルに出てくることはないが、Unified Memoryの範囲のアドレスの場合は、両方のページテーブルに、同じVAに対するページテーブルエントリを持たせる。ただし、それぞれのページテーブルのValidビットは、どちらか一方がValid、他方はInvalidとなるように制御する。

したがって、CPU側のページテーブルエントリがValidの場合はシステムメモリが使われ、GPU側のページテーブルエントリがValidの場合はGPUメモリが使われる。

しかし、GPUメモリが使われている状態で、CPUがそのVAをアクセスしようとすると、CPU側のページテーブルエントリはInvalidであるので、ページフォールトの割り込みが発生し、GPUのメモリドライバが介入してGPU側のページテーブルエントリをInvalidにしてGPUからアクセスできないようにする。そして、GPUメモリのそのページのデータをシステムメモリにコピーして、次にCPU側のページテーブルエントリをValidに変更して、CPUのアクセスをやり直させる。今度は、CPU側のページテーブルエントリはValidなので、システムメモリをアクセスして処理を続ける。

このようにすれば、CPUがアクセスしたときには、GPUに有ったページが自動的にシステムメモリに移動したように動く。メデタシメデタシであるが、実はPascal以前のKepler GPUでは、問題があった。