Jeg postede det på Github .
Det er en simpel GPU-hash-tabel, der er i stand til at behandle hundredvis af millioner af inserts i sekundet. På min NVIDIA GTX 1060 bærbare computer indsætter koden 64 millioner tilfældigt genererede nøgleværdi-par på omkring 210 ms og fjerner 32 millioner par på omkring 64 ms.
Det vil sige, at hastigheden på en bærbar computer er cirka 300 millioner indsættelser/sek. og 500 millioner sletninger/sek.
Tabellen er skrevet i CUDA, selvom den samme teknik kan anvendes på HLSL eller GLSL. Implementeringen har flere begrænsninger for at sikre høj ydeevne på et videokort:
- Kun 32-bit nøgler og de samme værdier behandles.
- Hashbordet har en fast størrelse.
- Og denne størrelse skal være lig med to til magten.
For nøgler og værdier skal du reservere en simpel afgrænsningsmarkør (i ovenstående kode er dette 0xffffffff).
Hashbord uden låse
Hash-tabellen bruger åben adressering med , det vil sige, at det simpelthen er en række nøgleværdi-par, der er gemt i hukommelsen og har overlegen cache-ydeevne. Det samme kan ikke siges om chaining, som involverer søgning efter en pointer i en sammenkædet liste. En hash-tabel er et simpelt array, der gemmer elementer KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Tabellens størrelse er en potens af to, ikke et primtal, fordi én hurtig instruktion er nok til at anvende pow2/AND-masken, men modulusoperatoren er meget langsommere. Dette er vigtigt i tilfælde af lineær sondering, da i et lineært tabelopslag skal slotindekset pakkes ind i hver slot. Og som et resultat tilføjes omkostningerne ved operationen modulo i hver slot.
Tabellen gemmer kun nøglen og værdien for hvert element, ikke en hash af nøglen. Da tabellen kun gemmer 32-bit nøgler, beregnes hashen meget hurtigt. Ovenstående kode bruger Murmur3-hash, som kun udfører nogle få skift, XOR'er og multiplikationer.
Hash-tabellen bruger låsebeskyttelsesteknikker, der er uafhængige af hukommelsesrækkefølge. Selvom nogle skriveoperationer forstyrrer rækkefølgen af andre sådanne operationer, vil hash-tabellen stadig bevare den korrekte tilstand. Vi vil tale om dette nedenfor. Teknikken fungerer godt med videokort, der kører tusindvis af tråde samtidigt.
Nøglerne og værdierne i hash-tabellen initialiseres til tomme.
Koden kan modificeres til også at håndtere 64-bit nøgler og værdier. Nøgler kræver atomare læse-, skrive- og sammenligne-og-bytte-operationer. Og værdier kræver atomare læse- og skriveoperationer. Heldigvis, i CUDA, er læse-skriveoperationer for 32- og 64-bit værdier atomare, så længe de er naturligt justeret (se nedenfor). ), og moderne videokort understøtter 64-bit atomic sammenligning-og-udveksling operationer. Selvfølgelig, når du flytter til 64 bit, vil ydeevnen falde lidt.
Hash-tabeltilstand
Hvert nøgle-værdi-par i en hash-tabel kan have en af fire tilstande:
- Nøgle og værdi er tomme. I denne tilstand initialiseres hash-tabellen.
- Nøglen er skrevet ned, men værdien er endnu ikke skrevet. Hvis en anden tråd i øjeblikket læser data, vender den tilbage tom. Dette er normalt, det samme ville være sket, hvis en anden udførelsestråd havde fungeret lidt tidligere, og vi taler om en samtidig datastruktur.
- Både nøglen og værdien registreres.
- Værdien er tilgængelig for andre udførelsestråde, men nøglen er ikke endnu. Dette kan ske, fordi CUDA-programmeringsmodellen har en løst ordnet hukommelsesmodel. Dette er normalt; under alle omstændigheder er nøglen stadig tom, selvom værdien ikke længere er det.
En vigtig nuance er, at når nøglen først er skrevet til spalten, bevæger den sig ikke længere - selvom nøglen slettes, vil vi tale om dette nedenfor.
Hashtabelkoden fungerer endda med løst ordnede hukommelsesmodeller, hvor rækkefølgen, hvori hukommelsen læses og skrives, er ukendt. Når vi ser på indsættelse, opslag og sletning i en hash-tabel, så husk, at hvert nøgle-værdi-par er i en af de fire tilstande beskrevet ovenfor.
Indsættelse i en hash-tabel
CUDA-funktionen, der indsætter nøgle-værdi-par i en hash-tabel, ser sådan ud:
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);
}
}
For at indsætte en nøgle, gentager koden gennem hash-tabel-arrayet, der starter med hashen for den indsatte nøgle. Hvert slot i arrayet udfører en atomisk sammenligning-og-swap-operation, der sammenligner nøglen i denne plads med tom. Hvis der opdages en uoverensstemmelse, opdateres nøglen i spalten med den indsatte nøgle, og derefter returneres den originale slotnøgle. Hvis denne originale nøgle var tom eller matchede den indsatte nøgle, fandt koden en passende plads til indsættelse og indsatte den indsatte værdi i spalten.
Hvis i et kernekald gpu_hashtable_insert() der er flere elementer med den samme nøgle, så kan enhver af deres værdier skrives til nøgleslottet. Dette anses for normalt: en af nøgleværdiskrivningerne under opkaldet vil lykkes, men da alt dette sker parallelt inden for flere udførelsestråde, kan vi ikke forudsige, hvilken hukommelsesskrivning der vil være den sidste.
Hash tabelopslag
Kode til søgenøgler:
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);
}
}
For at finde værdien af en nøgle, der er gemt i en tabel, gentager vi arrayet, der starter med hashen for den nøgle, vi leder efter. I hver slot tjekker vi, om nøglen er den, vi leder efter, og i så fald returnerer vi dens værdi. Vi tjekker også om nøglen er tom, og i så fald afbryder vi søgningen.
Hvis vi ikke kan finde nøglen, returnerer koden en tom værdi.
Alle disse søgeoperationer kan udføres samtidigt gennem indsættelser og sletninger. Hvert par i tabellen vil have en af de fire tilstande beskrevet ovenfor for flowet.
Sletter i en hash-tabel
Kode til sletning af nøgler:
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);
}
}
Sletning af en nøgle sker på en usædvanlig måde: vi efterlader nøglen i tabellen og markerer dens værdi (ikke selve nøglen) som tom. Denne kode minder meget om lookup(), bortset fra at når der findes et match på en nøgle, bliver dens værdi tom.
Som nævnt ovenfor, når en nøgle først er skrevet til et slot, flyttes den ikke længere. Selv når et element slettes fra tabellen, forbliver nøglen på plads, dens værdi bliver simpelthen tom. Det betyder, at vi ikke behøver at bruge en atomisk skriveoperation til slotværdien, for det er lige meget om den aktuelle værdi er tom eller ej, den bliver stadig tom.
Ændre størrelsen på en hash-tabel
Du kan ændre størrelsen på en hash-tabel ved at oprette en større tabel og indsætte ikke-tomme elementer fra den gamle tabel i den. Jeg implementerede ikke denne funktionalitet, fordi jeg ønskede at holde prøvekoden enkel. Desuden, i CUDA-programmer, udføres hukommelsesallokering ofte i værtskoden snarere end i CUDA-kernen.
I artiklen beskriver, hvordan man ændrer en sådan låsebeskyttet datastruktur.
Konkurrenceevne
I ovenstående funktionskodestykker gpu_hashtable_insert(), _lookup() и _delete() behandle et nøgleværdi-par ad gangen. Og lavere gpu_hashtable_insert(), _lookup() и _delete() behandle et array af par parallelt, hvert par i en separat GPU-udførelsestråd:
// 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);
}
}
Den låsesikre hash-tabel understøtter samtidige indsættelser, opslag og sletninger. Fordi nøgle-værdi-par altid er i en af fire tilstande, og tasterne ikke bevæger sig, garanterer tabellen korrekthed, selv når forskellige typer operationer bruges samtidigt.
Men hvis vi behandler en batch af indsættelser og sletninger parallelt, og hvis input-arrayet af par indeholder duplikerede nøgler, så vil vi ikke være i stand til at forudsige, hvilke par der vil "vinde" - vil blive skrevet til hash-tabellen sidst. Lad os sige, at vi kaldte indsættelseskoden med et input-array af par A/0 B/1 A/2 C/3 A/4. Når koden er færdig, parres B/1 и C/3 er garanteret til stede i tabellen, men samtidig vil et hvilket som helst af parrene optræde i den A/0, A/2 eller A/4. Dette kan være et problem eller ikke - det hele afhænger af applikationen. Du ved måske på forhånd, at der ikke er nogen duplikerede nøgler i input-arrayet, eller du er måske ligeglad med, hvilken værdi der blev skrevet sidst.
Hvis dette er et problem for dig, så skal du adskille dubletparrene i forskellige CUDA-systemkald. I CUDA fuldføres enhver operation, der kalder kernen, altid før det næste kernekald (mindst inden for én tråd. I forskellige tråde udføres kerner parallelt). I eksemplet ovenfor, hvis du kalder en kerne med A/0 B/1 A/2 C/3, og den anden med A/4, derefter nøglen A vil få værdien 4.
Lad os nu tale om, hvorvidt funktioner skal lookup() и delete() brug en almindelig eller flygtig pointer til en række par i hash-tabellen. Oplyser, at:
Compileren kan vælge at optimere læsning og skrivning til global eller delt hukommelse... Disse optimeringer kan deaktiveres ved hjælp af nøgleordet
volatile: ... enhver reference til denne variabel kompileres til en læse- eller skriveinstruktion i ægte hukommelse.
Rigtighedshensyn kræver ikke ansøgning volatile. Hvis udførelsestråden bruger en cachelagret værdi fra en tidligere læseoperation, vil den bruge lidt forældet information. Men stadigvæk er dette information fra den korrekte tilstand af hash-tabellen på et bestemt tidspunkt af kernekaldet. Hvis du skal bruge de seneste oplysninger, kan du bruge indekset volatile, men så vil ydeevnen falde en anelse: ifølge mine tests, ved sletning af 32 millioner elementer, faldt hastigheden fra 500 millioner sletninger/sek. til 450 millioner sletninger/sek.
Ydelse
I testen for at indsætte 64 millioner elementer og slette 32 millioner af dem, konkurrence mellem std::unordered_map og der er stort set ingen hash-tabel til GPU'en:
std::unordered_map brugte 70 ms på at indsætte og fjerne elementer og derefter frigøre dem unordered_map (at slippe af med millioner af elementer tager meget tid, fordi indeni unordered_map der foretages flere hukommelsestildelinger). Ærligt talt, std:unordered_map helt andre restriktioner. Det er en enkelt CPU-udførelsestråd, understøtter nøgleværdier af enhver størrelse, fungerer godt ved høje udnyttelsesgrader og viser stabil ydeevne efter flere sletninger.
Varigheden af hash-tabellen for GPU'en og kommunikation mellem programmer var 984 ms. Dette inkluderer den tid, der bruges på at placere tabellen i hukommelsen og slette den (tildeling af 1 GB hukommelse én gang, hvilket tager noget tid i CUDA), indsættelse og sletning af elementer og iteration over dem. Alle kopier til og fra videokortets hukommelse tages også i betragtning.
Selve hashtabellen tog 271 ms at færdiggøre. Dette inkluderer den tid, som videokortet bruger på at indsætte og slette elementer, og tager ikke højde for den tid, der bruges på at kopiere ind i hukommelsen og gentage den resulterende tabel. Hvis GPU-tabellen lever i lang tid, eller hvis hash-tabellen er indeholdt helt i hukommelsen på videokortet (for eksempel for at oprette en hash-tabel, der vil blive brugt af anden GPU-kode og ikke den centrale processor), så testresultatet er relevant.
Hash-tabellen til et videokort demonstrerer høj ydeevne på grund af høj gennemstrømning og aktiv parallelisering.
Begrænsninger
Hash-tabelarkitekturen har et par problemer, du skal være opmærksom på:
- Lineær sondering er hæmmet af klyngedannelse, hvilket bevirker, at nøglerne i tabellen placeres mindre end perfekt.
- Nøgler fjernes ikke ved hjælp af funktionen
deleteog med tiden roder de bordet.
Som et resultat kan ydeevnen af en hash-tabel gradvist forringes, især hvis den eksisterer i lang tid og har adskillige indsættelser og sletninger. En måde at afbøde disse ulemper på er at genhash til en ny tabel med en ret lav udnyttelsesgrad og filtrere de fjernede nøgler fra under genhasningen.
For at illustrere de beskrevne problemer, vil jeg bruge ovenstående kode til at oprette en tabel med 128 millioner elementer og gå gennem 4 millioner elementer, indtil jeg har udfyldt 124 millioner pladser (udnyttelsesgrad på ca. 0,96). Her er resultattabellen, hver række er et CUDA-kernekald til at indsætte 4 millioner nye elementer i én hash-tabel:
Brugsrate
Indsættelsesvarighed 4 elementer
0,00
11,608448 ms (361,314798 millioner nøgler/sek.)
0,03
11,751424 ms (356,918799 millioner nøgler/sek.)
0,06
11,942592 ms (351,205515 millioner nøgler/sek.)
0,09
12,081120 ms (347,178429 millioner nøgler/sek.)
0,12
12,242560 ms (342,600233 millioner nøgler/sek.)
0,16
12,396448 ms (338,347235 millioner nøgler/sek.)
0,19
12,533024 ms (334,660176 millioner nøgler/sek.)
0,22
12,703328 ms (330,173626 millioner nøgler/sek.)
0,25
12,884512 ms (325,530693 millioner nøgler/sek.)
0,28
13,033472 ms (321,810182 millioner nøgler/sek.)
0,31
13,239296 ms (316,807174 millioner nøgler/sek.)
0,34
13,392448 ms (313,184256 millioner nøgler/sek.)
0,37
13,624000 ms (307,861434 millioner nøgler/sek.)
0,41
13,875520 ms (302,280855 millioner nøgler/sek.)
0,44
14,126528 ms (296,909756 millioner nøgler/sek.)
0,47
14,399328 ms (291,284699 millioner nøgler/sek.)
0,50
14,690304 ms (285,515123 millioner nøgler/sek.)
0,53
15,039136 ms (278,892623 millioner nøgler/sek.)
0,56
15,478656 ms (270,973402 millioner nøgler/sek.)
0,59
15,985664 ms (262,379092 millioner nøgler/sek.)
0,62
16,668673 ms (251,627968 millioner nøgler/sek.)
0,66
17,587200 ms (238,486174 millioner nøgler/sek.)
0,69
18,690048 ms (224,413765 millioner nøgler/sek.)
0,72
20,278816 ms (206,831789 millioner nøgler/sek.)
0,75
22,545408 ms (186,038058 millioner nøgler/sek.)
0,78
26,053312 ms (160,989275 millioner nøgler/sek.)
0,81
31,895008 ms (131,503463 millioner nøgler/sek.)
0,84
42,103294 ms (99,619378 millioner nøgler/sek.)
0,87
61,849056 ms (67,815164 millioner nøgler/sek.)
0,90
105,695999 ms (39,682713 millioner nøgler/sek.)
0,94
240,204636 ms (17,461378 millioner nøgler/sek.)
Efterhånden som udnyttelsen stiger, falder ydeevnen. Dette er i de fleste tilfælde ikke ønskeligt. Hvis en applikation indsætter elementer i en tabel og derefter kasserer dem (f.eks. når man tæller ord i en bog), så er dette ikke et problem. Men hvis applikationen bruger en hash-tabel med lang levetid (f.eks. i en grafikeditor til at gemme ikke-tomme dele af billeder, hvor brugeren ofte indsætter og sletter information), så kan denne adfærd være problematisk.
Og målte hash-tabellens sonderingsdybde efter 64 millioner inserts (udnyttelsesfaktor 0,5). Den gennemsnitlige dybde var 0,4774, så de fleste nøgler var enten i den bedst mulige plads eller et slot væk fra den bedste position. Den maksimale lyddybde var 60.
Derefter målte jeg sonderingsdybden på et bord med 124 millioner skær (udnyttelsesfaktor 0,97). Den gennemsnitlige dybde var allerede 10,1757, og den maksimale - 6474 (!!). Lineær sensing ydeevne falder betydeligt ved høje udnyttelsesgrader.
Det er bedst at holde denne hash-tabels udnyttelsesgrad lav. Men så øger vi ydeevnen på bekostning af hukommelsesforbruget. Heldigvis, i tilfælde af 32-bit nøgler og værdier, kan dette retfærdiggøres. Hvis vi i eksemplet ovenfor, i en tabel med 128 millioner elementer, beholder udnyttelsesfaktoren på 0,25, så kan vi ikke placere mere end 32 millioner elementer i den, og de resterende 96 millioner slots vil gå tabt - 8 bytes for hvert par , 768 MB tabt hukommelse.
Bemærk venligst, at vi taler om tab af videokorthukommelse, som er en mere værdifuld ressource end systemhukommelse. Selvom de fleste moderne desktop-grafikkort, der understøtter CUDA, har mindst 4 GB hukommelse (i skrivende stund har NVIDIA 2080 Ti 11 GB), ville det stadig ikke være den klogeste beslutning at miste sådanne beløb.
Senere vil jeg skrive mere om at lave hashtabeller til videokort, der ikke har problemer med sonderingsdybden, samt måder at genbruge slettede slots.
Lyddybdemåling
For at bestemme sonderingsdybden af en nøgle kan vi udtrække nøglens hash (den ideelle tabelindeks) fra dens faktiske tabelindeks:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
På grund af magien ved tos to komplementære binære tal og det faktum, at kapaciteten af hash-tabellen er to potenseret af to, vil denne tilgang fungere, selv når nøgleindekset flyttes til begyndelsen af tabellen. Lad os tage en nøgle, der hash til 1, men er indsat i slot 3. Så får vi for et bord med kapacitet 4 (3 — 1) & 3, hvilket svarer til 2.
Konklusion
Hvis du har spørgsmål eller kommentarer, så send mig en mail på eller åbne et nyt emne i.
Denne kode er skrevet under inspiration fra fremragende artikler:
I fremtiden vil jeg fortsætte med at skrive om hash tabel implementeringer til videokort og analysere deres ydeevne. Mine planer omfatter chaining, Robin Hood-hashing og gøg-hashing ved hjælp af atomariske operationer i datastrukturer, der er GPU-venlige.
Kilde: www.habr.com