(前編はコチラ)

一方、GPUカーネルを起動のためにGPUメモリにデータをコピーすると、CPUアドレス変換機構は、赤の×で示したように、CPUメモリとの結合を切り離し、その仮想アドレスに対応するCPU側の物理メモリは存在しないという状態にする。つまり、GPUでカーネルが走っている間は、CPU側には、論理的にはその仮想アドレスのデータが存在せず、アクセスできないようになる。従って、その後、GPUがその仮想アドレスのデータを書き換えても矛盾は生じない。

そして、GPUカーネルでの処理が終わり、結果がGPUメモリに書き込まれると、CPUが統一メモリ領域に書き込まれたデータを読もうとする。しかし、CPUのアドレス変換機構では、その仮想アドレスに対応する物理メモリは存在しないことになっており、ページフォールトが発生する。CUDA 6では、このページフォールトで、アクセスした仮想アドレスのページのデータをGPUメモリからCPUメモリにPCI Express経由でコピーし、CPUのアドレス変換機構にその仮想アドレスに対応する物理ページが存在する状態にして、アクセスを行う。

GPUメモリからCPUメモリへのデータ転送は、ページ単位でデマンド転送を行う

このように、CPUからGPUへのデータの受け渡しはcudaMallocManaged( )で確保したすべての領域に対して一括して行われるが、GPUからCPUへのデータの受け渡しはCPUがアクセスした4KBのページ単位で転送が行われ、CPUがアクセスしないページのデータは転送されないようになっている。

CPUからGPUへの転送もオンデマンドの転送にしなかったのは、x86 CPUのメモリ管理のページサイズは4KBであるが、GPU側のアドレス変換機構は、より大きいサイズのページを使っているので、CPUメモリとGPUメモリのコヒーレンスが確実に維持できないケースがあるためであるという。なお、次世代のMaxwell GPUではページサイズを合わせるので、技術的には、双方向でオンデマンドのページ単位での転送が可能になるとのことである。

cudaMallocManaged( )で定義したメモリ領域は、GPUカーネルとの統一メモリとなり、カーネルを起動する直前にデータのコピーが行われる。しかし、複数のカーネルを使う場合は、あるメモリ領域はカーネルAとCPUで統一メモリであるが、カーネルBはそのメモリ領域は使わないという場合がある。このため、cudaMallocManaged( )は、デフォルトではすべてのカーネルとの統一メモリを確保してデータを一括転送するが、どのカーネルとの統一メモリであるかを指定することができるようになっており、不要なデータ転送を行わないようにできるようになっている。

なお、複数GPUを使う場合、複数のGPUのメモリに同じ仮想アドレスのコピーを作ると、GPUからの書き込みによって矛盾が生じてしまう。このため、cudaMallocManaged( )は1つのGPUのメモリだけにコピーを作り、他のGPUがその仮想アドレスのデータをアクセスする場合は、PCI Express経由のアクセスとなる。

このように、CUDA 6では、cudaMallocManaged( )で統一メモリに「見える」領域を定義すると、自動的にデータ転送を行ってくれ、プログラムに明示的にメモリ転送を書く必要がなくなる。また、データを転送する時間は当然必要であるが、一度転送を行ってしまえば、CPUは通常のCPUメモリのアクセス、GPUも通常のGPUメモリの速度でアクセスができ、UVAのように、PCI Expressの性能で制限されることはない。

AMDのHSAは物理的にも共有のメモリをCPUとGPUがアクセスする共有メモリであるが、この方式ではGDDRメモリの高バンド幅とDDR3/4メモリの容量を両立させることはできない。これに対して、NVIDIAのCUDA 6の方式は、GPU処理の開始と終了時にメモリ間のコピーが必要になるが、コピーが終われば、その後は、CPUはDDR3/4、GPUはGDDR5のバンド幅とレーテンシでのアクセスが可能になる。

AMDの方式は、1チップにCPUとGPUを搭載する小規模な構成ではメリットがあるが、CPUには大容量のメモリが必要で、GPUには高バンド幅のGDDRメモリを付けて高い演算性能が必要というスパコンなど規模の大きいシステムでは、CUDA 6のやり方が適していると思われる。

しかし、Micron Technologyが推進する高メモリバンド幅で、メモリ容量の拡張も可能というHMC(Hybrid Memory Cube)というテクノロジを使えばAMDのHSAの問題である高バンド幅とメモリ容量の両立という問題を解決できる可能性もある。CPUとGPUを使って電力効率の高いスパコンを実現するというのは、スパコン業界の流れであり、AMDの物理的にも共通のメモリを使う方式と、NVIDIAのオンデマンド転送方式のいずれが主流になるのか、興味深いところである。