NVIDIA PascalのNVLinkとUnified Memory

@Vengineerの戯言 : Twitter
SystemVerilogの世界へようこそ、すべては、SystemC v0.9公開から始まった

8月21日の暑い日曜日に、
関東GPGPU勉強会 #4、GTX1080を使い倒す会に行ってきました。

外はめっちゃ暑かったけど、
セミナー会場はクーラーガンガンに効いていて最初は良かったのです。
が、後半寒かったです。
クーラーの時は長袖着ていますが、それでも寒かった。
短パン半袖で来ていた人、風邪ひかなかっただろうか。。。

さて本題。
最初はNVIDIAの森野さんの「NVIDIA Pascal GPUのご紹介」ということで、
NVIDIAについて、技術的な詳しい説明を。
最近のNVIDIAのセミナーは技術セミナーというより
ビジネスよりのセミナー内容になっていますね。
今回の森野さんのお話は、久しぶりに技術を聞けて非常に良かったです。
当日の資料はアップされてませんが、
NVIDIAが既に公開しているPascalの資料をピックアップしてみました。

なんといっても、

NVIDIA Tesla P100 Whitepaper

まずは、コレを読まないと始まりません。

英語版(登録要)、日本語(登録要)

なお、ここからは、英語版が直接読めます。

また、Tesla P100のホワイトペーパを読んでその特徴をまとめる(概要、NVLinkについて)というのがありました。

GPGPUの計算能力よりもH/Wアーキテクチャの方が気になったので、

　・NVLink
　・Unified Memory

の2点について、さくっと、書いてみます。

NVLink

Tesla P100は、NVLinkとPCIe Gen3 x16のインターフェースを持ち、
x86の場合は、PCIe Gen3 x16で接続し、POWER8の場合は、NVLinkで接続します。

NVLinkについては、
NVIDIAのブログ、
データの高速車線――NVLinkでアプリケーション・パフォーマンスを改善にあるように、
PCIe Gen3 x16の16GB/sの転送レートに対して、
NVLinkは1リンク当り20GB/sで4リンク搭載しているので最大40GB/sになります。
POWER9では、NVLink 2.0とPCIe Gen4になり、
もっと転送レートが良くなるようです。
その時のGPUは、Volta(ここに図がありました)のようです。

Whitepaperの24頁(図18)にP100のざっくりとした内部ブロック図があります。
これによると、HSHUB経由でNVLinkとPCIe Gen3が接続されています。
HSHUBにはHSCE(High-Speed Copy Engines)も接続していて、
このHSCEを使うことでNVLinkのピーク性能でデータの移動ができるようです。

Pascal Unified Memory

CUDA 6(Kepler)のUnified Memoryでは、
GPU側のメモリ(Device Memory)サイズまでしか利用できませんでしたが、
CUDA 8(Pascal)のUnified Memoryでは、
システムメモリのサイズまで利用できるようになった。
これにより、
Device Memoryサイズ以上のデータが必要なアプリケーションのコードが
非常に綺麗に書けるようになります。

プログラム的には、28頁(図22)にあるように

次のようなCPUコード

引用
void sortfile( FILE *fp, int N ) {
  char *data;
  data = (char *)malloc(N);

  fread(data, 1, N, fp);

  qsort(data, N, 1, compare);

  use_data(data);

  free(data);
}

が、Pascal Unified Memoryを使ったコードは、

void sortfile( FILE *fp, int N ) {
  char *data;
  data = (char *)malloc(N);

  fread(data, 1, N, fp);

  qsort<<<...>>>(data, N, 1, compare);
  cudaDevieSynchronize();

  use_data(data);

  free(data);
}

になり、メモリ管理がCPUとGPUのコードで同じになり、非常に綺麗になります。

コードとしては綺麗になって嬉しいですが、
裏ではそうとうなことをやっているはずです。

ちょっと調べてみましたが、どうやら、Linux HMMというものを使っているようです。
NVIDIAのGTC2016の資料、THE FUTURE OF UNIFIED MEMORYの43頁、「HETEROGENEOUS MEMORY MANAGER」には、

引用
　・HMM will manage a GPU page table and keep it synchronize with the CPU page table 
    Also handle DMA mapping on behalf of the device 
　・HMM allows migration of process memory to device memory CPU access will trigger fault 
    that will migrate memory back 
　・HMM is not only for GPUs, network devices can use it 
    as well Mellanox has on-demand paging mechanism, so RDMA will work in future

とあり、GPUのページテーブルとCPUのページテーブルを上手く同期するとあります。
また、GPUだけでなく、ネットワークデバイスでも利用できると。。

これって、実は、IBM CAPIがやっていることと同じだと思うんですが。。
アプリケーションが使っているユーザ空間のメモリに、
デバイス(GPUやFPGA)が簡単にアクセスできる仕組みということで。
また、AMDのHSAも同じだと思います。。。

ということで、2点、さくっと、書いてみました。

その他の資料は、
PASCAL AND CUDA 8.0
All aboard the PCIe bus for Nvidia's Tesla P100 supercomputer grunt
があります。

P.S
NVIDIAでは、まだまだ人を募集しているようです。
てか、どこも人材不足のようですね。