Githubに投稿しました 新しいプロジェクト シンプルな GPU ハッシュ テーブル.

これは、1060 秒あたり数億件の挿入を処理できるシンプルな GPU ハッシュ テーブルです。 私の NVIDIA GTX 64 ラップトップでは、コードはランダムに生成された 210 万のキーと値のペアを約 32 ミリ秒で挿入し、64 万のペアを約 XNUMX ミリ秒で削除します。

つまり、ラップトップの速度は約 300 億挿入/秒、500 億削除/秒です。

テーブルは CUDA で作成されていますが、同じ手法を HLSL または GLSL に適用することもできます。 ビデオ カードで高いパフォーマンスを確保するには、実装にはいくつかの制限があります。

• 32 ビットのキーと同じ値のみが処理されます。
• ハッシュ テーブルのサイズは固定です。
• そして、このサイズは XNUMX のべき乗に等しくなければなりません。

キーと値については、単純な区切りマーカーを予約する必要があります (上記のコードでは、これは 0xffffffff です)。

ロックのないハッシュ テーブル

ハッシュ テーブルはオープン アドレス指定を使用します。 リニアプロービングつまり、メモリに保存され、優れたキャッシュ パフォーマンスを備えた単なるキーと値のペアの配列です。 同じことは、リンクされたリスト内のポインターの検索を伴う連鎖の場合には言えません。 ハッシュ テーブルは要素を格納する単純な配列です KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

pow2/AND マスクを適用するには XNUMX つの高速命令で十分ですが、モジュラス演算子ははるかに遅いため、テーブルのサイズは素数ではなく XNUMX のべき乗です。 線形テーブル ルックアップではスロット インデックスを各スロットでラップする必要があるため、これは線形プローブの場合に重要です。 その結果、操作のコストがスロットごとにモジュロで追加されます。

テーブルには各要素のキーと値のみが保存され、キーのハッシュは保存されません。 テーブルには 32 ビットのキーのみが格納されるため、ハッシュは非常に迅速に計算されます。 上記のコードは Murmur3 ハッシュを使用しており、数回のシフト、XOR、乗算のみを実行します。

ハッシュ テーブルは、メモリの順序に依存しないロック保護技術を使用します。 一部の書き込み操作によって他のそのような操作の順序が混乱した場合でも、ハッシュ テーブルは依然として正しい状態を維持します。 これについては以下で説明します。 この技術は、数千のスレッドを同時に実行するビデオ カードでうまく機能します。

ハッシュ テーブルのキーと値は空に初期化されます。

コードを変更して、64 ビットのキーと値も処理することもできます。 キーにはアトミックな読み取り、書き込み、比較および交換の操作が必要です。 また、値にはアトミックな読み取りおよび書き込み操作が必要です。 幸いなことに、CUDA では、32 ビット値と 64 ビット値の読み取り/書き込み操作は、それらが自然にアライメントされている限りアトミックです (以下を参照)。 ここで)、最新のビデオ カードは 64 ビットのアトミック比較および交換操作をサポートしています。 もちろん、64 ビットに移行すると、パフォーマンスは若干低下します。

ハッシュテーブルの状態

ハッシュ テーブル内の各キーと値のペアは、次の XNUMX つの状態のいずれかになります。

• キーと値が空です。 この状態でハッシュテーブルが初期化されます。
• キーは書き留められていますが、値はまだ書き込まれていません。 別のスレッドが現在データを読み取っている場合、そのスレッドは空を返します。 これは正常なことであり、別の実行スレッドがもう少し早く動作していれば同じことが起こっていたでしょう。ここで話しているのは同時データ構造です。
• キーと値の両方が記録されます。
• 値は他の実行スレッドで使用できますが、キーはまだ使用できません。 これは、CUDA プログラミング モデルのメモリ モデルの順序が緩いために発生する可能性があります。 これは正常であり、いずれにしても、値が空でなくなっても、キーは空のままです。

