Я выклаў на Github .
Гэта простая хэш-табліца для GPU, здольная апрацоўваць у секунду сотні мільёнаў уставак. На маім наўтбуку з NVIDIA GTX 1060 код устаўляе 64 мільёна выпадкова згенераваных пар ключ-значэнне прыкладна за 210 мс і выдаляе 32 мільёны пар прыкладна за 64 мс.
Гэта значыць хуткасць на наўтбуку складае прыкладна 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.
У артыкуле апісана, як змяняць такую структуру дадзеных, абароненую ад блакіровак.
Канкурэнтнасць
У прыведзеных вышэй фрагментах кода функцыі 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) паказальнік на масіў пар у хэш-табліцы. сцвярджае, што:
Кампілятар можа па сваім меркаванні аптымізаваць аперацыі чытання і запісы ў глабальную ці агульную памяць… Гэтыя аптымізацыі можна адключыць з дапамогай ключавога слова
volatile: … любая спасылка на гэтую зменную кампілюецца ў сапраўдную інструкцыю чытання ці запісы ў памяць.
Меркаванні карэктнасці не патрабуюць ужыванні volatile. Калі паток выканання выкарыстоўвае закэшаваныя значэнне з больш ранняй аперацыі чытання, то гэта азначае, што ён будзе выкарыстоўваць крыху састарэлую інфармацыю. Але ўсё ж гэта інфармацыя з карэктнага стану хэш-табліцы ў пэўны момант выкліку ядра. Калі вам трэба выкарыстоўваць самую свежую інфармацыю, то можна прымяняць паказальнік volatile, Але тады крыху знізіцца прадукцыйнасць: па маіх тэстах – пры выдаленні 32 элементаў хуткасць знізілася з 500 выдаленняў / сек да 450 млн выдаленняў / сек.
Proizvoditelnost
У тэсце на ўстаўку 64 элементаў і выдаленне 32 з іх канкурэнцыя паміж std::unordered_map і хэш-табліцай для GPU фактычна адсутнічае:
std::unordered_map выдаткавала 70 мс на ўстаўку і выдаленне элементаў з наступным вызваленнем unordered_map (вызваленне ад мільёнаў элементаў займае нямала часу, таму што ўнутры unordered_map выконваюцца шматлікія вылучэнні памяці). Шчыра кажучы, у std:unordered_map зусім іншыя абмежаванні. Гэта адзіны CPU-струмень выканання, ён падтрымлівае ключы-значэнні любога памеру, добра працуе пры высокіх каэфіцыентах выкарыстання і паказвае стабільную прадукцыйнасць пасля шматлікіх выдаленняў.
Працягласць працы хэш-табліцы для GPU і міжпраграмнага ўзаемадзеяння склала 984 мс. Сюды ўваходзіць час, затрачаны на размяшчэнне табліцы ў памяці і яе выдаленне (аднаразовае вылучэнне 1 Гб памяці, якое ў CUDA займае нейкі час), устаўка і выдаленне элементаў, а таксама ітэраванне па іх. Таксама ўлічаны ўсе капіяванні ў памяць і з памяці відэакарты.
Праца самой хэш-табліцы заняла 271 мс. Сюды ўваходзіць час, выдаткаваны відэакартай на ўстаўку і выдаленне элементаў, і не ўлічваецца час на капіраванне ў памяць і ітэраванне па табліцы. Калі GPU-табліца жыве доўга, ці калі хэш-табліца ўтрымоўваецца цалкам у памяці відэакарты (напрыклад, для стварэння хэш-табліцы, якая будзе выкарыстоўвацца іншым GPU-кодам, а не цэнтральным працэсарам), то вынік тэставання рэлевантны.
Хэш-табліца для відэакарты дэманструе высокую прадукцыйнасць дзякуючы вялікай прапускной здольнасці і актыўнаму распаралельванню.
Недахопы
У архітэктуры хэш-табліцы ёсць некалькі праблем, пра якія трэба памятаць:
- Лінейнаму зандзіравання перашкаджае кластарызацыя, з-за якой ключы ў табліцы размяшчаюцца далёка не ідэальна.
- Ключы не выдаляюцца з дапамогай функцыі
deleteі з часам загрувашчваюць табліцу.
У выніку прадукцыйнасць хэш-табліцы можа паступова змяншацца, асабліва калі яна існуе доўга і ў ёй выконваюцца шматлікія ўстаўкі і выдаленні. Адным са спосабаў змякчэння гэтых недахопаў з'яўляецца перахэшаванне ў новую табліцу з досыць нізкім каэфіцыентам выкарыстання і фільтраваннем выдаленых ключоў пры перахэшаванні.
Каб праілюстраваць апісаныя праблемы, выкарыстоўваю вышэйпрыведзены код для стварэння табліцы на 128 млн. элементаў, цыклічна буду ўстаўляць 4 млн. элементаў, пакуль не запаўню 124 млн. слотаў (каэфіцыент выкарыстання каля 0,96). Вось табліца вынікаў, кожны радок - гэта выклік ядра CUDA са ўстаўкай 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 Мб страчанай памяці.
Звярніце ўвагу, што гаворка ідзе аб страце памяці відэакарты, якая з'яўляецца больш каштоўным рэсурсам, чым сістэмная памяць. Хоць большасць сучасных настольных відэакартай, якія падтрымліваюць 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.
Заключэнне
Калі ў вас ёсць пытанні ці каментары, напішыце мне ў ці адкрыйце новую тэму ў.
Гэты код напісаны пад натхненнем ад выдатных артыкулаў:
У будучыні я працягну пісаць аб рэалізацыях хэш-табліц для відэакарт і буду аналізаваць іх прадукцыйнасць. У планах у мяне звязванне ў ланцужок, хэшаванне Робіна Гуда і зязюліна хэшаванне з выкарыстаннем атамарных аперацый у структурах дадзеных, якія зручныя для відэакартай.
Крыніца: habr.com