【CUDA】Warp Sum & Warp Scan

Warp Sum

Warpの持つ値の合計を計算する。 例:[1, 2, 3, 4] -> [10, 10, 10, 10]

for (int i = 1; i < warpSize; i *= 2)
  value += __shfl_xor(value, i);

Warp Scan

Warpの持つ値の累積を計算する。 例:[1, 2, 3, 4] -> [1, 3, 6, 10]

for (int i = 1; i < warpSize; i *= 2) {
  int n = __shfl_up(value, i);
  if (laneId >= i) value += n;
}

サンプルコード

CUDA Warp's Sum and Scan

2017/12/11追記

warpSizeはどうやらコンパイル時定数ではないみたいです。 そのため、サンプルコードのように書くとループ展開が行われない可能性があります。 パフォーマンスを気にするのであれば warpSize32 とした方が良さそうです。

基数を取得する

最下位ビット/バイトからk番目の基数を取得する関数です。

  • radix1bit → 2進数表現で右からk番目の値を取り出す (0 or 1)
  • radix4bit → 16進数表現で右からk番目の値を取り出す (0 to 15)
  • radix8bit → 1バイト区切りで右からk番目の値を取り出す (0 to 255)

Select k-th (from LSB) radix of the number.

CUDAのハッシュテーブル CUDPP編

インストール

まずCUDPPのインストール手順です。事前にcmakeのインストールが必要です。

$ git clone -b 2.2 https://github.com/cudpp/cudpp.git
$ cd cudpp
$ git submodule init
$ git submodule update
$ mkdir build && cd build
$ cmake ..
$ make
$ sudo make install

自分の環境だとgit submodule updateに失敗してしまったので、 以下のコマンドを追加で実行しました。

$ cd ext/cub
$ git reset --hard HEAD

更に、ビルド時にcudpp_config.hが無いと言われてしまったので、以下のコマンドも。

$ sudo cp include/cudpp_config.h /usr/local/include

サンプルコード

CUDPPは全体的にサンプルコードが不足しているような気がするので、 まず普通にハッシュテーブルを使うためのサンプルコードを載せておきます。

Example of CUDPP Hash Table

上記のサンプルではエラーチェックをしていません。 CUDPPのエラーチェックをする場合は以下のようになります。

CUDPPResult result = cudppCreate(&cudpp);
if (result != CUDPP_SUCCESS) {
  fprintf(stderr, "Error in cudppCreate.\n");
  return 1;
}

コンパイルは以下のようなコマンドで行います。

$ nvcc cudpp_hash_sample.cc -lcudpp -lcudpp_hash

ハッシュテーブル

CUDPPには以下の3種類のハッシュテーブルがあります。

  • CUDPP_BASIC_HASH_TABLE … 通常のキー・値のハッシュテーブル
  • CUDPP_COMPACTING_HASH_TABLE … キーに対する一意なIDを生成するハッシュテーブル
  • CUDPP_MULTIVALUE_HASH_TABLE … 一つのキーに対して複数の値を格納できるハッシュテーブル

どのハッシュテーブルを使うかは、CUDPPHashTableConfigで設定します。

CUDPPHashTableConfig config;
config.type = CUDPP_BASIC_HASH_TABLE;

ハッシュテーブルへのデータの挿入および取得は ホストからのみ可能で、 複数のキー・値をまとめて挿入、まとめて取得します。 データの挿入、取得の際に与えるキー・値の配列は バイス上に確保 されている必要があります。

// d_keys → 挿入するキーの配列、d_vals → 挿入する値の配列、N → 挿入するキー・値の数
cudppHashInsert(hash_table_handle, d_keys, d_vals, N);

// d_input → 値を取得するキーの配列、d_output → 取得した値の書き込み先、N → 取得する値の数
cudppHashRetrieve(hash_table_handle, d_input, d_output, N);

キーおよび値には32ビットのデータしか使えない ので注意してください。ここからは各ハッシュテーブルの使い方を見ていきます。

CUDPP_BASIC_HASH_TABLE

一つのキーに対して一つの値を格納できるハッシュテーブルです。 キーは重複してはいけません。また、格納していないキーの値を取得しようとすると CUDPP_HASH_KEY_NOT_FOUNDが返ってきます。

// keys = {2, 4, 0, 3, 1}, vals = {0, 1, 2, 3, 4};
cudppHashInsert(hash_table_handle, keys, vals, 5);

