Opublikowałem to na Githubie .
Jest to prosta tablica skrótów GPU zdolna do przetwarzania setek milionów wstawek na sekundę. Na moim laptopie NVIDIA GTX 1060 kod wstawia 64 miliony losowo wygenerowanych par klucz-wartość w ciągu około 210 ms i usuwa 32 miliony par w około 64 ms.
Oznacza to, że prędkość laptopa wynosi około 300 milionów wstawień/s i 500 milionów usunięć/s.
Tabela jest napisana w CUDA, chociaż tę samą technikę można zastosować do HLSL lub GLSL. Implementacja ma kilka ograniczeń, aby zapewnić wysoką wydajność karty graficznej:
- Przetwarzane są tylko klucze 32-bitowe i te same wartości.
- Tabela mieszająca ma stały rozmiar.
- A ten rozmiar musi być równy dwa do potęgi.
Dla kluczy i wartości należy zarezerwować prosty znacznik ogranicznika (w powyższym kodzie jest to 0xffffffff).
Tabela mieszająca bez zamków
Tabela mieszająca używa otwartego adresowania za pomocą oznacza to, że jest to po prostu tablica par klucz-wartość przechowywana w pamięci i charakteryzująca się doskonałą wydajnością pamięci podręcznej. Tego samego nie można powiedzieć o łańcuchowaniu, które polega na szukaniu wskaźnika na połączonej liście. Tablica mieszająca to prosta tablica przechowująca elementy KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Rozmiar tablicy to potęga dwójki, a nie liczba pierwsza, ponieważ do zastosowania maski pow2/AND wystarczy jedna szybka instrukcja, ale operator modułu jest znacznie wolniejszy. Jest to ważne w przypadku sondowania liniowego, ponieważ przy przeszukiwaniu tabeli liniowej indeks szczeliny musi być zawinięty w każdą szczelinę. W rezultacie koszt operacji jest dodawany modulo w każdym gnieździe.
Tabela przechowuje tylko klucz i wartość każdego elementu, a nie skrót klucza. Ponieważ tabela przechowuje tylko klucze 32-bitowe, skrót jest obliczany bardzo szybko. Powyższy kod używa skrótu Murmur3, który wykonuje tylko kilka przesunięć, XOR i mnożeń.
Tabela mieszająca wykorzystuje techniki ochrony przed blokowaniem, które są niezależne od kolejności pamięci. Nawet jeśli niektóre operacje zapisu zaburzą kolejność innych tego typu operacji, tablica mieszająca nadal zachowa prawidłowy stan. Porozmawiamy o tym poniżej. Technika ta doskonale sprawdza się w przypadku kart graficznych obsługujących tysiące wątków jednocześnie.
Klucze i wartości w tabeli mieszającej są inicjowane jako puste.
Kod można zmodyfikować tak, aby obsługiwał również klucze i wartości 64-bitowe. Klucze wymagają niepodzielnych operacji odczytu, zapisu oraz porównania i zamiany. A wartości wymagają atomowych operacji odczytu i zapisu. Na szczęście w CUDA operacje odczytu i zapisu dla wartości 32- i 64-bitowych są niepodzielne, o ile są naturalnie wyrównane (patrz poniżej). ), a nowoczesne karty graficzne obsługują 64-bitowe, atomowe operacje porównywania i wymiany. Oczywiście po przejściu na wersję 64-bitową wydajność nieznacznie spadnie.
Stan tablicy mieszającej
Każda para klucz-wartość w tabeli skrótów może mieć jeden z czterech stanów:
- Klucz i wartość są puste. W tym stanie inicjowana jest tablica mieszająca.
- Klucz został zapisany, ale wartość nie została jeszcze zapisana. Jeśli inny wątek aktualnie odczytuje dane, zwraca wartość pustą. To normalne, to samo by się stało, gdyby nieco wcześniej zadziałał inny wątek wykonania, a mówimy o współbieżnej strukturze danych.
- Rejestrowany jest zarówno klucz, jak i wartość.
- Wartość jest dostępna dla innych wątków wykonania, ale klucz jeszcze nie jest. Może się tak zdarzyć, ponieważ model programowania CUDA ma luźno uporządkowany model pamięci. Jest to normalne; w każdym przypadku klucz jest nadal pusty, nawet jeśli wartość już nie jest.
Ważnym niuansem jest to, że po zapisaniu klucza w gnieździe już się nie porusza - nawet jeśli klucz zostanie usunięty, porozmawiamy o tym poniżej.
Kod tablicy mieszającej działa nawet z luźno uporządkowanymi modelami pamięci, w których kolejność odczytywania i zapisywania pamięci jest nieznana. Kiedy przyglądamy się wstawieniu, wyszukiwaniu i usuwaniu tabeli skrótów, pamiętajmy, że każda para klucz-wartość znajduje się w jednym z czterech stanów opisanych powyżej.
Wstawianie do tablicy mieszającej
Funkcja CUDA, która wstawia pary klucz-wartość do tablicy mieszającej, 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 wykonuje iterację po tablicy tabeli skrótów, zaczynając od skrótu wstawionego klucza. Każde gniazdo w tablicy wykonuje niepodzielną operację porównywania i zamiany, która porównuje klucz w tym gnieździe z pustym. W przypadku wykrycia niezgodności klucz w gnieździe jest aktualizowany o włożony klucz, a następnie zwracany jest oryginalny klucz do gniazda. Jeśli ten oryginalny klucz był pusty lub pasował do wstawionego klucza, wówczas kod znalazł odpowiednie miejsce do wstawienia i wstawił wprowadzoną wartość do gniazda.
Jeśli w jednym wywołaniu jądra gpu_hashtable_insert() istnieje wiele elementów o tym samym kluczu, wówczas dowolną z ich wartości można zapisać w gnieździe klucza. Uważa się to za normalne: jeden z zapisów klucz-wartość podczas wywołania zakończy się sukcesem, ale ponieważ wszystko to dzieje się równolegle w kilku wątkach wykonania, nie możemy przewidzieć, który zapis w pamięci będzie ostatnim.
Wyszukiwanie tabeli mieszającej
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 przechowywanego w tabeli, iterujemy po tablicy, zaczynając od skrótu klucza, którego szukamy. W każdym slocie sprawdzamy czy klucz jest tym którego szukamy i jeśli tak to zwracamy jego wartość. Sprawdzamy również, czy klucz jest pusty i jeśli tak, przerywamy wyszukiwanie.
Jeżeli nie uda nam się znaleźć klucza, kod zwróci pustą wartość.
Wszystkie te operacje wyszukiwania można wykonywać jednocześnie poprzez wstawianie i usuwanie. Każda para w tabeli będzie miała jeden z czterech stanów opisanych powyżej dla przepływu.
Usuwanie w tabeli mieszającej
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);
}
}
Usuwanie klucza odbywa się w nietypowy sposób: zostawiamy klucz w tabeli i zaznaczamy jego wartość (a nie sam klucz) jako pustą. Ten kod jest bardzo podobny do lookup(), z tą różnicą, że gdy w kluczu zostanie znalezione dopasowanie, jego wartość staje się pusta.
Jak wspomniano powyżej, po zapisaniu klucza w gnieździe nie można go już przenosić. Nawet gdy element zostanie usunięty z tabeli, klucz pozostaje na swoim miejscu, jego wartość po prostu staje się pusta. Oznacza to, że nie musimy stosować atomowej operacji zapisu dla wartości slotu, ponieważ nie ma znaczenia, czy bieżąca wartość jest pusta, czy nie – i tak pozostanie pusta.
Zmiana rozmiaru tabeli mieszającej
Możesz zmienić rozmiar tabeli mieszającej, tworząc większą tabelę i wstawiając do niej niepuste elementy ze starej tabeli. Nie zaimplementowałem tej funkcji, ponieważ chciałem, aby przykładowy kod był prosty. Co więcej, w programach CUDA alokacja pamięci często odbywa się w kodzie hosta, a nie w jądrze CUDA.
Artykuł opisuje, jak modyfikować taką strukturę danych chronioną blokadą.
Konkurencyjność
W powyższych fragmentach kodu funkcji gpu_hashtable_insert(), _lookup() и _delete() przetwarzać jedną parę klucz-wartość na raz. I niżej gpu_hashtable_insert(), _lookup() и _delete() przetwarzaj równolegle tablicę par, każdą parę w oddzielnym wątku wykonawczym 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);
}
}
Odporna na blokady tabela skrótów obsługuje jednoczesne wstawianie, przeglądanie i usuwanie. Ponieważ pary klucz-wartość znajdują się zawsze w jednym z czterech stanów, a klucze się nie poruszają, tabela gwarantuje poprawność nawet w przypadku jednoczesnego stosowania różnych typów operacji.
Jeśli jednak równolegle przetworzymy serię wstawień i usunięć, a wejściowa tablica par zawiera zduplikowane klucze, wówczas nie będziemy w stanie przewidzieć, które pary „wygrają” – zostaną zapisane do tablicy mieszającej jako ostatnie. Powiedzmy, że wywołaliśmy kod wstawiania z tablicą wejściową par A/0 B/1 A/2 C/3 A/4. Po zakończeniu kodu następuje parowanie B/1 и C/3 z pewnością będą obecne w tabeli, ale jednocześnie pojawi się w niej dowolna para A/0, A/2 lub A/4. Może to stanowić problem, ale nie musi - wszystko zależy od zastosowania. Możesz wiedzieć z góry, że w tablicy wejściowej nie ma zduplikowanych kluczy, lub możesz nie zwracać uwagi na to, która wartość została zapisana jako ostatnia.
Jeśli stanowi to dla Ciebie problem, musisz rozdzielić zduplikowane pary na różne wywołania systemowe CUDA. W CUDA każda operacja wywołująca jądro zawsze kończy się przed następnym wywołaniem jądra (przynajmniej w obrębie jednego wątku. W różnych wątkach jądra są wykonywane równolegle). W powyższym przykładzie, jeśli wywołasz jedno jądro za pomocą A/0 B/1 A/2 C/3, a drugi z A/4, a następnie klucz A otrzyma wartość 4.
Porozmawiajmy teraz o tym, czy funkcje powinny lookup() и delete() użyj zwykłego lub ulotnego wskaźnika do tablicy par w tablicy mieszającej. Stwierdza, że:
Kompilator może wybrać optymalizację odczytu i zapisu w pamięci globalnej lub współdzielonej... Te optymalizacje można wyłączyć za pomocą słowa kluczowego
volatile: ... wszelkie odniesienia do tej zmiennej są kompilowane w instrukcję odczytu lub zapisu pamięci rzeczywistej.
Względy poprawności nie wymagają zastosowania volatile. Jeśli wątek wykonawczy korzysta z buforowanej wartości z wcześniejszej operacji odczytu, będzie korzystał z nieco nieaktualnych informacji. Ale nadal jest to informacja z prawidłowego stanu tablicy mieszającej w pewnym momencie wywołania jądra. Jeśli chcesz skorzystać z najświeższych informacji, możesz skorzystać z indeksu volatile, ale wtedy wydajność nieznacznie spadnie: według moich testów przy usuwaniu 32 milionów elementów prędkość spadła z 500 milionów usunięć/s do 450 milionów usunięć/s.
produktywność
W teście na wstawienie 64 milionów elementów i usunięcie 32 milionów z nich dochodzi do rywalizacji pomiędzy std::unordered_map i praktycznie nie ma tablicy mieszającej dla procesora graficznego:
std::unordered_map spędził 70 691 ms na wstawianiu i usuwaniu elementów, a następnie ich uwalnianiu unordered_map (pozbycie się milionów elementów zajmuje dużo czasu, bo inside unordered_map dokonywanych jest wielokrotnych alokacji pamięci). Szczerze mówiąc, std:unordered_map zupełnie inne ograniczenia. Jest to wątek wykonawczy pojedynczego procesora, obsługuje pary klucz-wartość dowolnego rozmiaru, działa dobrze przy wysokim stopniu wykorzystania i wykazuje stabilną wydajność po wielokrotnym usunięciu.
Czas trwania tablicy mieszającej dla procesora graficznego i komunikacji między programami wyniósł 984 ms. Obejmuje to czas poświęcony na umieszczenie tabeli w pamięci i jej usunięcie (jednorazowe przydzielenie 1 GB pamięci, co w CUDA zajmuje trochę czasu), wstawianie i usuwanie elementów oraz iterację po nich. Uwzględniane są także wszystkie kopie do i z pamięci karty graficznej.
Ukończenie samej tablicy skrótów zajęło 271 ms. Obejmuje to czas spędzony przez kartę graficzną na wstawianiu i usuwaniu elementów i nie uwzględnia czasu spędzonego na kopiowaniu do pamięci i iterowaniu po wynikowej tabeli. Jeśli tablica GPU działa przez długi czas lub jeśli tablica mieszająca jest w całości zawarta w pamięci karty graficznej (na przykład w celu utworzenia tablicy mieszającej, która będzie używana przez inny kod GPU, a nie centralny procesor), wówczas wynik testu jest istotny.
Tabela skrótów dla karty graficznej wykazuje wysoką wydajność dzięki dużej przepustowości i aktywnej równoległości.
Ograniczenia
Architektura tablicy mieszającej wiąże się z kilkoma problemami, o których należy pamiętać:
- Sondowanie liniowe jest utrudnione przez grupowanie, które powoduje, że klucze w tabeli są rozmieszczone mniej niż idealnie.
- Za pomocą tej funkcji klucze nie są usuwane
deletei z biegiem czasu zaśmiecają stół.
W rezultacie wydajność tabeli mieszającej może stopniowo spadać, szczególnie jeśli istnieje ona przez długi czas i zawiera liczne operacje wstawiania i usuwania. Jednym ze sposobów złagodzenia tych wad jest ponowne utworzenie nowej tabeli o dość niskim współczynniku wykorzystania i odfiltrowanie usuniętych kluczy podczas ponownego mieszania.
Aby zilustrować opisane problemy, użyję powyższego kodu, aby utworzyć tabelę zawierającą 128 milionów elementów i przeglądać 4 miliony elementów, aż zapełnię 124 miliony miejsc (wskaźnik wykorzystania około 0,96). Oto tabela wyników, każdy wiersz jest wywołaniem jądra CUDA w celu wstawienia 4 milionów nowych elementów do jednej tabeli mieszającej:
Wskaźnik użycia
Czas wstawiania 4 194 304 elementów
0,00
11,608448 ms (361,314798 milionów kluczy/s)
0,03
11,751424 ms (356,918799 milionów kluczy/s)
0,06
11,942592 ms (351,205515 milionów kluczy/s)
0,09
12,081120 ms (347,178429 milionów kluczy/s)
0,12
12,242560 ms (342,600233 milionów kluczy/s)
0,16
12,396448 ms (338,347235 milionów kluczy/s)
0,19
12,533024 ms (334,660176 milionów kluczy/s)
0,22
12,703328 ms (330,173626 milionów kluczy/s)
0,25
12,884512 ms (325,530693 milionów kluczy/s)
0,28
13,033472 ms (321,810182 milionów kluczy/s)
0,31
13,239296 ms (316,807174 milionów kluczy/s)
0,34
13,392448 ms (313,184256 milionów kluczy/s)
0,37
13,624000 ms (307,861434 milionów kluczy/s)
0,41
13,875520 ms (302,280855 milionów kluczy/s)
0,44
14,126528 ms (296,909756 milionów kluczy/s)
0,47
14,399328 ms (291,284699 milionów kluczy/s)
0,50
14,690304 ms (285,515123 milionów kluczy/s)
0,53
15,039136 ms (278,892623 milionów kluczy/s)
0,56
15,478656 ms (270,973402 milionów kluczy/s)
0,59
15,985664 ms (262,379092 milionów kluczy/s)
0,62
16,668673 ms (251,627968 milionów kluczy/s)
0,66
17,587200 ms (238,486174 milionów kluczy/s)
0,69
18,690048 ms (224,413765 milionów kluczy/s)
0,72
20,278816 ms (206,831789 milionów kluczy/s)
0,75
22,545408 ms (186,038058 milionów kluczy/s)
0,78
26,053312 ms (160,989275 milionów kluczy/s)
0,81
31,895008 ms (131,503463 milionów kluczy/s)
0,84
42,103294 ms (99,619378 milionów kluczy/s)
0,87
61,849056 ms (67,815164 milionów kluczy/s)
0,90
105,695999 ms (39,682713 milionów kluczy/s)
0,94
240,204636 ms (17,461378 milionów kluczy/s)
Wraz ze wzrostem wykorzystania spada wydajność. W większości przypadków nie jest to pożądane. Jeśli aplikacja wstawia elementy do tabeli, a następnie je odrzuca (na przykład podczas liczenia słów w książce), nie stanowi to problemu. Jeśli jednak aplikacja korzysta z długotrwałej tablicy mieszającej (na przykład w edytorze graficznym do przechowywania niepustych części obrazów, w których użytkownik często wstawia i usuwa informacje), wówczas takie zachowanie może być problematyczne.
I zmierzył głębokość sondowania tabeli mieszającej po 64 milionach wstawek (współczynnik wykorzystania 0,5). Średnia głębokość wynosiła 0,4774, więc większość klawiszy znajdowała się albo w najlepszym możliwym slocie, albo o jedno miejsce od najlepszej pozycji. Maksymalna głębokość sondowania wynosiła 60.
Następnie zmierzyłem głębokość sondowania na stole ze 124 milionami płytek (współczynnik wykorzystania 0,97). Średnia głębokość wynosiła już 10,1757, a maksymalna - 6474 (!!). Wydajność wykrywania liniowego znacznie spada przy wysokim stopniu wykorzystania.
Najlepiej jest utrzymywać niski stopień wykorzystania tej tabeli mieszającej. Ale wtedy zwiększamy wydajność kosztem zużycia pamięci. Na szczęście w przypadku kluczy i wartości 32-bitowych można to uzasadnić. Jeśli w powyższym przykładzie w tabeli zawierającej 128 milionów elementów zachowamy współczynnik wykorzystania 0,25, to możemy umieścić w niej nie więcej niż 32 miliony elementów, a pozostałe 96 milionów miejsc zostanie utracone - po 8 bajtów na parę , 768 MB utraconej pamięci.
Należy pamiętać, że mówimy o utracie pamięci karty graficznej, która jest cenniejszym zasobem niż pamięć systemowa. Chociaż większość nowoczesnych kart graficznych do komputerów stacjonarnych obsługujących CUDA ma co najmniej 4 GB pamięci (w momencie pisania tego tekstu NVIDIA 2080 Ti ma 11 GB), w dalszym ciągu nie byłoby najmądrzejszą decyzją rezygnowanie z takich ilości.
Później napiszę więcej o tworzeniu tablic skrótów dla kart graficznych, które nie mają problemów z głębokością sondowania, a także o sposobach ponownego wykorzystania usuniętych gniazd.
Pomiar głębokości sondowania
Aby określić głębokość sondowania klucza, możemy wyodrębnić skrót klucza (jego idealny indeks tabeli) z jego rzeczywistego indeksu tabeli:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Ze względu na magię liczb binarnych z uzupełnieniem do dwóch i fakt, że pojemność tablicy mieszającej wynosi dwa do potęgi drugiej, to podejście będzie działać nawet wtedy, gdy kluczowy indeks zostanie przeniesiony na początek tabeli. Weźmy klucz, który ma skrót do 1, ale jest włożony do gniazda 3. Następnie dla stołu o pojemności 4 otrzymamy (3 — 1) & 3, co jest równoważne 2.
wniosek
Jeśli masz pytania lub uwagi, napisz do mnie na adres lub otwórz nowy temat w.
Ten kod został napisany zainspirowany świetnymi artykułami:
W przyszłości będę nadal pisać o implementacjach tablic skrótów dla kart graficznych i analizować ich wydajność. Moje plany obejmują łączenie łańcuchowe, mieszanie Robin Hooda i mieszanie z kukułką przy użyciu operacji atomowych w strukturach danych przyjaznych dla GPU.
Źródło: www.habr.com