Поставио сам га на Гитхуб .
То је једноставна ГПУ хеш табела способна да обради стотине милиона уметања у секунди. На мом НВИДИА ГТКС 1060 лаптопу, код убацује 64 милиона насумично генерисаних парова кључ-вредност за око 210 мс и уклања 32 милиона парова за око 64 мс.
То јест, брзина на лаптопу је приближно 300 милиона уметања у секунди и 500 милиона брисања у секунди.
Табела је написана у ЦУДА-и, иако се иста техника може применити на ХЛСЛ или ГЛСЛ. Имплементација има неколико ограничења како би се осигурале високе перформансе на видео картици:
- Обрађују се само 32-битни кључеви и исте вредности.
- Хеш табела има фиксну величину.
- И ова величина мора бити једнака два на степен.
За кључеве и вредности, потребно је да резервишете једноставан граничник (у горњем коду ово је 0кфффффффф).
Хасх табела без брава
Хеш табела користи отворено адресирање са , то јест, то је једноставно низ парова кључ-вредност који се чува у меморији и има супериорне перформансе кеша. Исто се не може рећи за уланчавање, које укључује тражење показивача на повезаној листи. Хеш табела је једноставан низ за складиштење елемената KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Величина табеле је степен два, а не прост број, јер је једна брза инструкција довољна за примену пов2/АНД маске, али је оператор модула много спорији. Ово је важно у случају линеарног сондирања, пошто у линеарном тражењу табеле индекс слота мора бити умотан у сваки слот. И као резултат тога, цена операције се додаје модулом у сваки слот.
Табела чува само кључ и вредност за сваки елемент, а не хеш кључа. Пошто табела чува само 32-битне кључеве, хеш се израчунава веома брзо. Код изнад користи Мурмур3 хеш, који изводи само неколико померања, КСОР-ова и множења.
Хеш табела користи технике заштите закључавања које су независне од редоследа меморије. Чак и ако неке операције писања поремете редослед других таквих операција, хеш табела ће и даље одржавати исправно стање. О овоме ћемо причати у наставку. Техника одлично функционише са видео картицама које истовремено покрећу хиљаде нити.
Кључеви и вредности у хеш табели су иницијализовани на празне.
Код се такође може модификовати за руковање 64-битним кључевима и вредностима. Кључеви захтевају атомске операције читања, писања и поређења и замене. А вредности захтевају атомске операције читања и писања. На срећу, у ЦУДА-и, операције читања и писања за 32- и 64-битне вредности су атомске све док су природно усклађене (погледајте доле). ), а модерне видео картице подржавају 64-битне атомске операције упоређивања и размене. Наравно, када се пређе на 64 бита, перформансе ће се мало смањити.
Стање хеш табеле
Сваки пар кључ/вредност у хеш табели може имати једно од четири стања:
- Кључ и вредност су празни. У овом стању, хеш табела је иницијализована.
- Кључ је записан, али вредност још није записана. Ако друга нит тренутно чита податке, онда се враћа празно. Ово је нормално, иста ствар би се десила да је друга нит извршења прорадила мало раније, а говоримо о конкурентној структури података.
- И кључ и вредност се снимају.
- Вредност је доступна другим нитима извршења, али кључ још није. Ово се може десити зато што ЦУДА модел програмирања има лабаво уређен меморијски модел. Ово је нормално; у сваком случају, кључ је и даље празан, чак и ако вредност више није тачна.
Важна нијанса је да када је кључ уписан у слот, он се више не помера - чак и ако је кључ избрисан, о томе ћемо говорити у наставку.
Код хеш табеле функционише чак и са лабаво уређеним меморијским моделима у којима је непознат редослед читања и уписивања меморије. Док гледамо уметање, тражење и брисање у хеш табели, запамтите да је сваки пар кључ-вредност у једном од четири горе описана стања.
Уметање у хеш табелу
Функција ЦУДА која убацује парове кључ/вредност у хеш табелу изгледа овако:
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(), осим што када се нађе подударање на кључу, чини његову вредност празном.
Као што је горе поменуто, када се кључ упише у слот, он се више не помера. Чак и када се елемент избрише из табеле, кључ остаје на месту, његова вредност једноставно постаје празна. То значи да не треба да користимо атомску операцију писања за вредност слота, јер није битно да ли је тренутна вредност празна или не – она ће и даље постати празна.
Промена величине хеш табеле
Можете променити величину хеш табеле тако што ћете креирати већу табелу и у њу убацити непразне елементе из старе табеле. Нисам имплементирао ову функционалност јер сам желео да узорак кода буде једноставан. Штавише, у ЦУДА програмима, додела меморије се често врши у коду домаћина, а не у ЦУДА кернелу.
У чланку описује како модификовати такву структуру података заштићену закључавањем.
Конкурентност
У горњим исечцима кода функције 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. Ово може или не мора бити проблем - све зависи од апликације. Можда знате унапред да у низу уноса нема дуплих кључева или вам можда није важно која је вредност последња уписана.
Ако вам ово представља проблем, онда морате да раздвојите дуплиране парове у различите ЦУДА системске позиве. У ЦУДА-и, свака операција која позива кернел се увек завршава пре следећег позива кернела (барем унутар једне нити. У различитим нитима, језгра се извршавају паралелно). У горњем примеру, ако позовете једно језгро са A/0 B/1 A/2 C/3, а други са A/4, затим кључ A добиће вредност 4.
Хајде сада да разговарамо о томе да ли функције треба lookup() и delete() користите обичан или променљив показивач на низ парова у хеш табели. Наводи да:
Преводилац може изабрати да оптимизује читање и уписивање у глобалну или заједничку меморију... Ове оптимизације се могу онемогућити помоћу кључне речи
volatile: ... свака референца на ову променљиву се компајлира у праву инструкцију за читање или писање меморије.
Разматрања исправности не захтевају примену volatile. Ако извршна нит користи кеширану вредност из раније операције читања, тада ће користити мало застареле информације. Али ипак, ово је информација из исправног стања хеш табеле у одређеном тренутку позива кернела. Ако желите да користите најновије информације, можете користити индекс volatile, али онда ће се перформансе мало смањити: према мојим тестовима, при брисању 32 милиона елемената брзина је смањена са 500 милиона брисања/сек на 450 милиона брисања/сек.
Перформансе
У тесту за убацивање 64 милиона елемената и брисање њих 32 милиона, конкуренција између std::unordered_map и практично не постоји хеш табела за ГПУ:
std::unordered_map провео 70 мс уметајући и уклањајући елементе и затим их ослобађајући unordered_map (да бисте се ослободили милиона елемената, потребно је много времена, јер унутра unordered_map врши се вишеструка додела меморије). Искрено говорећи, std:unordered_map сасвим друга ограничења. То је једна ЦПУ нит извршења, подржава кључ/вредност било које величине, добро ради при високим стопама коришћења и показује стабилне перформансе након вишеструких брисања.
Трајање хеш табеле за ГПУ и међупрограмску комуникацију је било 984 мс. Ово укључује време утрошено на стављање табеле у меморију и њено брисање (додељивање 1 ГБ меморије једном, за шта је потребно неко време у ЦУДА-и), уметање и брисање елемената и понављање преко њих. Такође се узимају у обзир све копије у и из меморије видео картице.
Сама хеш табела се завршила 271 мс. Ово укључује време које је видео картица потрошила на уметање и брисање елемената и не узима у обзир време утрошено на копирање у меморију и понављање преко резултујуће табеле. Ако ГПУ табела живи дуго времена или је хеш табела у потпуности садржана у меморији видео картице (на пример, да би се направила хеш табела коју ће користити други ГПУ код, а не централни процесор), онда резултат теста је релевантан.
Хеш табела за видео картицу показује високе перформансе због велике пропусности и активне паралелизације.
Ограничења
Архитектура хеш табеле има неколико проблема којих треба да будете свесни:
- Линеарно испитивање отежава груписање, што узрокује да кључеви у табели буду постављени мање него савршено.
- Тастери се не уклањају коришћењем функције
deleteа временом затрпају сто.
Као резултат тога, перформансе хеш табеле могу постепено да деградирају, посебно ако постоји дуже време и има бројна уметања и брисања. Један од начина да се ублаже ови недостаци је да се поново хаширају у нову табелу са прилично ниском стопом искоришћења и филтрирају уклоњени кључеви током рехаширања.
Да бих илустровао описане проблеме, користићу горњи код да направим табелу са 128 милиона елемената и прођем кроз 4 милиона елемената док не попуним 124 милиона слотова (стопа искоришћења од око 0,96). Ево табеле резултата, сваки ред је позив ЦУДА кернела за уметање 4 милиона нових елемената у једну хеш табелу:
Стопа кориштења
Трајање уметања 4 елемента
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 МБ изгубљене меморије.
Имајте на уму да говоримо о губитку меморије видео картице, која је вреднији ресурс од системске меморије. Иако већина модерних десктоп графичких картица које подржавају ЦУДА имају најмање 4 ГБ меморије (у тренутку писања НВИДИА 2080 Ти има 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.
Закључак
Ако имате питања или коментара, пошаљите ми е-пошту на или отворите нову тему.
Овај код је написан инспирисан одличним чланцима:
У будућности ћу наставити да пишем о имплементацији хеш табеле за видео картице и анализирам њихове перформансе. Моји планови укључују уланчавање, хеширање Робина Худа и хеширање кукавице користећи атомске операције у структурама података које су прилагођене ГПУ-у.
Извор: ввв.хабр.цом