// input = {0, 1, 2, 3, 4, 0, 2, 4} --> output = {2, 4, 0, 3, 1, 2, 0, 1}
cudppHashRetrieve(hash_table_handle, input, output, 8);

CUDPP_COMPACTING_HASH_TABLE

キーに対する一意なIDを生成するハッシュテーブルです。 例えば{16, 34, 81, 16, 52}という入力に対して、{16=>0, 34=>1, 81=>2, 52=>3}のような ハッシュテーブルを作ることができます。キーは重複していても問題ありません。 また、値の配列は不要なのでNULLを指定すればOKです。

// keys = {16, 34, 81, 16, 52}
cudppHashInsert(hash_table_handle, keys, NULL, 5);

// input = {16, 34, 81, 16, 52} --> output = {2, 3, 1, 2, 0}
cudppHashRetrieve(hash_table_handle, input, output, 5);

CUDPP_MULTIVALUE_HASH_TABLE

一つのキーに対して複数の値を格納できるハッシュテーブルです。 キーが重複していた場合に、同じキーに対して複数の値が設定されます。

// keys = {3, 1, 3, 3, 2, 1}, vals = {0, 1, 2, 3, 4, 5};
cudppHashInsert(hash_table_handle, keys, vals, 6);

このハッシュテーブルにはcudppMultivalueHashGetAllValuescudppMultivalueHashGetValuesSizeという特別な関数が用意されています。 cudppMultivalueHashGetAllValuesでは全ての値が格納された配列を返します。 値は以下の図のように、各キー毎にまとめて並べられています。

|     key1    | key2 |        key3        | ...
| val1 | val2 | val1 | val1 | val2 | val3 | ...

cudppMultivalueHashGetValuesSizeでは、その配列のサイズを取得できます。

// hash_table = {1 => {1, 5}, 2 => {4}, 3 => {0, 2, 3}}

unsigned int size;
cudppMultivalueHashGetValuesSize(hash_table_handle, &size);
// size = 6

unsigned int *d_all_vals;
cudppMultivalueHashGetAllValues(hash_table_handle, &d_all_vals);
// d_all_vals = {1, 5, 4, 0, 2, 3}

このハッシュテーブルでは他の二つと異なり、cudppHashRetrieveで64ビットのデータを返します。 前半32ビットには先ほど取得した配列上のインデックス、後半32ビットには値の数が表されています。

// hash_table = {1 => {1, 5}, 2 => {4}, 3 => {0, 2, 3}}

uint2 *d_output;
cudaMalloc((void**)&d_output, sizeof(uint2) * 3);
cudaMemset(d_output, 0, sizeof(uint2) * 3);

// input = {1, 2, 3} --> output = {{0, 2}, {2, 1}, {3, 3}}
cudppHashRetrieve(hash_table_handle, d_input, d_output, 3);

// 値を表示
uint2 h_output[3];
cudaMemcpy(h_output, d_output, sizeof(int) * 3, cudaMemcpyDeviceToHost);
for (int i = 0; i < 3; ++i) {
  for (int j = 0; j < output[i].y; ++j) {
    printf("key = %d, val = %d\n", input[i], all_vals[output[i].x + j]);
  }
}

まとめ

CUDPPの3種類のハッシュテーブルの使い方を見ていきました。 少し使い方にクセがあるかもしれませんが、貴重なハッシュテーブルのCUDA実装なので ぜひ試してみてください。

Google Homeでのやり取りの履歴を見る

Google Homeに話しかけた内容、また、Google Homeの返答の履歴を見る方法です。

  1. こちらのページを開く
  2. 検索の下の"日付とサービスでフィルタ"を選択
  3. アシスタントにチェック
  4. 検索の虫眼鏡ボタンを押すと、Google Homeとのやり取りの履歴が出てきます

スマートフォン等で利用したGoogleアシスタントの履歴も含まれます。

f:id:dfukunaga:20171029224454p:plain

f:id:dfukunaga:20171029224513p:plain

こんな感じで履歴が表示されます。自分の声も聞くことができます。

f:id:dfukunaga:20171029224445p:plain

【CUDA】カーネル内での動的メモリ確保

バイス上のメモリは、ホストからのcudaMalloc()で確保するのが通常だが、 カーネル内で動的にメモリ確保することもできる*1カーネル内での動的メモリ確保には、C/C++と同じようにmalloc()/free()new/deleteが使える。

