Мен оны Github сайтында жарияладым .
Бұл секундына жүздеген миллион кірістірулерді өңдеуге қабілетті қарапайым GPU хэш кестесі. Менің NVIDIA GTX 1060 ноутбугімде код шамамен 64 мс-те кездейсоқ құрылған 210 миллион кілт-мән жұбын енгізеді және шамамен 32 мс-те 64 миллион жұпты жояды.
Яғни, ноутбуктағы жылдамдық шамамен 300 миллион кірістіру/сек және 500 миллион өшіру/сек.
Кесте CUDA тілінде жазылған, дегенмен бірдей әдісті HLSL немесе GLSL үшін қолдануға болады. Іске асыруда бейне картада жоғары өнімділікті қамтамасыз ету үшін бірнеше шектеулер бар:
- Тек 32 биттік кілттер және бірдей мәндер өңделеді.
- Хэш кестесінің белгіленген өлшемі бар.
- Және бұл өлшем қуатқа екіге тең болуы керек.
Кілттер мен мәндер үшін қарапайым бөлгіш маркерді сақтау қажет (жоғарыдағы кодта бұл 0xffffffff).
Құлыптары жоқ хэш кестесі
Хэш кестесі ашық адрестеуді пайдаланады , яғни бұл жадта сақталатын және жоғары кэш өнімділігі бар кілт-мән жұптарының массиві. Байланыстырылған тізімдегі көрсеткішті іздеуді қамтитын тізбектеу үшін де дәл осылай айту мүмкін емес. Хэш-кесте элементтерді сақтайтын қарапайым массив KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Кесте өлшемі жай санның емес, екінің дәрежесі болып табылады, өйткені pow2/AND маскасын қолдану үшін бір жылдам нұсқау жеткілікті, бірақ модуль операторы әлдеқайда баяу. Бұл сызықтық зондтау жағдайында маңызды, өйткені сызықтық кестені іздеуде ұяшық индексі әрбір ұяшыққа оралуы керек. Нәтижесінде операция құны әрбір ұяшыққа модуль қосылады.
Кесте кілттің хэшін емес, әрбір элемент үшін кілт пен мәнді ғана сақтайды. Кесте тек 32-биттік кілттерді сақтайтындықтан, хэш өте жылдам есептеледі. Жоғарыдағы код бірнеше ауысымды, XOR және көбейтуді ғана орындайтын Murmur3 хэшін пайдаланады.
Хэш кестесі жад тәртібіне тәуелсіз құлыптан қорғау әдістерін пайдаланады. Кейбір жазу операциялары басқа осындай операциялардың тәртібін бұзса да, хэш кестесі әлі де дұрыс күйді сақтайды. Бұл туралы төменде айтатын боламыз. Техника мыңдаған ағындарды бір уақытта іске қосатын видеокарталармен тамаша жұмыс істейді.
Хэш кестесіндегі кілттер мен мәндер бос күйге инициализацияланады.
Кодты 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() хэш кестесіндегі жұптар массивіне қарапайым немесе өзгермелі көрсеткішті пайдаланыңыз. мынаны мәлімдейді:
Компилятор жаһандық немесе ортақ жадқа оқу мен жазуды оңтайландыруды таңдауы мүмкін... Бұл оңтайландыруларды кілт сөзді пайдаланып өшіруге болады.
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 коды пайдаланатын хэш кестесін жасау үшін), онда сынақ нәтижесі өзекті болып табылады.
Бейнекартаға арналған хэш-кесте жоғары өткізу қабілеті мен белсенді параллелизацияның арқасында жоғары өнімділікті көрсетеді.
кемшіліктер
Хэш кестесінің архитектурасында білуге болатын бірнеше мәселелер бар:
- Сызықтық зондтауға кластерлеу кедергі келтіреді, бұл кестедегі кілттердің мінсіз орналасуына әкеліп соғады.
- Функция арқылы пернелер жойылмайды
deleteжәне уақыт өте келе олар үстелді шатастырады.
Нәтижесінде хэш кестесінің өнімділігі, әсіресе ол ұзақ уақыт бойы бар болса және көптеген кірістірулер мен жоюлар болса, бірте-бірте нашарлауы мүмкін. Бұл кемшіліктерді азайтудың бір жолы - пайдалану жылдамдығы өте төмен жаңа кестеге қайта қосу және қайта өңдеу кезінде жойылған кілттерді сүзгілеу.
Сипатталған мәселелерді суреттеу үшін мен 128 миллион элементі бар кестені құру және 4 миллион слотты толтырғанша (пайдалану көрсеткіші шамамен 124) 0,96 миллион элементті айналдыру үшін жоғарыдағы кодты қолданамын. Міне, нәтижелер кестесі, әрбір жол бір хэш кестесіне 4 миллион жаңа элементтерді кірістіру үшін CUDA ядросының шақыруы:
Пайдалану жылдамдығы
Енгізу ұзақтығы 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-ге тең.
қорытынды
Сұрақтарыңыз немесе пікірлеріңіз болса, маған электрондық пошта арқылы жіберіңіз немесе жаңа тақырыпты ашыңыз.
Бұл код тамаша мақалалардан шабытпен жазылған:
Болашақта мен видеокарталарға арналған хэш-кестелерді енгізу туралы жазуды және олардың өнімділігін талдауды жалғастырамын. Менің жоспарларыма GPU қолайлы деректер құрылымдарында атомдық операцияларды пайдалана отырып тізбектеу, Робин Гуд хэшинг және көкек хэшинг кіреді.
Ақпарат көзі: www.habr.com