Jag la upp det på Github .
Det är en enkel GPU-hashtabell som kan bearbeta hundratals miljoner inlägg per sekund. På min NVIDIA GTX 1060 bärbara dator infogar koden 64 miljoner slumpmässigt genererade nyckel-värdepar på cirka 210 ms och tar bort 32 miljoner par på cirka 64 ms.
Det vill säga hastigheten på en bärbar dator är cirka 300 miljoner infogar/sek och 500 miljoner raderingar/sek.
Tabellen är skriven i CUDA, även om samma teknik kan tillämpas på HLSL eller GLSL. Implementeringen har flera begränsningar för att säkerställa hög prestanda på ett grafikkort:
- Endast 32-bitars nycklar och samma värden bearbetas.
- Hashbordet har en fast storlek.
- Och denna storlek måste vara lika med två till makten.
För nycklar och värden måste du reservera en enkel avgränsningsmarkör (i ovanstående kod är detta 0xffffffff).
Hastbord utan lås
Hashtabellen använder öppen adressering med , det vill säga det är helt enkelt en uppsättning nyckel-värdepar som lagras i minnet och har överlägsen cacheprestanda. Detsamma kan inte sägas om chaining, som innebär att man söker efter en pekare i en länkad lista. En hashtabell är en enkel array som lagrar element KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Storleken på tabellen är en potens av två, inte ett primtal, eftersom en snabb instruktion räcker för att applicera pow2/AND-masken, men moduloperatorn är mycket långsammare. Detta är viktigt i fallet med linjär sondering, eftersom i en linjär tabelluppslagning måste luckindexet lindas in i varje lucka. Och som ett resultat läggs kostnaden för operationen till modulo i varje lucka.
Tabellen lagrar bara nyckeln och värdet för varje element, inte en hash av nyckeln. Eftersom tabellen bara lagrar 32-bitars nycklar, beräknas hashen mycket snabbt. Koden ovan använder Murmur3-hash, som bara utför ett fåtal skift, XOR och multiplikationer.
Hashtabellen använder låsningsskyddstekniker som är oberoende av minnesordning. Även om vissa skrivoperationer stör ordningen för andra sådana operationer, kommer hashtabellen fortfarande att behålla det korrekta tillståndet. Vi kommer att prata om detta nedan. Tekniken fungerar utmärkt med grafikkort som kör tusentals trådar samtidigt.
Nycklarna och värdena i hash-tabellen initieras till tomma.
Koden kan modifieras för att hantera 64-bitars nycklar och värden också. Nycklar kräver atomära läs-, skriv- och jämför-och-byte-operationer. Och värden kräver atomära läs- och skrivoperationer. Lyckligtvis, i CUDA är läs-skrivoperationer för 32- och 64-bitars värden atomära så länge de är naturligt inriktade (se nedan). ), och moderna grafikkort stöder 64-bitars atomic compare-and-exchange operationer. Naturligtvis, när du flyttar till 64 bitar, kommer prestandan att minska något.
Hash-tabelltillstånd
Varje nyckel-värdepar i en hashtabell kan ha ett av fyra tillstånd:
- Nyckel och värde är tomma. I detta tillstånd initieras hashtabellen.
- Nyckeln har skrivits ned, men värdet har ännu inte skrivits ned. Om en annan tråd för närvarande läser data, returneras den tom. Detta är normalt, samma sak skulle ha hänt om en annan exekveringstråd hade fungerat lite tidigare, och vi pratar om en samtidig datastruktur.
- Både nyckeln och värdet registreras.
- Värdet är tillgängligt för andra exekveringstrådar, men nyckeln är inte ännu. Detta kan hända eftersom CUDA-programmeringsmodellen har en löst ordnad minnesmodell. Detta är normalt, i alla händelser är nyckeln fortfarande tom, även om värdet inte längre är det.
En viktig nyans är att när nyckeln väl har skrivits till luckan så rör den sig inte längre - även om nyckeln raderas kommer vi att prata om detta nedan.
Hashtabellkoden fungerar till och med med löst ordnade minnesmodeller där det inte är känt i vilken ordning minnet läses och skrivs. När vi tittar på infogning, uppslagning och radering i en hashtabell, kom ihåg att varje nyckel-värdepar är i ett av de fyra tillstånden som beskrivs ovan.
Infogar i en hashtabell
CUDA-funktionen som infogar nyckel-värdepar i en hashtabell ser ut så här:
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);
}
}
För att infoga en nyckel, itererar koden genom hashtabellen med början med hashen för den infogade nyckeln. Varje fack i arrayen utför en atomär jämförelse-och-byte-operation som jämför nyckeln i den fack med tom. Om en felmatchning upptäcks uppdateras nyckeln i luckan med den insatta nyckeln, och sedan returneras den ursprungliga lucknyckeln. Om denna ursprungliga nyckel var tom eller matchade den infogade nyckeln, så hittade koden en lämplig plats för insättning och infogade det infogade värdet i luckan.
Om i ett kärnanrop gpu_hashtable_insert() det finns flera element med samma nyckel, då kan vilket som helst av deras värden skrivas till nyckelfacket. Detta anses normalt: en av nyckel-värdesskrivningarna under samtalet kommer att lyckas, men eftersom allt detta sker parallellt inom flera exekveringstrådar kan vi inte förutsäga vilken minnesskrivning som kommer att vara den sista.
Uppslag i hashtabell
Kod för att söka nycklar:
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);
}
}
För att hitta värdet på en nyckel lagrad i en tabell, itererar vi genom arrayen och börjar med hashen för nyckeln vi letar efter. I varje slot kontrollerar vi om nyckeln är den vi letar efter, och i så fall returnerar vi dess värde. Vi kontrollerar även om nyckeln är tom och i så fall avbryter vi sökningen.
Om vi inte kan hitta nyckeln returnerar koden ett tomt värde.
Alla dessa sökoperationer kan utföras samtidigt genom infogning och radering. Varje par i tabellen kommer att ha ett av de fyra tillstånden som beskrivs ovan för flödet.
Ta bort i en hashtabell
Kod för att radera nycklar:
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);
}
}
Att ta bort en nyckel görs på ett ovanligt sätt: vi lämnar nyckeln i tabellen och markerar dess värde (inte själva nyckeln) som tom. Denna kod är väldigt lik lookup(), förutom att när en matchning hittas på en nyckel, gör den dess värde tomt.
Som nämnts ovan, när en nyckel väl har skrivits till en plats, flyttas den inte längre. Även när ett element raderas från tabellen förblir nyckeln på plats, dess värde blir helt enkelt tomt. Det betyder att vi inte behöver använda en atomskrivoperation för slotvärdet, eftersom det inte spelar någon roll om det aktuella värdet är tomt eller inte - det kommer fortfarande att bli tomt.
Ändra storlek på en hashtabell
Du kan ändra storleken på en hashtabell genom att skapa en större tabell och infoga icke-tomma element från den gamla tabellen i den. Jag implementerade inte den här funktionen eftersom jag ville hålla provkoden enkel. Dessutom, i CUDA-program görs minnesallokering ofta i värdkoden snarare än i CUDA-kärnan.
Artikeln beskriver hur man ändrar en sådan låsskyddad datastruktur.
Konkurrenskraft
I ovanstående funktionskodavsnitt gpu_hashtable_insert(), _lookup() и _delete() bearbeta ett nyckel-värdepar i taget. Och lägre gpu_hashtable_insert(), _lookup() и _delete() bearbeta en array av par parallellt, varje par i en separat GPU-exekveringstrå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åssäkra hashtabellen stöder samtidiga infogningar, uppslagningar och raderingar. Eftersom nyckel-värdepar alltid är i ett av fyra tillstånd och nycklarna inte rör sig, garanterar tabellen korrekthet även när olika typer av operationer används samtidigt.
Men om vi bearbetar en sats av infogningar och borttagningar parallellt, och om inmatningsarrayen av par innehåller dubbletter av nycklar, kommer vi inte att kunna förutsäga vilka par som kommer att "vinna" - kommer att skrivas till hashtabellen sist. Låt oss säga att vi kallade insättningskoden med en inmatningsarray av par A/0 B/1 A/2 C/3 A/4. När koden är klar, parar B/1 и C/3 är garanterat närvarande i tabellen, men samtidigt kommer vilket som helst av paren att synas i den A/0, A/2 eller A/4. Detta kan eller kanske inte är ett problem - allt beror på applikationen. Du kanske vet i förväg att det inte finns några dubbletter av nycklar i inmatningsmatrisen, eller så bryr du dig kanske inte vilket värde som skrevs sist.
Om detta är ett problem för dig, måste du separera dubblettparen i olika CUDA-systemanrop. I CUDA slutförs alltid alla operationer som anropar kärnan före nästa kärnanrop (åtminstone inom en tråd. I olika trådar exekveras kärnor parallellt). I exemplet ovan, om du anropar en kärna med A/0 B/1 A/2 C/3, och den andra med A/4, sedan nyckeln A kommer att få värdet 4.
Låt oss nu prata om funktioner bör lookup() и delete() använd en vanlig eller flyktig pekare till en array av par i hashtabellen. Stater som:
Kompilatorn kan välja att optimera läsning och skrivning till globalt eller delat minne... Dessa optimeringar kan inaktiveras med nyckelordet
volatile: ... alla referenser till denna variabel kompileras till en läs- eller skrivinstruktion för ett riktigt minne.
Korrekthetshänsyn kräver ingen ansökan volatile. Om exekveringstråden använder ett cachat värde från en tidigare läsoperation, kommer den att använda något föråldrad information. Men ändå är detta information från det korrekta tillståndet för hashtabellen vid ett visst ögonblick av kärnanropet. Om du behöver använda den senaste informationen kan du använda indexet volatile, men då kommer prestandan att minska något: enligt mina tester, vid radering av 32 miljoner element, minskade hastigheten från 500 miljoner raderingar/sek till 450 miljoner raderingar/sek.
Производительность
I testet för att infoga 64 miljoner element och ta bort 32 miljoner av dem, konkurrens mellan std::unordered_map och det finns praktiskt taget ingen hashtabell för GPU:n:
std::unordered_map spenderade 70 691 ms på att sätta in och ta bort element och sedan frigöra dem unordered_map (att bli av med miljontals element tar mycket tid, för inuti unordered_map flera minnestilldelningar görs). Ärligt talat, std:unordered_map helt andra restriktioner. Det är en enda CPU-tråd för exekvering, stöder nyckel-värden av alla storlekar, presterar bra vid höga utnyttjandegrader och visar stabil prestanda efter flera raderingar.
Längden på hashtabellen för GPU:n och kommunikation mellan program var 984 ms. Detta inkluderar den tid som ägnas åt att placera tabellen i minnet och radera den (tilldela 1 GB minne en gång, vilket tar lite tid i CUDA), att infoga och ta bort element och iterera över dem. Alla kopior till och från grafikkortets minne beaktas också.
Själva hashtabellen tog 271 ms att slutföra. Detta inkluderar den tid som ägnas åt att grafikkortet sätter i och raderar element, och tar inte hänsyn till tiden för att kopiera in i minnet och iterera över den resulterande tabellen. Om GPU-tabellen lever under en lång tid, eller om hashtabellen helt och hållet finns i grafikkortets minne (till exempel för att skapa en hashtabell som kommer att användas av annan GPU-kod och inte centralprocessorn), då testresultatet är relevant.
Hashtabellen för ett grafikkort visar hög prestanda på grund av hög genomströmning och aktiv parallellisering.
Begränsningar
Hashtabellarkitekturen har några problem att vara medveten om:
- Linjär sondering försvåras av klustring, vilket gör att nycklarna i tabellen placeras mindre än perfekt.
- Nycklar tas inte bort med funktionen
deleteoch med tiden rör de bordet.
Som ett resultat kan prestandan för en hashtabell gradvis försämras, särskilt om den existerar under en lång tid och har många infogningar och borttagningar. Ett sätt att mildra dessa nackdelar är att omhasha till en ny tabell med en ganska låg utnyttjandegrad och filtrera bort de borttagna nycklarna under omhasningen.
För att illustrera problemen som beskrivs kommer jag att använda ovanstående kod för att skapa en tabell med 128 miljoner element och gå igenom 4 miljoner element tills jag har fyllt 124 miljoner platser (utnyttjandegrad på cirka 0,96). Här är resultattabellen, varje rad är ett CUDA-kärnanrop för att infoga 4 miljoner nya element i en hashtabell:
Användningsgrad
Insättningslängd 4 194 304 element
0,00
11,608448 ms (361,314798 miljoner nycklar/sek.)
0,03
11,751424 ms (356,918799 miljoner nycklar/sek.)
0,06
11,942592 ms (351,205515 miljoner nycklar/sek.)
0,09
12,081120 ms (347,178429 miljoner nycklar/sek.)
0,12
12,242560 ms (342,600233 miljoner nycklar/sek.)
0,16
12,396448 ms (338,347235 miljoner nycklar/sek.)
0,19
12,533024 ms (334,660176 miljoner nycklar/sek.)
0,22
12,703328 ms (330,173626 miljoner nycklar/sek.)
0,25
12,884512 ms (325,530693 miljoner nycklar/sek.)
0,28
13,033472 ms (321,810182 miljoner nycklar/sek.)
0,31
13,239296 ms (316,807174 miljoner nycklar/sek.)
0,34
13,392448 ms (313,184256 miljoner nycklar/sek.)
0,37
13,624000 ms (307,861434 miljoner nycklar/sek.)
0,41
13,875520 ms (302,280855 miljoner nycklar/sek.)
0,44
14,126528 ms (296,909756 miljoner nycklar/sek.)
0,47
14,399328 ms (291,284699 miljoner nycklar/sek.)
0,50
14,690304 ms (285,515123 miljoner nycklar/sek.)
0,53
15,039136 ms (278,892623 miljoner nycklar/sek.)
0,56
15,478656 ms (270,973402 miljoner nycklar/sek.)
0,59
15,985664 ms (262,379092 miljoner nycklar/sek.)
0,62
16,668673 ms (251,627968 miljoner nycklar/sek.)
0,66
17,587200 ms (238,486174 miljoner nycklar/sek.)
0,69
18,690048 ms (224,413765 miljoner nycklar/sek.)
0,72
20,278816 ms (206,831789 miljoner nycklar/sek.)
0,75
22,545408 ms (186,038058 miljoner nycklar/sek.)
0,78
26,053312 ms (160,989275 miljoner nycklar/sek.)
0,81
31,895008 ms (131,503463 miljoner nycklar/sek.)
0,84
42,103294 ms (99,619378 miljoner nycklar/sek.)
0,87
61,849056 ms (67,815164 miljoner nycklar/sek.)
0,90
105,695999 ms (39,682713 miljoner nycklar/sek.)
0,94
240,204636 ms (17,461378 miljoner nycklar/sek.)
När utnyttjandet ökar minskar prestandan. Detta är inte önskvärt i de flesta fall. Om ett program infogar element i en tabell och sedan kasserar dem (till exempel när man räknar ord i en bok), är detta inget problem. Men om applikationen använder en hashtabell med lång livslängd (till exempel i en grafikredigerare för att lagra icke-tomma delar av bilder där användaren ofta infogar och raderar information), kan detta beteende vara problematiskt.
Och mätte hashtabellens sonderingsdjup efter 64 miljoner skär (utnyttjandefaktor 0,5). Det genomsnittliga djupet var 0,4774, så de flesta nycklar var antingen i den bästa möjliga luckan eller en lucka från den bästa positionen. Det maximala ljuddjupet var 60.
Jag mätte sedan sonderingsdjupet på ett bord med 124 miljoner skär (utnyttjandefaktor 0,97). Det genomsnittliga djupet var redan 10,1757, och det maximala - 6474 (!!). Linjära avkänningsprestanda sjunker avsevärt vid höga utnyttjandegrader.
Det är bäst att hålla denna hashtabells utnyttjandegrad låg. Men då ökar vi prestandan på bekostnad av minnesförbrukningen. Lyckligtvis, när det gäller 32-bitars nycklar och värden, kan detta motiveras. Om vi i exemplet ovan, i en tabell med 128 miljoner element, behåller utnyttjandefaktorn på 0,25, kan vi inte placera mer än 32 miljoner element i den, och de återstående 96 miljoner platserna kommer att gå förlorade - 8 byte för varje par , 768 MB förlorat minne.
Observera att vi talar om förlust av grafikkortsminne, vilket är en mer värdefull resurs än systemminne. Även om de flesta moderna skrivbordsgrafikkort som stöder CUDA har minst 4 GB minne (i skrivande stund har NVIDIA 2080 Ti 11 GB) så vore det ändå inte det klokaste beslutet att förlora sådana mängder.
Senare kommer jag att skriva mer om att skapa hashtabeller för grafikkort som inte har problem med sonderingsdjup, samt sätt att återanvända borttagna slots.
Ljuddjupmätning
För att bestämma sonderingsdjupet för en nyckel kan vi extrahera nyckelns hash (dess idealiska tabellindex) från dess faktiska tabellindex:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
På grund av magin med tvås två komplement binära tal och det faktum att hashtabellens kapacitet är två till två, kommer detta tillvägagångssätt att fungera även när nyckelindexet flyttas till början av tabellen. Låt oss ta en nyckel som hashas till 1, men som sätts in i fack 3. Sedan får vi för ett bord med kapacitet 4 (3 — 1) & 3, vilket motsvarar 2.
Slutsats
Om du har frågor eller kommentarer, vänligen maila mig på eller öppna ett nytt ämne i.
Denna kod skrevs med inspiration från utmärkta artiklar:
I framtiden kommer jag att fortsätta skriva om hashtabellimplementationer för grafikkort och analysera deras prestanda. Mina planer inkluderar chaining, Robin Hood-hashing och gökhashning med hjälp av atomoperationer i datastrukturer som är GPU-vänliga.
Källa: will.com