__global__
void testKernel() {
  size_t size = 1024;

  // Dynamic allocation with malloc/free
  char *a = (char*)malloc(size);
  free(a);

  // Dynamic allocation with new/delete
  char *b = new char[size];
  delete[] b;
}

int main() {
  testKernel<<<1, 1>>>();
  cudaDeviceSynchronize();
  return 0;
}

ただし、ホストからのcudaMalloc()カーネル内でのmalloc()はメモリ確保方法が異なる。 ホストからのcudaMalloc()はデバイスメモリの未使用の領域にメモリを確保するのに対して、 カーネル内でのmalloc()はあらかじめ用意されたデバイス上のヒープにメモリを確保する。 そのため、ホストからのcudaMalloc()で確保したメモリはカーネル内では解放できないし、 カーネル内で確保したメモリはホストからは解放できない。 以下のコードはエラーになるパターンを表している。

__global__
void mallocInKernel(char **ptr, size_t size) {
  *ptr = (char*)malloc(size);
}

__global__
void freeInKernel(char *ptr) {
  free(ptr);
}

int main() {
  size_t size = 1024;

  // ホスト上でcudaMalloc()したメモリをカーネル内でfree()
  char *a;
  cudaMalloc(&a, size);
  freeInKernel<<<1, 1>>>(a);  // Error!!
  cudaDeviceSynchronize();

  // カーネル内でmalloc()したメモリをホストからcudaFree()
  char *hb, **db;
  cudaMalloc(&db, sizeof(char*));
  mallocInKernel<<<1, 1>>>(db, size);
  cudaMemcpy(&hb, db, sizeof(char*), cudaMemcpyDeviceToHost);
  cudaFree(hb);  // Error!!
  cudaDeviceSynchronize();

  return 0;
}

バイスのヒープはデフォルトで8MBしか確保されない。 つまり、カーネル内のmalloc()で8MB以上のメモリを確保するとエラーになる。 この設定を変えるにはcudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)を使えばよい。 例えば、以下のように書くと1GBまでカーネル内でmalloc()できる。

size_t size = 1024 * 1024 * 1024;
cudaDeviceSetLimit(cudaLimitMallocHeapSize, size);

この設定は全てのカーネル実行が行われる前に書かれている必要があり、途中で設定を変えることはできない。 また、現在のヒープサイズの設定がどうなっているか確認するためには cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)を使えばよい。

size_t size;
cudaDeviceGetLimit(&size, cudaLimitMallocHeapSize);
printf("Heap Size=%ld\n", size);

Compute Capability 3.5以降ではカーネル内でcudaMalloc()/cudaFree()を使うことができるが、 これはカーネル内でのmalloc()/free()と同じ効果をもたらす

バイスのヒープサイズを無駄に拡げないためにも、できるだけホストからのcudaMalloc()で メモリを確保したいが、場合によってはカーネル内での動的メモリ確保が有効なケースもあるだろう。

*1:Compute Capability 2.x以降でのみサポート

NVIDIA Jetson TX2を買ったのでセットアップ

届きました!

f:id:dfukunaga:20170902213325p:plain

意外と大きいです。一辺20cmくらいでしょうか。 ディスプレイ(HDMI)、キーボード、マウス、電源を接続します。USBポートは一基しかないので、キーボードとマウスを両方接続する場合は USBハブを使うか、付属のMicroUSB-USB変換ケーブルを使います。

f:id:dfukunaga:20170902213757p:plain