重要なニュアンスは、キーがスロットに書き込まれると、キーが削除されても移動しなくなるということです。これについては以下で説明します。

ハッシュ テーブル コードは、メモリが読み書きされる順序が不明な、緩やかに順序付けされたメモリ モデルでも機能します。 ハッシュ テーブルでの挿入、検索、削除を確認するときは、各キーと値のペアが上記の XNUMX つの状態のいずれかにあることに注意してください。

ハッシュテーブルへの挿入

キーと値のペアをハッシュ テーブルに挿入する CUDA 関数は次のようになります。

void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
        if (prev == kEmpty || prev == key)
        {
            hashtable[slot].value = value;
            break;
        }
        slot = (slot + 1) & (kHashTableCapacity-1);
    }
}

キーを挿入するには、コードは挿入されたキーのハッシュから始まるハッシュ テーブル配列を反復処理します。 配列内の各スロットは、そのスロット内のキーを空と比較するアトミックな比較交換操作を実行します。 不一致が検出された場合、スロット内のキーは挿入されたキーで更新され、元のスロット キーが返されます。 この元のキーが空であるか、挿入されたキーと一致する場合、コードは挿入に適したスロットを見つけて、挿入された値をそのスロットに挿入します。

XNUMX つのカーネル呼び出しの場合 gpu_hashtable_insert() 同じキーを持つ要素が複数ある場合、それらの値のいずれかをキー スロットに書き込むことができます。 これは正常であると考えられます。呼び出し中のキーと値の書き込みの XNUMX つは成功しますが、これらすべてが複数の実行スレッド内で並行して行われるため、どのメモリへの書き込みが最後の書き込みになるかを予測できません。

ハッシュテーブルルックアップ

キーを検索するコード:

uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
        uint32_t slot = hash(key);

        while (true)
        {
            if (hashtable[slot].key == key)
            {
                return hashtable[slot].value;
            }
            if (hashtable[slot].key == kEmpty)
            {
                return kEmpty;
            }
            slot = (slot + 1) & (kHashTableCapacity - 1);
        }
}

テーブルに格納されているキーの値を見つけるには、探しているキーのハッシュから始まる配列を反復処理します。 各スロットで、キーが探しているものであるかどうかを確認し、そうである場合はその値を返します。 また、キーが空かどうかも確認し、空の場合は検索を中止します。

キーが見つからない場合、コードは空の値を返します。

これらの検索操作はすべて、挿入と削除を通じて同時に実行できます。 テーブル内の各ペアは、フローに関して上で説明した XNUMX つの状態のうちの XNUMX つを持ちます。

ハッシュテーブルでの削除

キーを削除するコード:

void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        if (hashtable[slot].key == key)
        {
            hashtable[slot].value = kEmpty;
            return;
        }
        if (hashtable[slot].key == kEmpty)
        {
            return;
        }
        slot = (slot + 1) & (kHashTableCapacity - 1);
    }
}

キーの削除は通常とは異なる方法で行われます。キーをテーブルに残し、その値 (キー自体ではなく) を空としてマークします。 このコードは以下に非常によく似ています lookup()ただし、キーで一致が見つかった場合、その値は空になります。

前述したように、キーがスロットに書き込まれると、それ以降は移動されなくなります。 要素がテーブルから削除されても、キーはそのまま残り、その値は単に空になります。 これは、現在の値が空かどうかは関係なく、依然として空になるため、スロット値に対してアトミックな書き込み操作を使用する必要がないことを意味します。

ハッシュテーブルのサイズ変更

ハッシュ テーブルのサイズを変更するには、より大きなテーブルを作成し、古いテーブルから空でない要素をそこに挿入します。 サンプル コードをシンプルに保ちたかったので、この機能は実装しませんでした。 さらに、CUDA プログラムでは、メモリ割り当ては CUDA カーネルではなくホスト コードで行われることがよくあります。

