Го објавив на Github нов проект Едноставна хеш табела на графичкиот процесор.

Тоа е едноставна табела за хаш на графичкиот процесор способна да обработува стотици милиони инсерти во секунда. На мојот лаптоп 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() обработувајте низа од парови паралелно, секој пар во посебна нишка за извршување на графичкиот процесор:

// 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 и практично нема хеш-табела за графичкиот процесор:

std::unordered_map потрошил 70 ms вметнувајќи и отстранувајќи елементи и потоа ослободувајќи ги unordered_map (да се ослободиме од милиони елементи потребно е многу време, бидејќи внатре unordered_map се прават повеќе мемориски распределби). Искрено кажано, std:unordered_map сосема различни ограничувања. Тоа е една нишка за извршување на процесорот, поддржува клучни вредности од која било големина, добро функционира при високи стапки на искористеност и покажува стабилни перформанси по повеќекратни бришења.

Времетраењето на хаш-табелата за графичкиот процесор и комуникацијата меѓу програмите беше 984 ms. Ова го вклучува времето поминато за поставување на табелата во меморијата и нејзино бришење (одвојување 1 GB меморија еднократно, за што е потребно извесно време во CUDA), вметнување и бришење елементи и повторување над нив. Се земаат предвид и сите копии до и од меморијата на видео картичката.

На самата хаш-табела и беа потребни 271 ms за да се комплетира. Ова го вклучува времето поминато од видео картичката за вметнување и бришење елементи и не го зема предвид времето поминато за копирање во меморијата и повторување преку добиената табела. Ако табелата на графичкиот процесор живее долго време, или ако табелата за хаш е целосно содржана во меморијата на видео картичката (на пример, да се создаде хеш-табела што ќе се користи од друг код на графичкиот процесор, а не од централниот процесор), тогаш резултатот од тестот е релевантен.

Табелата за хаш за видео картичка покажува високи перформанси поради високата пропусност и активната паралелизација.

Ограничувања

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

• Линеарното сондирање е попречено од групирањето, што предизвикува клучевите во табелата да бидат поставени помалку од совршено.
• Копчињата не се отстрануваат со помош на функцијата delete и со текот на времето ја натрупуваат масата.

Како резултат на тоа, перформансите на хеш-табела може постепено да се влошуваат, особено ако постои долго време и има бројни вметнувања и бришења. Еден начин да се ублажат овие недостатоци е да се преиспита во нова табела со прилично ниска стапка на искористеност и да се филтрираат отстранетите клучеви за време на повторното пробивање.

За да ги илустрирам опишаните проблеми, ќе го користам горенаведениот код за да создадам табела со 128 милиони елементи и да вртам низ 4 милиони елементи додека не пополнам 124 милиони слотови (стапка на искористеност од околу 0,96). Еве ја табелата со резултати, секој ред е повик на CUDA кернелот за вметнување 4 милиони нови елементи во една хаш табела:

Стапка на користење
Времетраење на вметнување 4 елементи

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 или отворете нова тема во складишта.

Овој код е напишан под инспирација од одлични статии:

• Наједноставната хеш табела во светот без заклучување
• Хеш табела без чекање без заклучување

Во иднина, ќе продолжам да пишувам за имплементации на хеш-табели за видео картички и да ги анализирам нивните перформанси. Моите планови вклучуваат поврзување со синџири, хеширање на Робин Худ и хаширање со кукавици со користење на атомски операции во структури на податоци кои се пријателски за графичкиот процесор.

Извор: www.habr.com