Публикувах в Github нов проект A Simple GPU Hash Table.

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

Тоест, скоростта на лаптопа е около 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 691 мс за вмъкване и изтриване на елементи с последващо освобождаване на unordered_map (освобождаването от милиони елементи отнема време, тъй като вътре в unordered_map се извършват многобройни операции по алокация на памет). Честно казано, при std:unordered_map съществуват съвсем различни ограничения. Това е единствена CPU потока на изпълнение, която поддържа ключове-стойности с произволен размер, добре работи при високи коефициенти на използване и показва стабилна производителност след многобройни изтривания.

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

Работата на самата хеш-таблица отне 271 мс. Това включва времето, което графичната карта изразходва за вставка и изтриване на елементи, и не включва времето за копиране в паметта и итериране през получената таблица. Ако GPU таблицата живее дълго или ако хеш-таблицата е изцяло в паметта на графичната карта (например за създаване на хеш-таблица, която ще се използва от друг GPU код, а не от централния процесор), то резултатът от теста е релевантен.

Хеш-таблицата за графичната карта демонстрира висока производителност благодарение на голямата пропускна способност и активно разпаралеляване.

Недостатъци

Архитектурата на хеш-таблицата има няколко проблема, които трябва да се вземат предвид:

• Линейното зондирване се затруднява от клъстериране, поради което ключовете в таблицата не са разположени идеално.
• Ключовете не се изтриват с функцията delete и с времето запълват таблицата.

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

За да илюстрирам разглежданите проблеми, използвам горепосочения код за създаване на таблица с 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 или отворете нова тема в репозитории.

Този код е написан вдъхновен от прекрасни статии:

• The World’s Simplest Lock-Free Hash Table
• A Lock-Free Wait-Free Hash Table

В бъдеще ще продължа да пиша за реализациите на хеш-таблици за видеокарти и ще анализирам тяхната производителност. В плановете ми е свързването на вериги, хеширане на Робин Худ и кукушко хеширане с използване на атомарни операции в структури от данни, удобни за видеокарти.

Източник: habr.com