Opublikowałem na Githubie .
To prosta tabela haszująca dla GPU, zdolna do przetwarzania setek milionów wstawek na sekundę. Na moim laptopie z kartą NVIDIA GTX 1060 kod wstawia 64 miliony losowo wygenerowanych par klucz-wartość w około 210 ms i usuwa 32 miliony par w około 64 ms.
Oznacza to, że prędkość na laptopie wynosi około 300 mln wstawek/s i 500 mln usunięć/s.
Tabela została napisana w CUDA, chociaż tę samą metodę można zastosować w HLSL lub GLSL. Implementacja ma kilka ograniczeń, które zapewniają wysoką wydajność na karcie graficznej:
- Przetwarzane są tylko 32-bitowe klucze i odpowiadające im wartości.
- Tabela haszująca ma stały rozmiar.
- A ten rozmiar musi być równy potędze liczby dwa.
Dla kluczy i wartości należy zarezerwować prosty separator (w zaprezentowanym kodzie to 0xffffffff).
Tabela haszująca bez blokad
W tabeli haszującej zastosowano otwartą adresację z , czyli to po prostu tablica par klucz-wartość, która jest przechowywana w pamięci i ma doskonałą wydajność pamięci podręcznej. Nie można tego powiedzieć o łańcuchowaniu (chaining), które zakłada wyszukiwanie wskaźnika w powiązanej liście. Tabela haszująca jest prostą tablicą przechowującą elementy KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Rozmiar tabeli jest potęgą liczby dwa, a nie liczbą pierwszą, ponieważ zastosowanie maski pow2/AND wymaga tylko jednej szybkiej instrukcji, podczas gdy operator modulo działa znacznie wolniej. To istotne w przypadku linear probing, ponieważ przy liniowym wyszukiwaniu w tabeli indeks slotu musi być owinięty w każdym slocie. W efekcie dodawany jest koszt operacji modulo w każdym slocie.
Tabela przechowuje tylko klucz i wartość dla każdego elementu, a nie hasz klucza. Ponieważ tabela przechowuje jedynie 32-bitowe klucze, hasz oblicza się bardzo szybko. W zaprezentowanym kodzie stosuje się hasz Murmur3, który wykonuje jedynie kilka przesunięć, operacji XOR i mnożeń.
W tablicy haszowej stosuje się metody ochrony przed blokadami, które nie zależą od kolejności umieszczania w pamięci. Nawet jeśli niektóre operacje zapisu naruszają kolejność innych takich operacji, tablica haszowa wciąż zachowa poprawny stan. O tym porozmawiamy poniżej. Metoda doskonale działa z kartami graficznymi, w których konkurencyjnie wykonuje się tysiące wątków.
Klucze i wartości w tablicy haszowej są inicjowane jako puste.
Kod można zmodyfikować, aby obsługiwał zarówno 64-bitowe klucze, jak i wartości. Klucze wymagają atomowych operacji odczytu, zapisu i porównania z wymianą (compare-and-swap). A dla wartości potrzebne są atomowe operacje odczytu i zapisu. Na szczęście, w CUDA operacje odczytu-zapisu dla wartości 32- i 64-bitowych są atomowe, o ile są naturalnie wyrównane (patrz ), a nowoczesne karty graficzne wspierają 64-bitowe atomowe operacje porównania z wymianą. Oczywiście, przy przejściu na 64 bity wydajność nieco spadnie.
Stan tablicy haszowej
Każda para klucz-wartość w tablicy haszowej może mieć jeden z czterech stanów:
- Klucz i wartość są puste. W takim stanie tablica haszowa jest inicjowana.
- Klucz został zapisany, a wartość jeszcze nie. Jeśli inny wątek w tym momencie odczytuje dane, zwróci wtedy pustą wartość. To normalne, to samo wydarzyłoby się, gdyby inny wątek wykonał operację nieco wcześniej, a mówimy o współbieżnej strukturze danych.
- Zapisano zarówno klucz, jak i wartość.
- Wartość jest dostępna dla innych wątków, a klucz jeszcze nie. Tak może być, ponieważ model programowania w CUDA zakłada słabo uporządkowany model pamięci. To normalne, w każdym przypadku klucz wciąż jest pusty, nawet jeśli wartość już nie jest.
Ważny szczegół polega na tym, że jak tylko klucz został zapisany w slocie, nie jest już przenoszony — nawet jeśli klucz zostanie usunięty, o tym porozmawiamy poniżej.
Kod tablicy haszowej działa nawet w słabo uporządkowanych modelach pamięci, w których nie znany jest porządek odczytu i zapisów w pamięci. Gdy będziemy omawiać wstawianie, wyszukiwanie i usuwanie w tablicy haszowej, pamiętajcie, że każda para klucz-wartość znajduje się w jednym z czterech opisanych powyżej stanów.
Wstawianie do tablicy haszowej
Funkcja CUDA, która wstawia pary klucz-wartość do tabeli haszowej, wygląda następująco:
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);
}
}
Aby wstawić klucz, kod iteruje przez tablicę tabeli haszowej, zaczynając od hasza wstawianego klucza. W każdej komórce tablicy wykonywana jest atomowa operacja porównania z wymianą, w której klucz w tej komórce jest porównywany z pustym. Jeśli napotkano niespójność, klucz w komórce jest aktualizowany na wstawiany klucz, a następnie zwracany jest oryginalny klucz komórki. Jeśli ten oryginalny klucz był pusty lub odpowiadał kluczowi wstawianemu, to kod znalazł odpowiednią komórkę do wstawienia i umieszcza w niej wstawianą wartość.
Jeśli w jednym wywołaniu jądra gpu_hashtable_insert() jest kilka elementów o tym samym kluczu, wtedy dowolna z ich wartości może zostać zapisana w komórce klucza. To jest uznawane za normalne: jedna z operacji zapisu klucz-wartość podczas wywołania będzie udana, ale ponieważ wszystko to dzieje się równolegle w ramach wielu wątków wykonawczych, nie możemy przewidzieć, która operacja zapisu do pamięci będzie ostatnia.
Wyszukiwanie w tabeli haszowej
Kod do wyszukiwania kluczy:
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);
}
}
Aby znaleźć wartość klucza przechowywaną w tabeli, iterujemy przez tablicę, zaczynając od hasza poszukiwanego klucza. W każdej komórce sprawdzamy, czy klucz jest tym, którego szukamy, a jeśli tak, zwracamy jego wartość. Sprawdzamy również, czy klucz jest pusty, a jeśli tak, przerywamy wyszukiwanie.
Jeśli nie uda nam się znaleźć klucza, kod zwraca pustą wartość.
Wszystkie te operacje wyszukiwania mogą być wykonywane równolegle w trakcie wstawień i usunięć. Każda para w tabeli będzie miała dla wątku jeden z czterech powyżej opisanych stanów.
Usuwanie w tabeli haszowej
Kod do usuwania kluczy:
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);
}
}
Usunięcie klucza odbywa się w nietypowy sposób: zostawiamy klucz w tabeli i oznaczamy jego wartość (nie sam klucz) jako pustą. Ten kod jest bardzo podobny do lookup(), z tą różnicą, że w przypadku wykrycia pasowania klucza ustawia jego wartość na pustą.
Jak już wspomniano, gdy klucz zostanie zapisany w slocie, nie jest już przesuwany. Nawet po usunięciu elementu z tabeli, klucz pozostaje na swoim miejscu, a jego wartość staje się tylko pusta. Oznacza to, że nie musimy używać atomowej operacji zapisu wartości slotu, ponieważ nie ma znaczenia, czy bieżąca wartość jest pusta, czy nie — i tak stanie się pusta.
Zwiększanie rozmiaru tabeli haszującej
Zwiększenie rozmiaru tabeli haszującej można osiągnąć poprzez utworzenie większej tabeli i wstawienie do niej niepustych elementów ze starej tabeli. Nie zaimplementowałem tej funkcjonalności, aby zachować kod prostym. Ponadto w programach CUDA, przydzielanie pamięci często odbywa się w kodzie hosta, a nie w jądrze CUDA.
W artykule opisuje, jak modyfikować taką strukturę danych, chronioną przed blokadami.
Konkurencyjność
W przedstawionych powyżej fragmentach kodu, funkcje gpu_hashtable_insert(), _lookup() i _delete() przetwarzają jedną parę klucz-wartość na raz. A poniżej gpu_hashtable_insert(), _lookup() i _delete() przetwarzają tablicę par równolegle, każdą parę w oddzielnym wątku 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);
}
}
Tabela haszująca z ochroną przed blokadami obsługuje równoległe wstawienia, wyszukiwania i usunięcia. Ponieważ pary klucz-wartość zawsze znajdują się w jednym z czterech stanów, a klucze nie są przesuwane, tabela gwarantuje poprawność nawet podczas współczesnego wykonywania różnych rodzajów operacji.
Jednak jeśli równolegle przetwarzamy pakiet wstawek i usunięć, a jeśli w wejściowej tablicy par znajdują się zduplikowane klucze, nie będziemy mogli przewidzieć, które pary ‘wygrają’ — zostaną zapisane w tabeli haszującej ostatnimi. Załóżmy, że wywołaliśmy kod wstawiania z wejściową tablicą par A/0 B/1 A/2 C/3 A/4. Kiedy kod zakończy działanie, pary B/1 i C/3 na pewno będą obecne w tabeli, ale w niej może być dowolna z par A/0, A/2 lub A/4Może to być problem, ale nie musi — wszystko zależy od zastosowania. Możesz z góry wiedzieć, że w wejściowej tablicy nie ma duplikujących się kluczy, albo może nie być dla Ciebie istotne, jaka wartość została zapisana ostatnio.
Jeśli to dla Ciebie problem, musisz rozdzielić duplikujące się pary na różne wywołania systemowe CUDA. W CUDA każda operacja przy wywołaniu jądra kończy się przed kolejnym wywołaniem jądra (przynajmniej w obrębie jednego wątku. W różnych wątkach jądra są wykonywane równolegle). Jeśli w powyższym przykładzie wywołasz jedno jądro z A/0 B/1 A/2 C/3, a drugie z A/4, wtedy klucz A otrzyma wartość 4.
Teraz porozmawiajmy o tym, czy funkcje lookup() i delete() powinny używać prostego (plain) lub zmiennego (volatile) wskaźnika do tablicy par w haszowanej tabeli. stwierdza, że:
Kompilator może według własnego uznania optymalizować operacje odczytu i zapisu w pamięci globalnej lub wspólnej … Te optymalizacje można dezaktywować za pomocą słowa kluczowego
volatile: … każde odniesienie do tej zmiennej jest kompilowane w rzeczywistą instrukcję odczytu lub zapisu w pamięci.
Kwestie poprawności nie wymagają stosowania volatile. Jeśli wątek wykonawczy używa zapisanej wartości z wcześniejszej operacji odczytu, oznacza to, że będzie korzystał z nieco przestarzałych informacji. Mimo to są to informacje w poprawnym stanie haszowanej tabeli w pewnym momencie wywołania jądra. Jeśli potrzebujesz użyć najbardziej aktualnych informacji, możesz zastosować wskaźnik volatile, ale wtedy wydajność nieco spadnie: według moich testów — przy usuwaniu 32 mln elementów prędkość spadła z 500 mln usunięć/s do 450 mln usunięć/s.
Wydajność
W teście z wstawianiem 64 mln elementów i usunięciem 32 mln z nich, rywalizacja pomiędzy std::unordered_map a haszowaną tabelą dla GPU praktycznie nie istnieje:
std::unordered_map poświęciła 70 691 ms na wstawianie i usuwanie elementów z późniejszym zwolnieniem unordered_map (zwolnienie milionów elementów zajmuje sporo czasu, ponieważ w obrębie unordered_map wykonuje się liczne alokacje pamięci). Szczerze mówiąc, w przypadku std::unordered_map zupełnie inne ograniczenia. To jeden wątek CPU, obsługujący klucze-wartości dowolnej wielkości, dobrze działający przy wysokim wykorzystaniu i wykazujący stabilną wydajność po licznych usunięciach.
Czas działania tablicy haszującej dla GPU i interakcji międzyprogramowej wyniósł 984 ms. Wlicza się w to czas potrzebny na umieszczenie tabeli w pamięci i jej usunięcie (jednorazowe przydzielenie 1 GB pamięci, które w CUDA zajmuje pewien czas), wstawianie i usuwanie elementów, a także iterowanie po nich. Uwzględniono również wszystkie kopie do i z pamięci karty graficznej.
Działanie samej tablicy haszującej zajęło 271 ms. Wlicza się w to czas, który karta graficzna poświęciła na wstawienie i usunięcie elementów, nie uwzględniając czasu na kopiowanie do pamięci i iterowanie po powstałej tabeli. Jeśli tabela GPU żyje długo, lub jeśli tablica haszująca znajduje się w całości w pamięci karty graficznej (na przykład do stworzenia tablicy haszującej, która będzie używana przez inny kod GPU, a nie procesor centralny), to wynik testu jest istotny.
Tablica haszująca dla karty graficznej wykazuje wysoką wydajność dzięki dużej przepustowości i aktywnemu równoległemu przetwarzaniu.
Wady
Architektura tablicy haszującej ma kilka problemów, o których należy pamiętać:
- Zgrupowanie utrudnia sondowanie liniowe, co powoduje, że klucze w tabeli nie są rozmieszczone w idealny sposób.
- Klucze nie są usuwane za pomocą funkcji
usuńi z czasem zagracają tablicę.
W rezultacie wydajność tablicy haszującej może stopniowo spadać, szczególnie jeśli istnieje długo, a w niej wykonuje się liczne wstawienia i usunięcia. Jednym ze sposobów łagodzenia tych niedociągnięć jest ponowne haszowanie w nową tabelę z wystarczająco niskim wskaźnikiem wykorzystania i filtracja usuniętych kluczy przy ponownym haszowaniu.
Aby zilustrować opisane problemy, wykorzystam powyższy kod do stworzenia tabeli z 128 milionami elementów, cyklicznie wstawiając 4 miliony elementów, aż wypełnię 124 miliony slotów (wskaźnik wykorzystania około 0,96). Oto tabela wyników, każdy wiersz to wywołanie jądra CUDA z wstawieniem 4 milionów nowych elementów do jednej tablicy haszującej:
Wskaźnik wykorzystania
Czas wstawiania 4 194 304 elementów
0,00
11,608448 ms (361,314798 mln kluczy/sekw.)
0,03
11,751424 ms (356,918799 mln kluczy/sekw.)
0,06
11,942592 ms (351,205515 mln kluczy/sekw.)
0,09
12,081120 ms (347,178429 mln kluczy/sekw.)
0,12
12,242560 ms (342,600233 mln kluczy/sekw.)
0,16
12,396448 ms (338,347235 mln kluczy/sekw.)
0,19
12,533024 ms (334,660176 mln kluczy/sekw.)
0,22
12,703328 ms (330,173626 mln kluczy/sekw.)
0,25
12,884512 ms (325,530693 mln kluczy/sekw.)
0,28
13,033472 ms (321,810182 mln kluczy/sekw.)
0,31
13,239296 ms (316,807174 mln kluczy/sekw.)
0,34
13,392448 ms (313,184256 mln kluczy/sekw.)
0,37
13,624000 ms (307,861434 mln kluczy/sekw.)
0,41
13,875520 ms (302,280855 mln kluczy/sekw.)
0,44
14,126528 ms (296,909756 mln kluczy/sekw.)
0,47
14,399328 ms (291,284699 mln kluczy/sekw.)
0,50
14,690304 ms (285,515123 mln kluczy/sekw.)
0,53
15,039136 ms (278,892623 mln kluczy/sekw.)
0,56
15,478656 ms (270,973402 mln kluczy/sekw.)
0,59
15,985664 ms (262,379092 mln kluczy/sekw.)
0,62
16,668673 ms (251,627968 mln kluczy/sekw.)
0,66
17,587200 ms (238,486174 mln kluczy/sekw.)
0,69
18,690048 ms (224,413765 mln kluczy/sekw.)
0,72
20,278816 ms (206,831789 mln kluczy/sekw.)
0,75
22,545408 ms (186,038058 mln kluczy/sekw.)
0,78
26,053312 ms (160,989275 mln kluczy/sekw.)
0,81
31,895008 ms (131,503463 mln kluczy/sekw.)
0,84
42,103294 ms (99,619378 mln kluczy/sekw.)
0,87
61,849056 ms (67,815164 mln kluczy/sekw.)
0,90
105,695999 ms (39,682713 mln kluczy/sekw.)
0,94
240,204636 ms (17,461378 mln kluczy/sekw.)
Wraz ze wzrostem współczynnika użycia wydajność maleje. Jest to niepożądane w większości przypadków. Jeśli aplikacja wstawia elementy do tabeli, a następnie je odrzuca (na przykład przy liczeniu słów w książce), nie jest to problem. Jednak jeśli aplikacja używa długo żyjącej tabeli mieszającej (na przykład w edytorze graficznym do przechowywania niepustych części obrazów, gdy użytkownik często wstawia i usuwa dane), takie zachowanie może być uciążliwe.
Zmierzono głębokość sondowania tabeli mieszającej po 64 mln wstawień (współczynnik użycia 0,5). Średnia głębokość wyniosła 0,4774, więc większość kluczy znajdowała się w najlepszych możliwych slotach lub w jednym slocie od najlepszej pozycji. Maksymalna głębokość sondowania wyniosła 60.
Następnie zmierzono głębokość sondowania w tabeli po 124 mln wstawień (współczynnik użycia 0,97). Średnia głębokość wyniosła już 10,1757, a maksymalna — 6474 (!!). Wydajność liniowego sondowania znacznie spada przy dużych współczynnikach użycia.
Najlepiej jest utrzymywać niski współczynnik wykorzystania w tej tablicy haszy. Ale wtedy zwiększamy wydajność kosztem pamięci. Na szczęście w przypadku 32-bitowych kluczy i wartości, może to być uzasadnione. Jeśli w powyższym przykładzie w tabeli na 128 milionów elementów utrzymamy współczynnik wykorzystania na poziomie 0,25, to pomieścimy w niej nie więcej niż 32 miliony elementów, a pozostałe 96 milionów miejsc zostanie zmarnowane — po 8 bajtów na każdą parę, 768 MB straconej pamięci.
Zauważ, że mówimy tutaj o utracie pamięci karty graficznej, która jest cenniejszym zasobem niż pamięć systemowa. Chociaż większość nowoczesnych kart graficznych obsługujących CUDA ma co najmniej 4 GB pamięci (w momencie pisania tego artykułu NVIDIA 2080 Ti ma 11 GB), jednak utrata takich ilości nie będzie mądrym rozwiązaniem.
Później szczegółowo opiszę tworzenie tablic haszy dla kart graficznych, które nie mają problemów z głębokością sondowania, oraz sposoby ponownego wykorzystania usuniętych slotów.
Pomiar głębokości sondowania
Aby określić głębokość sondowania klucza, możemy wyciągnąć hasz klucza (jego idealny indeks w tabeli) z jego rzeczywistego indeksu tablicy:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Dzięki magii dwóch liczb binarnych w kodzie uzupełniającym i faktowi, że pojemność tablicy haszy jest potęgą dwójki, podejście to będzie działać nawet wtedy, gdy indeks klucza zostanie przesunięty na początek tabeli. Weźmy klucz, który haszuje się do 1, ale został włożony do slotu 3. Wtedy w tabeli o pojemności 4 otrzymamy (3 — 1) & 3, co odpowiada 2.
Podsumowanie
Jeśli masz pytania lub komentarze, napisz do mnie na lub otwórz nowy temat w.
Ten kod został napisany pod wpływem wspaniałych artykułów:
W przyszłości będę kontynuować pisanie o implementacjach tablic haszy dla kart graficznych i analizować ich wydajność. Planuję chain linking, haszowanie Robina Hooda oraz haszowanie kukułcze z wykorzystaniem operacji atomowych w strukturach danych, które są wygodne dla kart graficznych.
Źródło: habr.com
