Мен аны Githubга жарыяладым .
Бул секундасына жүз миллиондогон кошумчаларды иштетүүгө жөндөмдүү жөнөкөй GPU хэш таблицасы. Менин NVIDIA GTX 1060 ноутбугумда код 64 мс аралыгында 210 миллион кокусунан түзүлгөн ачкыч-маани жуптарын киргизет жана болжол менен 32 мс ичинде 64 миллион жупту жок кылат.
Башкача айтканда, ноутбуктун ылдамдыгы болжол менен 300 миллион киргизүү/сек жана 500 миллион өчүрүү/сек.
Таблица CUDAда жазылган, бирок ошол эле ыкманы HLSL же GLSL үчүн колдонсо болот. ишке ашыруу Video картада жогорку аткарууну камсыз кылуу үчүн бир нече чектөөлөр бар:
- 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() хэш таблицасындагы жуптар массивине жөнөкөй же туруксуз көрсөткүчтү колдонуңуз. деп айтылат:
Компилятор глобалдык же жалпы эстутумда окууларды жана жазууларды оптималдаштырууну тандашы мүмкүн... Бул оптималдаштырууларды ачкыч сөз менен өчүрсө болот.
volatile: ... бул өзгөрмөнүн каалаган шилтемеси реалдуу эстутумда окуу же жазуу инструкциясына түзүлөт.
Тууралыкты эске алуу өтүнмөнү талап кылбайт volatile. Эгерде аткаруу жиптери мурунку окуу операциясынан кэштелген маанини колдонсо, анда ал бир аз эскирген маалыматты колдонот. Ошентсе да, бул ядро чакырыгынын белгилүү бир учурда хэш таблицасынын туура абалынан алынган маалымат. Сиз акыркы маалыматты колдонуу керек болсо, индексти колдоно аласыз volatile, бирок анда өндүрүмдүүлүк бир аз төмөндөйт: менин тесттериме ылайык, 32 миллион элементти жок кылууда ылдамдык 500 миллион өчүрүү/сектен 450 миллион өчүрүү/сек чейин төмөндөдү.
кирешелүүлүк
Сынакта 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). Бул жерде жыйынтык таблицасы, ар бир сап бир хэш таблицага 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 менен ыңгайлуу болгон маалымат структураларында атомдук операцияларды колдонуу менен күкүк хэшинги кирет.
Source: www.habr.com