Качих на 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);
    }
}

Хеш-таблица с защита от блокировки поддържа конкурентни вмъквания, търсения и изтривания. Тъй като двойките ключ-стойност винаги са в едно от четирите състояния и ключовете не се преместват, таблицата гарантира коректност дори при simultaneus операции от различен вид.

Ако обаче паралелно обработваме пакет от вмъквания и изтривания, и ако входният масив от двойки съдържа дублиращи се ключове, няма да можем да предвидим кои двойки "побеждават" — кои ще бъдат записани в хеш-таблицата последни. Да предположим, че сме извикали кода за вмъкване с входен масив от двойки 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 код, а не от централния процесор), то резултатът от теста е актуален.

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

Недостатъци

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

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

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

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

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

• Най-простата безключова хеш-таблица в света
• A Lock-Free Wait-Free Hash Table

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

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