Публикувах го в Github нов проект Проста хеш таблица на GPU.

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

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

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

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

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

Хеш маса без брави

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

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Размерът на таблицата е степен на две, а не просто число, защото една бърза инструкция е достатъчна, за да се приложи маската pow2/AND, но модулният оператор е много по-бавен. Това е важно в случай на линейно сондиране, тъй като при търсене в линейна таблица индексът на слота трябва да бъде обвит във всеки слот. И в резултат цената на операцията се добавя по модул във всеки слот.

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

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

Ключовете и стойностите в хеш-таблицата се инициализират за празни.

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

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

Конкурентоспособност

В горните функционални кодови фрагменти 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() използвайте обикновен или непостоянен указател към масив от двойки в хеш-таблицата. CUDA документация Гласи че:

Компилаторът може да избере да оптимизира четенията и записите в глобална или споделена памет... Тези оптимизации могат да бъдат деактивирани чрез ключовата дума volatile: ... всяка препратка към тази променлива се компилира в инструкция за четене или запис в реална памет.

Съображенията за коректност не изискват прилагане volatile. Ако нишката за изпълнение използва кеширана стойност от по-ранна операция за четене, тогава тя ще използва малко остаряла информация. Но все пак това е информация от правилното състояние на хеш-таблицата в определен момент от извикването на ядрото. Ако трябва да използвате най-новата информация, можете да използвате индекса volatile, но тогава производителността ще намалее леко: според моите тестове, при изтриване на 32 милиона елемента, скоростта намаля от 500 милиона изтривания/сек на 450 милиона изтривания/сек.

продуктивност

В теста за вмъкване на 64 милиона елемента и изтриване на 32 милиона от тях, конкуренцията между std::unordered_map и на практика няма хеш таблица за GPU:

std::unordered_map изразходва 70 691 ms за вмъкване и премахване на елементи и след това за освобождаването им unordered_map (отърваването от милиони елементи отнема много време, защото вътре unordered_map правят се множество разпределения на памет). честно казано, std:unordered_map напълно различни ограничения. Това е нишка за изпълнение с един процесор, поддържа ключ-стойности от всякакъв размер, представя се добре при високи нива на използване и показва стабилна производителност след множество изтривания.

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

Завършването на самата хеш таблица отне 271 ms. Това включва времето, прекарано от видеокартата за вмъкване и изтриване на елементи, и не взема предвид времето, прекарано за копиране в паметта и повторение на получената таблица. Ако 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 MB загубена памет.

Имайте предвид, че говорим за загуба на памет на видеокартата, която е по-ценен ресурс от системната памет. Въпреки че повечето съвременни настолни графични карти, които поддържат 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 получаваме (3 — 1) & 3, което е еквивалентно на 2.

Заключение

Ако имате въпроси или коментари, моля, пишете ми на Twitter или отворете нова тема в хранилища.

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

• Най-простата хеш таблица без заключване в света
• Хеш таблица без заключване и изчакване

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

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