電源ボタンを押すとUbuntu 16.04のコンソールが開きます。 まずは以下のコマンドでドライバをインストールします。 (sudo時のパスワードはnvidia

$ cd NVIDIA-INSTALLER
$ sudo ./installer.sh
$ sudo reboot

再起動するとGUIが立ち上がります。 Wi-Fiもしくは有線でネットワークの設定を行っておきます。 (後で使うので、割り振られたIPアドレスを確認しておきます。)

ここから、JetPackのインストールです。

JetPackとは、CUDAやcuDNNを始め、Jetsonを使うのに必要そうなソフトウェアをまとめたものっぽいです。 筆者はMacbook Proを使っているので、VirtualBox上のUbuntu 16.04をホストとしてJetPackをインストールすることにしました。

こちらからUbuntu 16.04のイメージをダウンロードして、 VirtualBoxVMを作ります。JetPackのインストールに最低10GBのディスクスペースが必要とのことなのですが、 ディスクサイズは余裕を持って30GBに設定しました。

VMを起動してこちらからJetPackをダウンロードします。 NVIDIA Developer Programへの登録が必要になります。ダウンロードしたら以下のコマンドで実行。

$ chmod +x JetPack-L4T-3.1-linux-x64.run
$ ./JetPack-L4T-3.1-linux-x64.run

途中の"Component Manager"画面でFlash OS Image to TargetのActionをno actionに変更します。 各種パッケージをインストールした後、 “Device Information - Jetson TX2"画面でJetsonのIPアドレスとユーザ名、パスワード(共にnvidia)を入力します。

新しいウィンドウが開き、Jetson側にいろいろインストールされていき、 Installation of target components finished, close this window to continue.と表示されたら終わりです。 /home/nvidia/NVIDIA_CUDA-8.0_SamplesにCUDAのサンプルが入っているので、 これがうまく動けば少なくともCUDAはちゃんと入っているはずです。

ちなみに、JetsonではNVIDIAGPUでおなじみのnvidia-smiコマンドは対応していないみたいです。 代わりにホームディレクトリにあるtegrastatsGPUの状態を確認します。

nvidia@tegra-ubuntu:~$ sudo ./tegrastats 
RAM 1103/7851MB (lfb 1501x4MB) cpu [0%@1728,off,off,0%@1728,0%@1727,0%@1728] EMC 0%@1600 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [0%@345,off,off,1%@345,0%@346,5%@345] EMC 8%@40 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [1%@345,off,off,1%@345,0%@345,5%@345] EMC 8%@40 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [1%@345,off,off,2%@345,0%@345,4%@339] EMC 8%@40 APE 150 GR3D 0%@114

GR3Dの右がGPU使用率だそうです。GPUメモリはシステムと共有らしいので一番左の列を見ればOKです。

ZeroMQ: 複数のEndpointに接続する

f:id:dfukunaga:20170419234028p:plain

ちゃんとドキュメントを読めという話かもしれないが、ZeroMQ のソケットが 複数のEndpointにbind/connectできることを最近知った。そこで、複数のEndpointに対してのbind/connectを 使ったメッセージングのパターンをいくつか考えてみた。

1. ローカル/リモート共用ソケット

f:id:dfukunaga:20170418233138p:plain

同じプロセス内でのソケット間通信をするなら、inprocを使ったほうが ローカルループバックに対するtcpを使うよりも圧倒的に速い。このパターンでは、 一つのソケットに同一プロセスからの接続用のアドレス(ローカル: inproc)と、 別マシンからの接続用のアドレス(リモート: tcp)の二つをbindし、接続するソケットの位置 によってconnectするアドレスを使い分ける。これにより、ローカルのソケット間通信はより高速に行われ、また、 受信する側では、ローカルから来たメッセージとリモートから来たメッセージを区別することなく処理を行うことができる。

2. ブロードキャストクラスタ

f:id:dfukunaga:20170418233145p:plain

各ノードでZMQ_PUBZMQ_SUBを一つずつ用意します。 ZMQ_PUBはbindしてZMQ_SUBは各ノードのZMQ_PUBにconnectすると(逆でもたぶんOK)、 ZMQ_PUBにsendすることで他の全ノードにメッセージが送信されるようになります。 複数のソケットを用意すれば同じことが実現できますが、一つのソケットで複数のEndpointにconnectした方が、 recv/pollするのが一つのソケットで済むので楽になるかと思います。 このパターンを実装するときはZMQ_SUBSCRIBEを設定するのを忘れずに。

3. P2Pクラスタ

f:id:dfukunaga:20170418233158p:plain

bindとconnectを併用することもできます。このパターンでは、各ノードにZMQ_ROUTERを一つ用意し、 一つのアドレスに対してbindするのと同時に、他のノードのZMQ_ROUTERにconnectします。 このように繋ぐことで、ZMQ_ROUTERのルーティング機能を使えばクラスタ内のどのノードへも メッセージを送信することができます。ZMQ_ROUTER同士を繋ぐときはidentityの設定が厄介ですが、 bindする前にZMQ_IDENTITYを設定し、connectする際にZMQ_CONNECT_RIDを設定しておけば問題ないでしょう。

以上、複数のEndpointへのbind/connectを用いたメッセージングパターンでした。