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

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

Така че скоростта на лаптопа е около 300 милиона вмъквания/сек и 500 милиона изтривания/сек.

Таблицата е написана на CUDA, но същата методика може да се приложи и за HLSL или GLSL. Реализацията има няколко ограничения, осигуряващи висока производителност на видеокартата:

• Обработват се само 32-битови ключове и същите стойности.
• Хэш-таблицата има фиксиран размер.
• И този размер трябва да бъде равен на две на степен.

За ключове и стойности е необходимо да се резервира прост разделител (в предоставения код това е 0xffffffff).

Хэш-таблица без блокировки

В хэш-таблицата се използва отворена адресация с линеен зонд, тоест това е просто масив от двойки ключ-стойност, който се съхранява в паметта и има отлична производителност на кеша. Това не може да се каже за свързването в верига, което предполага търсене на указател в свързан списък. Хэш-таблицата е прост масив, който съхранява елементи 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 и с времето запълват таблицата.

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

За да илюстрирам описаните проблеми, използвам горепосочения код за създаване на таблица с 128 млн. елемента, циклично ще вмъквам 4 млн. елемента, докато не запълня 124 млн. слота (коефициент на запълване около 0,96). Ето таблица с резултати, всеки ред е извикване на CUDA ядро с вмъкване на 4 млн. нови елемента в една хеш таблица:

Коефициент на запълване
Продължителност на вмъкването на 4 194 304 елемента

0,00
11,608448 ms (361,314798 млн ключа/сек.)

0,03
11,751424 ms (356,918799 млн ключа/сек.)

0,06
11,942592 ms (351,205515 млн ключа/сек.)

0,09
12,081120 ms (347,178429 млн ключа/сек.)

0,12
12,242560 ms (342,600233 млн ключа/сек.)

0,16
12,396448 ms (338,347235 млн ключа/сек.)

0,19
12,533024 ms (334,660176 млн ключа/сек.)

0,22
12,703328 ms (330,173626 млн ключа/сек.)

0,25
12,884512 ms (325,530693 млн ключа/сек.)

0,28
13,033472 ms (321,810182 млн ключа/сек.)

0,31
13,239296 ms (316,807174 млн ключа/сек.)

0,34
13,392448 ms (313,184256 млн ключа/сек.)

0,37
13,624000 ms (307,861434 млн ключа/сек.)

0,41
13,875520 ms (302,280855 млн ключа/сек.)

0,44
14,126528 ms (296,909756 млн ключа/сек.)

0,47
14,399328 ms (291,284699 млн ключа/сек.)

0,50
14,690304 ms (285,515123 млн ключа/сек.)

0,53
15,039136 ms (278,892623 млн ключа/сек.)

0,56
15,478656 ms (270,973402 млн ключа/сек.)

0,59
15,985664 ms (262,379092 млн ключа/сек.)

0,62
16,668673 ms (251,627968 млн ключа/сек.)

0,66
17,587200 ms (238,486174 млн ключа/сек.)

0,69
18,690048 ms (224,413765 млн ключа/сек.)

0,72
20,278816 ms (206,831789 млн ключа/сек.)

0,75
22,545408 ms (186,038058 млн ключа/сек.)

0,78
26,053312 ms (160,989275 млн ключа/сек.)

0,81
31,895008 ms (131,503463 млн ключа/сек.)

0,84
42,103294 ms (99,619378 млн ключа/сек.)

0,87
61,849056 ms (67,815164 млн ключа/сек.)

0,90
105,695999 ms (39,682713 млн ключа/сек.)

0,94
240,204636 ms (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