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は全体的にサンプルコードが不足しているような気がするので、 まず普通にハッシュテーブルを使うためのサンプルコードを載せておきます。
上記のサンプルではエラーチェックをしていません。 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);
このハッシュテーブルにはcudppMultivalueHashGetAllValues
と
cudppMultivalueHashGetValuesSize
という特別な関数が用意されています。
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実装なので ぜひ試してみてください。