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実装なので ぜひ試してみてください。