記事 ロックフリー、ウェイトフリーのハッシュテーブル では、このようなロックで保護されたデータ構造を変更する方法について説明します。

競争力

上記の関数コードスニペットでは gpu_hashtable_insert(), _lookup() и _delete() 一度に XNUMX つのキーと値のペアを処理します。 そして、より低い gpu_hashtable_insert(), _lookup() и _delete() ペアの配列を並列に処理し、各ペアを個別の GPU 実行スレッドで処理します。

// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);

// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
    unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
    if (threadid < numkvs)
    {
        gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
    }
}

ロック耐性のあるハッシュ テーブルは、同時挿入、検索、および削除をサポートします。 キーと値のペアは常に XNUMX つの状態のいずれかにあり、キーは移動しないため、異なる種類の操作が同時に使用された場合でも、テーブルの正確性が保証されます。

ただし、挿入と削除のバッチを並行して処理し、ペアの入力配列に重複キーが含まれている場合、どのペアが「勝つ」か、つまりハッシュ テーブルに最後に書き込まれるペアを予測することはできません。 ペアの入力配列を使用して挿入コードを呼び出したとします。 A/0 B/1 A/2 C/3 A/4。 コードが完了したら、ペアリングします B/1 и C/3 テーブル内に存在することが保証されていますが、同時にいずれかのペアがテーブル内に表示されます。 A/0, A/2 または A/4。 これは問題になる場合もあれば、問題にならない場合もあります。すべてはアプリケーションによって異なります。 入力配列に重複キーがないことが事前にわかっている場合や、どの値が最後に書き込まれたかを気にしない場合もあります。

これが問題になる場合は、重複したペアを異なる CUDA システム コールに分割する必要があります。 CUDA では、カーネルを呼び出す操作はすべて、次のカーネル呼び出しの前に常に完了します (少なくとも XNUMX つのスレッド内では。異なるスレッドでは、カーネルは並列実行されます)。 上記の例で、XNUMX つのカーネルを次のように呼び出した場合、 A/0 B/1 A/2 C/3、および他の A/4、次にキー A 値を取得します 4.

次に、関数が次のようにすべきかどうかについて話しましょう。 lookup() и delete() ハッシュ テーブル内のペアの配列へのプレーンまたは揮発性ポインタを使用します。 CUDA ドキュメント 次のように述べています。

コンパイラは、グローバル メモリまたは共有メモリへの読み取りと書き込みを最適化することを選択できます。これらの最適化は、キーワードを使用して無効にできます。 volatile: ... この変数への参照はすべて、実メモリの読み取りまたは書き込み命令にコンパイルされます。

正確性に関する考慮事項には適用は必要ありません volatile。 実行スレッドが以前の読み取り操作でキャッシュされた値を使用する場合、わずかに古い情報を使用することになります。 ただし、これはカーネル呼び出しのある時点でのハッシュ テーブルの正しい状態からの情報です。 最新の情報を使用する必要がある場合は、インデックスを使用できます。 volatileただし、その後パフォーマンスはわずかに低下します。私のテストによると、32 万個の要素を削除すると、速度は 500 億削除/秒から 450 億 XNUMX 万削除/秒に低下しました。

Производительность

64 万個の要素を挿入し、そのうち 32 万個を削除するテストでは、 std::unordered_map そして、GPU のハッシュ テーブルは事実上ありません。

std::unordered_map 要素の挿入と削除、および要素の解放に 70 ミリ秒を費やしました unordered_map (何百万もの要素を削除するには多大な時間がかかります。 unordered_map 複数のメモリ割り当てが行われます)。 正直言えば、 std:unordered_map 全く異なる制限。 これは単一の CPU スレッドで実行され、あらゆるサイズの Key-Value をサポートし、高い使用率で優れたパフォーマンスを発揮し、複数の削除後も安定したパフォーマンスを示します。

