Jeg la det ut på Github .
Det er en enkel GPU-hash-tabell som er i stand til å behandle hundrevis av millioner av innlegg per sekund. På min NVIDIA GTX 1060 bærbare datamaskin setter koden inn 64 millioner tilfeldig genererte nøkkelverdi-par på omtrent 210 ms og fjerner 32 millioner par på omtrent 64 ms.
Det vil si at hastigheten på en bærbar PC er omtrent 300 millioner innlegg/sek og 500 millioner slettinger/sek.
Tabellen er skrevet i CUDA, selv om den samme teknikken kan brukes på HLSL eller GLSL. Implementeringen har flere begrensninger for å sikre høy ytelse på et skjermkort:
- Bare 32-bits nøkler og de samme verdiene behandles.
- Hash-bordet har en fast størrelse.
- Og denne størrelsen må være lik to til makten.
For nøkler og verdier må du reservere en enkel skillemarkør (i koden ovenfor er dette 0xffffffff).
Hashbord uten låser
Hash-tabellen bruker åpen adressering med , det vil si at det ganske enkelt er en rekke nøkkelverdi-par som er lagret i minnet og har overlegen hurtigbufferytelse. Det samme kan ikke sies om kjeding, som innebærer å søke etter en peker i en koblet liste. En hash-tabell er en enkel matrise som lagrer elementer KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Størrelsen på tabellen er en potens av to, ikke et primtall, fordi én rask instruksjon er nok til å bruke pow2/AND-masken, men modulusoperatoren er mye tregere. Dette er viktig i tilfellet med lineær sondering, siden i et lineært tabelloppslag må sporindeksen pakkes inn i hvert spor. Og som et resultat blir kostnaden for operasjonen lagt til modulo i hvert spor.
Tabellen lagrer bare nøkkelen og verdien for hvert element, ikke en hash av nøkkelen. Siden tabellen kun lagrer 32-biters nøkler, beregnes hashen veldig raskt. Koden ovenfor bruker Murmur3-hashen, som bare utfører noen få skift, XOR-er og multiplikasjoner.
Hash-tabellen bruker låsebeskyttelsesteknikker som er uavhengige av minnerekkefølge. Selv om noen skriveoperasjoner forstyrrer rekkefølgen til andre slike operasjoner, vil hashtabellen fortsatt opprettholde riktig tilstand. Vi snakker om dette nedenfor. Teknikken fungerer utmerket med skjermkort som kjører tusenvis av tråder samtidig.
Nøklene og verdiene i hash-tabellen initialiseres til tomme.
Koden kan modifiseres for å håndtere 64-bits nøkler og verdier også. Nøkler krever atomære lese-, skrive- og sammenligne-og-bytte operasjoner. Og verdier krever atomære lese- og skriveoperasjoner. Heldigvis, i CUDA, er lese-skriveoperasjoner for 32- og 64-bits verdier atomære så lenge de er naturlig justert (se nedenfor). ), og moderne skjermkort støtter 64-bit atomiske sammenligning-og-utvekslingsoperasjoner. Selvfølgelig, når du flytter til 64 bits, vil ytelsen reduseres litt.
Hash-tabelltilstand
Hvert nøkkelverdi-par i en hashtabell kan ha en av fire tilstander:
- Nøkkel og verdi er tomme. I denne tilstanden initialiseres hashtabellen.
- Nøkkelen er skrevet ned, men verdien er ennå ikke skrevet. Hvis en annen tråd for øyeblikket leser data, returnerer den tom. Dette er normalt, det samme ville ha skjedd hvis en annen utførelsestråd hadde fungert litt tidligere, og vi snakker om en samtidig datastruktur.
- Både nøkkelen og verdien registreres.
- Verdien er tilgjengelig for andre utførelsestråder, men nøkkelen er ikke ennå. Dette kan skje fordi CUDA-programmeringsmodellen har en løst ordnet minnemodell. Dette er normalt; i alle fall er nøkkelen fortsatt tom, selv om verdien ikke lenger er det.
En viktig nyanse er at når nøkkelen først er skrevet til sporet, beveger den seg ikke lenger - selv om nøkkelen slettes, vil vi snakke om dette nedenfor.
Hash-tabellkoden fungerer til og med med løst ordnede minnemodeller der rekkefølgen minnet leses og skrives i er ukjent. Når vi ser på innsetting, oppslag og sletting i en hash-tabell, husk at hvert nøkkelverdi-par er i en av de fire tilstandene beskrevet ovenfor.
Setter inn i en hash-tabell
CUDA-funksjonen som setter inn nøkkelverdi-par i en hash-tabell ser slik ut:
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 å sette inn en nøkkel, itererer koden gjennom hash-tabellmatrisen og starter med hashen til den innsatte nøkkelen. Hvert spor i arrayet utfører en atomisk sammenligning-og-bytt operasjon som sammenligner nøkkelen i det sporet med tom. Hvis en mismatch oppdages, oppdateres nøkkelen i sporet med den innsatte nøkkelen, og deretter returneres den originale spornøkkelen. Hvis denne originalnøkkelen var tom eller samsvarte med den innsatte nøkkelen, fant koden et passende spor for innsetting og satte inn den innsatte verdien i sporet.
Hvis i ett kjernekall gpu_hashtable_insert() det er flere elementer med samme nøkkel, så kan hvilken som helst av verdiene deres skrives til nøkkelsporet. Dette anses som normalt: en av nøkkelverdi-skrivingene under samtalen vil lykkes, men siden alt dette skjer parallelt innenfor flere utføringstråder, kan vi ikke forutsi hvilken minneskriving som vil være den siste.
Hash-tabelloppslag
Kode for søkenøkler:
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 å finne verdien av en nøkkel lagret i en tabell, itererer vi gjennom matrisen og starter med hashen til nøkkelen vi leter etter. I hvert spor sjekker vi om nøkkelen er den vi ser etter, og i så fall returnerer vi verdien. Vi sjekker også om nøkkelen er tom, og i så fall avbryter vi søket.
Hvis vi ikke finner nøkkelen, returnerer koden en tom verdi.
Alle disse søkeoperasjonene kan utføres samtidig gjennom innsettinger og slettinger. Hvert par i tabellen vil ha en av de fire tilstandene beskrevet ovenfor for flyten.
Sletter i en hash-tabell
Kode for sletting av nøkler:
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);
}
}
Sletting av en nøkkel gjøres på en uvanlig måte: vi lar nøkkelen stå i tabellen og markerer verdien (ikke selve nøkkelen) som tom. Denne koden er veldig lik lookup(), bortsett fra at når et samsvar blir funnet på en nøkkel, blir verdien tom.
Som nevnt ovenfor, når en nøkkel er skrevet til et spor, flyttes den ikke lenger. Selv når et element slettes fra tabellen, forblir nøkkelen på plass, verdien blir ganske enkelt tom. Dette betyr at vi ikke trenger å bruke en atomskriveoperasjon for sporverdien, fordi det spiller ingen rolle om gjeldende verdi er tom eller ikke - den vil fortsatt bli tom.
Endre størrelsen på en hash-tabell
Du kan endre størrelsen på en hash-tabell ved å lage en større tabell og sette inn ikke-tomme elementer fra den gamle tabellen i den. Jeg implementerte ikke denne funksjonaliteten fordi jeg ønsket å holde prøvekoden enkel. Dessuten, i CUDA-programmer, gjøres minneallokering ofte i vertskoden i stedet for i CUDA-kjernen.
Artikkelen beskriver hvordan du endrer en slik låsbeskyttet datastruktur.
Konkurranseevne
I funksjonskodebitene ovenfor gpu_hashtable_insert(), _lookup() и _delete() behandle ett nøkkelverdi-par om gangen. Og lavere gpu_hashtable_insert(), _lookup() и _delete() behandle en rekke par parallelt, hvert par i en separat GPU-utfø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 hashtabellen støtter samtidige innsettinger, oppslag og slettinger. Fordi nøkkel-verdi-par alltid er i en av fire tilstander og nøklene ikke beveger seg, garanterer tabellen korrekthet selv når forskjellige typer operasjoner brukes samtidig.
Imidlertid, hvis vi behandler en gruppe innsettinger og slettinger parallelt, og hvis inngangsarrayen med par inneholder dupliserte nøkler, vil vi ikke være i stand til å forutsi hvilke par som vil "vinne" – vil bli skrevet til hash-tabellen sist. La oss si at vi kalte innsettingskoden med en inngangsarray av par A/0 B/1 A/2 C/3 A/4. Når koden er fullført, pares B/1 и C/3 er garantert tilstede i tabellen, men samtidig vil hvilket som helst av parene vises i den A/0, A/2 eller A/4. Dette kan være et problem eller ikke - alt avhenger av applikasjonen. Du vet kanskje på forhånd at det ikke er noen dupliserte nøkler i inndatamatrisen, eller du bryr deg kanskje ikke om hvilken verdi som ble skrevet sist.
Hvis dette er et problem for deg, må du separere duplikatparene i forskjellige CUDA-systemanrop. I CUDA fullføres alltid enhver operasjon som kaller kjernen før neste kjernekall (minst innenfor én tråd. I forskjellige tråder utføres kjernene parallelt). I eksemplet ovenfor, hvis du kaller en kjerne med A/0 B/1 A/2 C/3, og den andre med A/4, deretter nøkkelen A vil få verdien 4.
La oss nå snakke om funksjoner bør lookup() и delete() bruk en vanlig eller flyktig peker til en rekke par i hashtabellen. Stater som:
Kompilatoren kan velge å optimere lesing og skriving til globalt eller delt minne... Disse optimaliseringene kan deaktiveres ved å bruke nøkkelordet
volatile: ... enhver referanse til denne variabelen kompileres til en lese- eller skriveinstruksjon i ekte minne.
Riktighetshensyn krever ikke søknad volatile. Hvis utførelsestråden bruker en bufret verdi fra en tidligere leseoperasjon, vil den bruke litt utdatert informasjon. Men likevel er dette informasjon fra den riktige tilstanden til hash-tabellen på et bestemt tidspunkt av kjernekallet. Hvis du trenger å bruke den nyeste informasjonen, kan du bruke indeksen volatile, men da vil ytelsen reduseres litt: ifølge mine tester, når du sletter 32 millioner elementer, sank hastigheten fra 500 millioner slettinger/sek. til 450 millioner slettinger/sek.
Производительность
I testen for å sette inn 64 millioner elementer og slette 32 millioner av dem, konkurranse mellom std::unordered_map og det er praktisk talt ingen hash-tabell for GPUen:
std::unordered_map brukte 70 691 ms på å sette inn og fjerne elementer og deretter frigjøre dem unordered_map (å bli kvitt millioner av elementer tar mye tid, fordi innvendig unordered_map flere minnetildelinger gjøres). Ærlig talt, std:unordered_map helt andre restriksjoner. Det er en enkelt CPU-tråd for utførelse, støtter nøkkelverdier av alle størrelser, yter godt ved høye utnyttelsesgrader og viser stabil ytelse etter flere slettinger.
Varigheten av hashtabellen for GPU og inter-program kommunikasjon var 984 ms. Dette inkluderer tiden brukt på å plassere tabellen i minnet og slette den (tildele 1 GB minne én gang, noe som tar litt tid i CUDA), sette inn og slette elementer og iterere over dem. Alle kopier til og fra skjermkortminnet tas også i betraktning.
Selve hashtabellen tok 271 ms å fullføre. Dette inkluderer tiden brukt på skjermkortet til å sette inn og slette elementer, og tar ikke hensyn til tiden brukt på å kopiere inn i minnet og iterere over den resulterende tabellen. Hvis GPU-tabellen lever lenge, eller hvis hash-tabellen er inneholdt i minnet på skjermkortet (for eksempel for å lage en hash-tabell som skal brukes av annen GPU-kode og ikke sentralprosessoren), testresultatet er relevant.
Hash-tabellen for et skjermkort viser høy ytelse på grunn av høy gjennomstrømning og aktiv parallellisering.
Begrensninger
Hash-tabellarkitekturen har noen problemer å være klar over:
- Lineær sondering hindres av klynging, noe som fører til at nøklene i tabellen blir plassert mindre enn perfekt.
- Nøkler fjernes ikke ved hjelp av funksjonen
deleteog over tid roter de bordet.
Som et resultat kan ytelsen til en hashtabell gradvis forringes, spesielt hvis den eksisterer i lang tid og har mange innsettinger og slettinger. En måte å redusere disse ulempene på er å rehash inn i en ny tabell med en ganske lav utnyttelsesgrad og filtrere ut de fjernede nøklene under rehashing.
For å illustrere problemene som er beskrevet, vil jeg bruke koden ovenfor til å lage en tabell med 128 millioner elementer og gå gjennom 4 millioner elementer til jeg har fylt 124 millioner spor (utnyttelsesgrad på ca. 0,96). Her er resultattabellen, hver rad er et CUDA-kjernekall for å sette inn 4 millioner nye elementer i én hashtabell:
Bruksrate
Innsettingsvarighet 4 194 304 elementer
0,00
11,608448 ms (361,314798 millioner nøkler/sek.)
0,03
11,751424 ms (356,918799 millioner nøkler/sek.)
0,06
11,942592 ms (351,205515 millioner nøkler/sek.)
0,09
12,081120 ms (347,178429 millioner nøkler/sek.)
0,12
12,242560 ms (342,600233 millioner nøkler/sek.)
0,16
12,396448 ms (338,347235 millioner nøkler/sek.)
0,19
12,533024 ms (334,660176 millioner nøkler/sek.)
0,22
12,703328 ms (330,173626 millioner nøkler/sek.)
0,25
12,884512 ms (325,530693 millioner nøkler/sek.)
0,28
13,033472 ms (321,810182 millioner nøkler/sek.)
0,31
13,239296 ms (316,807174 millioner nøkler/sek.)
0,34
13,392448 ms (313,184256 millioner nøkler/sek.)
0,37
13,624000 ms (307,861434 millioner nøkler/sek.)
0,41
13,875520 ms (302,280855 millioner nøkler/sek.)
0,44
14,126528 ms (296,909756 millioner nøkler/sek.)
0,47
14,399328 ms (291,284699 millioner nøkler/sek.)
0,50
14,690304 ms (285,515123 millioner nøkler/sek.)
0,53
15,039136 ms (278,892623 millioner nøkler/sek.)
0,56
15,478656 ms (270,973402 millioner nøkler/sek.)
0,59
15,985664 ms (262,379092 millioner nøkler/sek.)
0,62
16,668673 ms (251,627968 millioner nøkler/sek.)
0,66
17,587200 ms (238,486174 millioner nøkler/sek.)
0,69
18,690048 ms (224,413765 millioner nøkler/sek.)
0,72
20,278816 ms (206,831789 millioner nøkler/sek.)
0,75
22,545408 ms (186,038058 millioner nøkler/sek.)
0,78
26,053312 ms (160,989275 millioner nøkler/sek.)
0,81
31,895008 ms (131,503463 millioner nøkler/sek.)
0,84
42,103294 ms (99,619378 millioner nøkler/sek.)
0,87
61,849056 ms (67,815164 millioner nøkler/sek.)
0,90
105,695999 ms (39,682713 millioner nøkler/sek.)
0,94
240,204636 ms (17,461378 millioner nøkler/sek.)
Når utnyttelsen øker, reduseres ytelsen. Dette er ikke ønskelig i de fleste tilfeller. Hvis en applikasjon setter inn elementer i en tabell og deretter forkaster dem (for eksempel når man teller ord i en bok), er ikke dette et problem. Men hvis applikasjonen bruker en hashtabell med lang levetid (for eksempel i et grafikkredigeringsprogram for å lagre ikke-tomme deler av bilder der brukeren ofte setter inn og sletter informasjon), kan denne oppførselen være problematisk.
Og målte hashtabellens sonderingsdybde etter 64 millioner innlegg (utnyttelsesfaktor 0,5). Gjennomsnittlig dybde var 0,4774, så de fleste nøklene var enten i best mulig spor eller ett spor unna den beste posisjonen. Maksimal lyddybde var 60.
Deretter målte jeg sonderingsdybden på et bord med 124 millioner innsatser (utnyttelsesfaktor 0,97). Den gjennomsnittlige dybden var allerede 10,1757, og den maksimale - 6474 (!!). Lineær sensorytelse synker betydelig ved høye utnyttelsesgrader.
Det er best å holde denne hashtabellens utnyttelsesgrad lav. Men så øker vi ytelsen på bekostning av minneforbruket. Heldigvis, når det gjelder 32-biters nøkler og verdier, kan dette rettferdiggjøres. Hvis vi i eksemplet ovenfor, i en tabell med 128 millioner elementer, beholder utnyttelsesfaktoren på 0,25, kan vi ikke plassere mer enn 32 millioner elementer i den, og de resterende 96 millioner sporene vil gå tapt - 8 byte for hvert par , 768 MB tapt minne.
Vær oppmerksom på at vi snakker om tap av skjermkortminne, som er en mer verdifull ressurs enn systemminne. Selv om de fleste moderne stasjonære grafikkort som støtter CUDA har minst 4 GB minne (i skrivende stund har NVIDIA 2080 Ti 11 GB), vil det likevel ikke være den klokeste avgjørelsen å tape slike beløp.
Senere vil jeg skrive mer om å lage hashtabeller for skjermkort som ikke har problemer med sonderingsdybde, samt måter å gjenbruke slettede spor.
Lyddybdemåling
For å bestemme sonderingsdybden til en nøkkel, kan vi trekke ut nøkkelens hash (den ideelle tabellindeksen) fra den faktiske tabellindeksen:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
På grunn av magien med tos to komplementære binære tall og det faktum at kapasiteten til hash-tabellen er to i potens av to, vil denne tilnærmingen fungere selv når nøkkelindeksen flyttes til begynnelsen av tabellen. La oss ta en nøkkel som hash til 1, men som settes inn i spor 3. Så for et bord med kapasitet 4 får vi (3 — 1) & 3, som tilsvarer 2.
Konklusjon
Hvis du har spørsmål eller kommentarer, vennligst send meg en e-post på eller åpne et nytt emne i.
Denne koden ble skrevet under inspirasjon fra utmerkede artikler:
I fremtiden vil jeg fortsette å skrive om hashtabellimplementeringer for skjermkort og analysere ytelsen deres. Planene mine inkluderer kjetting, Robin Hood-hashing og gjøk-hashing ved å bruke atomoperasjoner i datastrukturer som er GPU-vennlige.
Kilde: www.habr.com