Я виклав на Github новий проект A Simple GPU Hash Table.

Це проста хеш-таблиця для GPU, здатна обробляти за секунду сотні мільйонів вставок. На моєму ноутбуці з NVIDIA GTX 1060 код вставляє 64 мільйони випадково згенерованих пар ключ-значення приблизно за 210 мс і видаляє 32 мільйони пар приблизно за 64 мс.

Тобто швидкість на ноутбуці становить приблизно 300 млн вставок/сек та 500 млн видалень/сек.

Таблиця написана на CUDA, хоча ту ж методику можна застосувати до HLSL чи GLSL. У реалізації є кілька обмежень, що забезпечують високу продуктивність відеокарти:

Обробляються тільки 32-бітові ключі та такі ж значення.
Хеш-таблиця має фіксований розмір.
І цей розмір повинен дорівнювати двом у ступеня.

Для ключів і значень потрібно зарезервувати простий маркер, що розмежовує (у наведеному коді це 0xffffffff).

Хеш-таблиця без блокувань

У хеш-таблиці використовується відкрита адресація з лінійним зондуваннямЦе просто масив пар ключ-значення, який зберігається в пам'яті і має чудову продуктивність кешу. Цього не скажеш про зв'язування в ланцюжок (chaining), що передбачає пошук покажчика у списку. Хеш-таблиця є простим масивом, що зберігає елементи KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Розмір таблиці дорівнює двійці в мірі, а не простому числу, тому що для застосування pow2/AND-маски достатньо однієї швидкої інструкції, а оператор модуля працює набагато повільніше. Це важливо у разі лінійного зондування, оскільки при лінійному пошуку по таблиці індекс слота має бути обернутий у кожен слот. І в результаті додається вартість операції за модулем у кожному слоті.

Таблиця зберігає лише ключ та значення для кожного елемента, а не хеш ключа. Оскільки таблиця зберігає лише 32-бітові ключі, хеш обчислюється дуже швидко. У наведеному коді використовується хеш Murmur3, який виконує лише кілька зрушень, XOR-ів та множень.

У хеш-таблиці застосовується методика захисту від блокувань, яка залежить від порядку розміщення у пам'яті. Навіть якщо деякі операції запису порушують черговість таких операцій, хеш-таблиця все одно збереже коректний стан. Про це ми поговоримо нижче. Методика чудово працює з відеокартами, в яких конкурентно виконуються тисячі потоків.

Ключі та значення у хеш-таблиці ініціалізуються порожніми.

Код можна модифікувати, щоб він міг обробляти і 64-бітові ключі та значення. Для ключів потрібні атомарні операції читання, запису та порівняння з обміном (compare-and-swap). А для значень потрібні атомарні операції читання та запису. На щастя, у CUDA операції читання-запису для 32- і 64-бітових значень є атомарними доти, доки вони вирівняні природним чином (див. тут), а сучасні відеокарти підтримують 64-бітові атомарні операції порівняння з обміном. Звичайно, при переході на 64 біта продуктивність дещо знизиться.

Стан хеш-таблиці

Кожна пара ключ-значення в хеш-таблиці може мати один із чотирьох станів:

Ключ та значення порожні. У такому стані хеш-таблиця ініціалізується.
Ключ був записаний, а значення ще немає. Якщо інший потік виконання у цей момент зчитує дані, потім він повертає порожнє значення. Це нормально, те саме сталося б, якби інший потік виконання відпрацював трохи раніше, а ми говоримо про конкурентну структуру даних.
Записано і ключ, і значення.
Значення доступне для інших потоків виконання, а ключ ще немає. Таке може статися, тому що модель програмування в CUDA має на увазі слабко впорядковану модель пам'яті. Це нормально, за будь-якої події ключ все ще порожній, навіть якщо значення вже не є.

Важливий нюанс полягає в тому, що як тільки ключ був записаний у слот, він уже не переміщається — навіть якщо ключ буде видалений, про це ми поговоримо нижче.

Код хеш-таблиці працює навіть із слабо впорядкованими моделями пам'яті, в яких не відомий порядок читання та записів у пам'ять. Коли ми розбиратимемо вставку, пошук і видалення в хеш-таблиці, пам'ятайте, що кожна пара ключ-значення знаходиться в одному з чотирьох вищеописаних станів.

Вставка в хеш-таблицю

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);
    }
}

Для вставки ключа код ітерує по масиву хеш-таблиці починаючи з хеша ключа, що вставляється. У кожному слоті масиву виконується атомарна операція порівняння з обміном, коли він ключ у цьому слоті порівнюється з порожнім. Якщо виявляється невідповідність, то ключ у слоті оновлюється на ключ, що вставляється, а потім повертається оригінальний ключ слота. Якщо цей оригінальний ключ був порожнім або відповідав ключу, що вставляється, значить код знайшов відповідний для вставки слот і вносить в слот значення, що вставляється.