GPU とプログラム間通信のハッシュ テーブルの継続時間は 984 ミリ秒でした。 これには、テーブルのメモリへの配置と削除 (一度に 1 GB のメモリを割り当てるため、CUDA では時間がかかります)、要素の挿入と削除、およびそれらの反復に費やされる時間が含まれます。 ビデオ カード メモリとの間のすべてのコピーも考慮されます。

ハッシュ テーブル自体が完了するまでに 271 ミリ秒かかりました。 これには、ビデオ カードが要素の挿入と削除に費やした時間が含まれますが、メモリへのコピーと結果のテーブルの反復に費やした時間は考慮されていません。 GPU テーブルが長期間存続する場合、またはハッシュ テーブル全体がビデオ カードのメモリ内に含まれている場合 (たとえば、中央プロセッサではなく他の GPU コードで使用されるハッシュ テーブルを作成する場合)、テスト結果が関係します。

ビデオ カードのハッシュ テーブルは、高スループットとアクティブな並列化により高いパフォーマンスを発揮します。

制限事項

ハッシュ テーブルのアーキテクチャには、注意すべき問題がいくつかあります。

• 線形プローブはクラスタリングによって妨げられ、テーブル内のキーが完全に配置されなくなります。
• 関数を使用してキーは削除されません delete そして時間が経つとテーブルが散らかってしまいます。

その結果、特にハッシュ テーブルが長期間存在し、多数の挿入と削除が行われた場合、ハッシュ テーブルのパフォーマンスが徐々に低下する可能性があります。 これらの欠点を軽減する XNUMX つの方法は、かなり低い使用率で新しいテーブルに再ハッシュし、再ハッシュ中に削除されたキーをフィルターで除外することです。

説明した問題を説明するために、上記のコードを使用して 128 億 4 万の要素を持つテーブルを作成し、124 億 0,96 万のスロットが埋まるまで 4 万の要素をループします (使用率は約 XNUMX)。 これは結果テーブルです。各行は、XNUMX 万の新しい要素を XNUMX つのハッシュ テーブルに挿入するための CUDA カーネル呼び出しです。

使用率
挿入時間 4 要素

0,00
11,608448 ミリ秒 (361,314798 万キー/秒)

0,03
11,751424 ミリ秒 (356,918799 万キー/秒)

0,06
11,942592 ミリ秒 (351,205515 万キー/秒)

0,09
12,081120 ミリ秒 (347,178429 万キー/秒)

0,12
12,242560 ミリ秒 (342,600233 万キー/秒)

0,16
12,396448 ミリ秒 (338,347235 万キー/秒)

0,19
12,533024 ミリ秒 (334,660176 万キー/秒)

0,22
12,703328 ミリ秒 (330,173626 万キー/秒)

0,25
12,884512 ミリ秒 (325,530693 万キー/秒)

0,28
13,033472 ミリ秒 (321,810182 万キー/秒)

0,31
13,239296 ミリ秒 (316,807174 万キー/秒)

0,34
13,392448 ミリ秒 (313,184256 万キー/秒)

0,37
13,624000 ミリ秒 (307,861434 万キー/秒)

0,41
13,875520 ミリ秒 (302,280855 万キー/秒)

0,44
14,126528 ミリ秒 (296,909756 万キー/秒)

0,47
14,399328 ミリ秒 (291,284699 万キー/秒)

0,50
14,690304 ミリ秒 (285,515123 万キー/秒)

0,53
15,039136 ミリ秒 (278,892623 万キー/秒)

0,56
15,478656 ミリ秒 (270,973402 万キー/秒)

0,59
15,985664 ミリ秒 (262,379092 万キー/秒)

0,62
16,668673 ミリ秒 (251,627968 万キー/秒)

0,66
17,587200 ミリ秒 (238,486174 万キー/秒)

0,69
18,690048 ミリ秒 (224,413765 万キー/秒)

0,72
20,278816 ミリ秒 (206,831789 万キー/秒)

0,75
22,545408 ミリ秒 (186,038058 万キー/秒)

0,78
26,053312 ミリ秒 (160,989275 万キー/秒)

0,81
31,895008 ミリ秒 (131,503463 万キー/秒)

0,84
42,103294 ミリ秒 (99,619378 万キー/秒)

0,87
61,849056 ミリ秒 (67,815164 万キー/秒)

0,90
105,695999 ミリ秒 (39,682713 万キー/秒)

0,94
240,204636 ミリ秒 (17,461378 万キー/秒)

使用率が増加すると、パフォーマンスが低下します。 これはほとんどの場合、望ましくないことです。 アプリケーションがテーブルに要素を挿入し、その後それらを破棄する場合 (たとえば、本の中の単語を数える場合)、これは問題ではありません。 ただし、アプリケーションが長期間存続するハッシュ テーブルを使用する場合 (たとえば、ユーザーが情報を頻繁に挿入および削除する画像の空でない部分を保存するグラフィック エディターなど)、この動作は問題となる可能性があります。

そして、64 万回の挿入後のハッシュ テーブルのプローブの深さを測定しました (利用率 0,5)。 平均の深さは 0,4774 であったため、ほとんどのキーは可能な限り最適なスロット、または最適な位置から 60 スロット離れた場所にありました。 最大深度はXNUMXでした。

次に、124 億 0,97 万個のインサートを備えたテーブルのプローブ深さを測定しました (利用率 10,1757)。 平均深度はすでに XNUMX であり、最大深度は - 6474 (!!)。 使用率が高くなると、リニア センシングのパフォーマンスが大幅に低下します。

このハッシュ テーブルの使用率を低く保つことが最善です。 ただし、メモリ消費を犠牲にしてパフォーマンスを向上させます。 幸いなことに、32 ビットのキーと値の場合、これは正当化できます。 上記の例で、128 億 0,25 万個の要素を含むテーブルで使用率を 32 に維持すると、そこに配置できる要素は 96 万個までとなり、残りの 8 万個のスロット (ペアごとに 768 バイト) が失われます。 , 失われた記憶はXNUMXMB。

システム メモリよりも貴重なリソースであるビデオ カード メモリの損失について話していることに注意してください。 CUDA をサポートするほとんどの最新のデスクトップ グラフィックス カードには少なくとも 4 GB のメモリが搭載されていますが (この記事の執筆時点では、NVIDIA 2080 Ti には 11 GB のメモリが搭載されています)、それでもそのような量を失うのは賢明な決定とは言えません。

後で、プローブの深さに問題がないビデオ カードのハッシュ テーブルの作成と、削除されたスロットを再利用する方法について詳しく書きます。

測深深度測定

キーのプローブの深さを決定するには、実際のテーブル インデックスからキーのハッシュ (理想的なテーブル インデックス) を抽出します。

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

1 の 3 の補数の 4 進数の魔法と、ハッシュ テーブルの容量が XNUMX の XNUMX 乗であるという事実により、このアプローチは、キー インデックスがテーブルの先頭に移動された場合でも機能します。 XNUMX にハッシュされ、スロット XNUMX に挿入されるキーを考えてみましょう。次に、容量 XNUMX のテーブルの場合、次の結果が得られます。 (3 — 1) & 3、これは 2 に相当します。

まとめ

ご質問やご意見がございましたら、次のアドレスにメールしてください。 Twitter またはで新しいトピックを開きます リポジトリ.

このコードは、優れた記事からインスピレーションを受けて作成されました。

• 世界で最もシンプルなロックフリーのハッシュ テーブル
• ロックフリー、ウェイトフリーのハッシュテーブル

今後も、ビデオ カードのハッシュ テーブルの実装について書き、そのパフォーマンスを分析していきます。 私の計画には、GPU に適したデータ構造でのアトミック操作を使用したチェーン化、ロビン フッド ハッシュ、およびカッコー ハッシュが含まれます。

出所： habr.com