Якщо в одному виклику ядра gpu_hashtable_insert() є кілька елементів з однаковим ключем, тоді будь-яке з значень може бути записано в слот ключа. Це вважається нормальним: одна з операцій запису ключа-значення в ході виклику буде успішною, але оскільки все це відбувається паралельно в рамках кількох потоків виконання, ми не можемо передбачити, яка операція запису в пам'ять буде останньою.

Пошук у хеш-таблиці

Код для пошуку ключів:

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);
        }
}

Щоб знайти значення ключа, що зберігається в таблиці, ми ітеруємо по масиву починаючи з хеша ключа, що шукається. У кожному слоті ми перевіряємо, чи є ключ тим, який шукаємо, і якщо так, то повертаємо його значення. Також ми перевіряємо, чи є ключ порожнім, і якщо так, то перериваємо пошук.

Якщо нам не вдається знайти ключ, код повертає порожнє значення.

Всі ці операції пошуку можуть виконуватися конкурентно під час вставок та видалень. Кожна пара в таблиці матиме для потоку один із чотирьох вищеописаних станів.

Видалення у хеш-таблиці

Код видалення ключів:

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.

У статті A Lock-Free Wait-Free Hash Table описано як змінювати таку структуру даних, захищену від блокувань.

Конкурентність

У наведених вище фрагментах коду функції gpu_hashtable_insert(), _lookup() и _delete() обробляють по одній парі ключ-значення за один раз. А нижче 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);
    }
}

Хеш-таблиця із захистом від блокувань підтримує конкурентні вставки, пошуки та видалення. Оскільки пари ключ-значення завжди перебувають у одному з чотирьох станів, а ключі не переміщуються, таблиця гарантує коректність навіть за одночасному використанні операцій різних видів.

Однак якщо ми паралельно обробляємо пакет зі вставок і видалень, і якщо у вхідному масиві пар містяться ключі, що дублюються, то ми не зможемо передбачити, які пари «переможуть» — будуть записані в хеш-таблицю останніми. Допустимо, ми викликали код вставки з вхідним масивом з пар A/0 B/1 A/2 C/3 A/4. Коли код завершиться, пари B/1 и C/3 гарантовано будуть присутні у таблиці, але при цьому в ній виявиться будь-яка з пар A/0, A/2 або A/4. Це може бути проблемою, а може й не бути все залежить від застосування. Ви можете заздалегідь знати, що у вхідному масиві немає ключів, що дублюються, або вам може бути не важливо, яке значення було записано останнім.

Якщо для вас це проблема, то потрібно розділити пари, що дублюються, по різних системних CUDA-викликах. У CUDA будь-яка операція із викликом ядра завжди завершується до наступного виклику ядра (принаймні, всередині одного потоку. У різних потоках ядра виконуються паралельно). Якщо у наведеному вище прикладі викликати одне ядро з A/0 B/1 A/2 C/3, а інше з A/4тоді ключ A отримає значення 4.

Тепер поговоримо про те, чи потрібні функції lookup() и delete() використовувати простий (plain) або змінний (volatile) покажчик на масив пар у хеш-таблиці. Документація CUDA стверджує, що:

Компілятор може на свій розсуд оптимізувати операції читання та запису в глобальну або загальну пам'ять … Ці оптимізації можна вимкнути за допомогою ключового слова volatile: … будь-яке посилання на цю змінну компілюється в інструкцію читання або запису в пам'ять.

Міркування коректності не вимагають застосування volatile. Якщо потік виконання використовує закешоване значення більш ранньої операції читання, це означає, що він використовуватиме трохи застарілу інформацію. Але все ж таки це інформація з коректного стану хеш-таблиці в певний момент виклику ядра. Якщо вам потрібно використовувати найсвіжішу інформацію, то можна застосовувати покажчик volatileАле тоді трохи знизиться продуктивність: за моїми тестами - при видаленні 32 млн елементів швидкість знизилася з 500 млн видалень/сек до 450 млн видалень/сек.

Продуктивність

У тесті на вставку 64 млн елементів і видалення 32 млн з них конкуренція між std::unordered_map та хеш-таблицею для GPU фактично відсутня:

std::unordered_map витратила 70 мс на вставку і видалення елементів з подальшим звільненням unordered_map (Звільнення від мільйонів елементів займає чимало часу, тому що всередині unordered_map виконуються численні виділення пам'яті). Чесно кажучи, у std:unordered_map зовсім інші обмеження. Це єдиний CPU-потік виконання, він підтримує ключі-значення будь-якого розміру, добре працює при високих коефіцієнтах використання та показує стабільну продуктивність після численних видалень.

Тривалість роботи хеш-таблиці для GPU та міжпрограмної взаємодії склала 984 мс. Сюди входить час, витрачений на розміщення таблиці в пам'яті та її видалення (одноразове виділення 1 Гб пам'яті, яке CUDA займає якийсь час), вставка і видалення елементів, а також ітерування по них. Також враховано всі копіювання в пам'ять та з пам'яті відеокарти.

Робота самої хеш-таблиці зайняла 271 мс. Сюди входить час, витрачений відеокартою на вставку і видалення елементів, і не враховується час на копіювання в пам'ять та ітерування таблиці, що вийшла. Якщо GPU-таблиця живе довго, або якщо хеш-таблиця міститься цілком у пам'яті відеокарти (наприклад, для створення хеш-таблиці, яка використовуватиметься іншим GPU-кодом, а не центральним процесором), то результат тестування є релевантним.

Хеш-таблиця для відеокарти демонструє високу продуктивність завдяки великій пропускній здатності та активному розпаралелювання.

Недоліки

У архітектури хеш-таблиці є кілька проблем, про які потрібно пам'ятати:

Лінійному зондуванню заважає кластеризація, через яку ключі у таблиці розміщуються далеко не ідеально.
Ключі не видаляються за допомогою функції delete і згодом захаращують таблицю.

В результаті продуктивність хеш-таблиці може поступово знижуватися, особливо якщо вона існує довго і в ній виконуються численні вставки та видалення. Одним із способів пом'якшення цих недоліків є перехешування в нову таблицю з досить низьким коефіцієнтом використання та фільтрацією віддалених ключів при перехешуванні.

Щоб проілюструвати описані проблеми, використовую наведений вище код для створення таблиці на 128 млн елементів, циклічно вставлятиму 4 млн елементів, поки не заповню 124 млн слотів (коефіцієнт використання близько 0,96). Ось таблиця результатів, кожен рядок – це виклик ядра CUDA із вставкою 4 млн нових елементів в одну хеш-таблицю:

Коефіцієнт використання
Тривалість вставки 4 194 304 елементів

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.

Потім я виміряв глибину зондування в таблиці зі 124 млн вставок (коефіцієнт використання 0,97). Середня глибина становила вже 10,1757, а максимальна - 6474 (!!). Продуктивність лінійного зондування сильно знижується при великих коефіцієнтах використання.

Найкраще зберігати у цій хеш-таблиці низький коефіцієнт використання. Але тоді ми підвищуємо продуктивність рахунок споживання пам'яті. На щастя, у випадку з 32-бітними ключами та значеннями, і це може бути виправдано. Якщо в наведеному вище прикладі в таблиці на 128 млн. елементів зберігати коефіцієнт використання 0,25, то ми зможемо розмістити в ній не більше 32 млн. елементів, а решту 96 млн. слотів будуть втрачені — по 8 байтів на кожну пару, 768 Мб втраченої пам'яті.

Зверніть увагу, що йдеться про втрату пам'яті відеокарти, яка є ціннішим ресурсом, ніж системна пам'ять. Хоча більшість сучасних настільних відеокарт, що підтримують CUDA, мають не менше 4 Гб пам'яті (на момент написання статті у NVIDIA 2080 Ti є 11 Гб), все ж таки втрачати такі обсяги буде не наймудрішим рішенням.

Пізніше я докладніше напишу про створення хеш-таблиць для відеокарт, які не мають проблем із глибиною зондування, а також про способи повторного використання віддалених слотів.

Вимірювання глибини зондування

Щоб визначити глибину зондування ключа, ми можемо витягти хеш ключа (його ідеальний індекс у таблиці) з фактичного табличного індексу:

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

Через магію двох двійкових чисел у додатковому коді і того факту, що ємність хеш-таблиці дорівнює двійці в міру, цей підхід працюватиме навіть тоді, коли індекс ключа переноситься на початок таблиці. Візьмемо ключ, який хешується в 1, але вставлений у слот 3. Тоді для таблиці з ємністю 4 ми отримаємо (3 — 1) & 3, Що еквівалентно 2.

Висновок

Якщо у вас є питання або коментарі, напишіть мені Twitter або відкрийте нову тему в репозиторії.

Цей код написаний під натхненням від прекрасних статей:

У майбутньому я продовжу писати про реалізації хеш-таблиць для відеокарт і аналізуватиму їх продуктивність. У планах у мене зв'язування в ланцюжок, хешування Робіна Гуда і зозуліне хешування з використанням атомарних операцій у структурах даних, які зручні для відеокарт.

Джерело: habr.com

Проста хеш-таблиця для GPU