Saya mempostingnya di Github .
Ini adalah tabel hash GPU sederhana yang mampu memproses ratusan juta sisipan per detik. Di laptop NVIDIA GTX 1060 saya, kode tersebut memasukkan 64 juta pasangan nilai kunci yang dibuat secara acak dalam waktu sekitar 210 mdtk dan menghapus 32 juta pasangan dalam waktu sekitar 64 mdtk.
Artinya, kecepatan pada laptop kurang lebih 300 juta insert/detik dan 500 juta hapus/detik.
Tabel ini ditulis dalam CUDA, meskipun teknik yang sama dapat diterapkan pada HLSL atau GLSL. Implementasinya memiliki beberapa keterbatasan untuk memastikan kinerja tinggi pada kartu video:
- Hanya kunci 32-bit dan nilai yang sama yang diproses.
- Tabel hash memiliki ukuran tetap.
- Dan ukuran ini harus sama dengan dua pangkat.
Untuk kunci dan nilai, Anda perlu memesan penanda pembatas sederhana (dalam kode di atas ini adalah 0xffffffff).
Tabel hash tanpa kunci
Tabel hash menggunakan pengalamatan terbuka dengan , artinya, ini hanyalah serangkaian pasangan nilai kunci yang disimpan dalam memori dan memiliki kinerja cache yang unggul. Hal yang sama tidak berlaku untuk chaining, yang melibatkan pencarian pointer dalam daftar tertaut. Tabel hash adalah array sederhana yang menyimpan elemen KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ukuran tabelnya adalah pangkat dua, bukan bilangan prima, karena satu instruksi cepat sudah cukup untuk menerapkan topeng pow2/AND, namun operator modulusnya jauh lebih lambat. Hal ini penting dalam kasus pemeriksaan linier, karena dalam pencarian tabel linier, indeks slot harus dibungkus dalam setiap slot. Dan alhasil, biaya pengoperasiannya bertambah modulo di setiap slot.
Tabel hanya menyimpan kunci dan nilai untuk setiap elemen, bukan hash dari kunci tersebut. Karena tabel hanya menyimpan kunci 32-bit, hash dihitung dengan sangat cepat. Kode di atas menggunakan hash Murmur3 yang hanya melakukan beberapa shift, XOR, dan perkalian.
Tabel hash menggunakan teknik perlindungan penguncian yang tidak bergantung pada urutan memori. Bahkan jika beberapa operasi tulis mengganggu urutan operasi lainnya, tabel hash akan tetap mempertahankan keadaan yang benar. Kami akan membicarakannya di bawah. Teknik ini bekerja sangat baik dengan kartu video yang menjalankan ribuan thread secara bersamaan.
Kunci dan nilai dalam tabel hash diinisialisasi menjadi kosong.
Kode dapat dimodifikasi untuk menangani kunci dan nilai 64-bit juga. Kunci memerlukan operasi baca, tulis, dan bandingkan-dan-swap atomik. Dan nilai memerlukan operasi baca dan tulis atom. Untungnya, di CUDA, operasi baca-tulis untuk nilai 32 dan 64-bit bersifat atomik selama keduanya selaras secara alami (lihat di bawah). ), dan kartu video modern mendukung operasi perbandingan dan pertukaran atom 64-bit. Tentu saja jika berpindah ke 64 bit, performanya akan sedikit menurun.
Status tabel hash
Setiap pasangan nilai kunci dalam tabel hash dapat memiliki salah satu dari empat status:
- Kunci dan nilai kosong. Dalam keadaan ini, tabel hash diinisialisasi.
- Kuncinya sudah tertulis, tapi nilainya belum tertulis. Jika thread lain sedang membaca data, thread tersebut akan kembali dalam keadaan kosong. Ini normal, hal yang sama akan terjadi jika rangkaian eksekusi lain bekerja sedikit lebih awal, dan kita berbicara tentang struktur data bersamaan.
- Kunci dan nilainya dicatat.
- Nilai tersedia untuk rangkaian eksekusi lainnya, namun kuncinya belum. Hal ini dapat terjadi karena model pemrograman CUDA memiliki model memori yang diurutkan secara longgar. Hal ini normal; dalam keadaan apa pun, kuncinya tetap kosong, meskipun nilainya sudah tidak kosong lagi.
Nuansa penting adalah bahwa setelah kunci ditulis ke slot, kunci tidak lagi bergerak - bahkan jika kunci dihapus, kita akan membicarakannya di bawah.
Kode tabel hash bahkan berfungsi dengan model memori yang diurutkan secara longgar di mana urutan pembacaan dan penulisan memori tidak diketahui. Saat kita melihat penyisipan, pencarian, dan penghapusan dalam tabel hash, ingatlah bahwa setiap pasangan nilai kunci berada di salah satu dari empat keadaan yang dijelaskan di atas.
Memasukkan ke dalam tabel hash
Fungsi CUDA yang menyisipkan pasangan nilai kunci ke dalam tabel hash terlihat seperti ini:
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);
}
}
Untuk memasukkan kunci, kode melakukan iterasi melalui array tabel hash yang dimulai dengan hash dari kunci yang dimasukkan. Setiap slot dalam array melakukan operasi perbandingan dan pertukaran atomik yang membandingkan kunci dalam slot tersebut dengan yang kosong. Jika ketidakcocokan terdeteksi, kunci dalam slot diperbarui dengan kunci yang dimasukkan, dan kemudian kunci slot asli dikembalikan. Jika kunci asli ini kosong atau cocok dengan kunci yang dimasukkan, maka kode menemukan slot yang cocok untuk dimasukkan dan memasukkan nilai yang dimasukkan ke dalam slot tersebut.
Jika dalam satu panggilan kernel gpu_hashtable_insert() ada beberapa elemen dengan kunci yang sama, maka nilainya apa pun dapat ditulis ke slot kunci. Hal ini dianggap normal: salah satu penulisan nilai kunci selama panggilan akan berhasil, tetapi karena semua ini terjadi secara paralel dalam beberapa rangkaian eksekusi, kami tidak dapat memprediksi penulisan memori mana yang akan menjadi yang terakhir.
Pencarian tabel hash
Kode untuk mencari kunci:
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);
}
}
Untuk menemukan nilai kunci yang disimpan dalam tabel, kita melakukan iterasi melalui array yang dimulai dengan hash dari kunci yang kita cari. Di setiap slot, kami memeriksa apakah kuncinya adalah yang kami cari, dan jika demikian, kami mengembalikan nilainya. Kami juga memeriksa apakah kuncinya kosong, dan jika demikian, kami membatalkan pencarian.
Jika kita tidak dapat menemukan kuncinya, kode mengembalikan nilai kosong.
Semua operasi pencarian ini dapat dilakukan secara bersamaan melalui penyisipan dan penghapusan. Setiap pasangan dalam tabel akan memiliki salah satu dari empat keadaan yang dijelaskan di atas untuk alirannya.
Menghapus di tabel hash
Kode untuk menghapus kunci:
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);
}
}
Menghapus kunci dilakukan dengan cara yang tidak biasa: kita membiarkan kunci di tabel dan menandai nilainya (bukan kunci itu sendiri) sebagai kosong. Kode ini sangat mirip dengan lookup(), kecuali ketika kecocokan ditemukan pada kunci, nilainya akan kosong.
Seperti disebutkan di atas, setelah kunci ditulis ke slot, kunci tersebut tidak lagi dipindahkan. Bahkan ketika sebuah elemen dihapus dari tabel, kuncinya tetap di tempatnya, nilainya menjadi kosong. Artinya kita tidak perlu menggunakan operasi penulisan atom untuk nilai slot, karena tidak masalah apakah nilai saat ini kosong atau tidak - nilai tersebut akan tetap kosong.
Mengubah ukuran tabel hash
Anda dapat mengubah ukuran tabel hash dengan membuat tabel yang lebih besar dan memasukkan elemen tidak kosong dari tabel lama ke dalamnya. Saya tidak menerapkan fungsi ini karena saya ingin membuat kode contoh tetap sederhana. Terlebih lagi, dalam program CUDA, alokasi memori sering kali dilakukan pada kode host daripada di kernel CUDA.
Artikel menjelaskan cara memodifikasi struktur data yang dilindungi kunci.
Daya saing
Dalam cuplikan kode fungsi di atas gpu_hashtable_insert(), _lookup() ΠΈ _delete() memproses satu pasangan nilai kunci pada satu waktu. Dan lebih rendah gpu_hashtable_insert(), _lookup() ΠΈ _delete() memproses serangkaian pasangan secara paralel, masing-masing pasangan dalam thread eksekusi GPU terpisah:
// 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);
}
}
Tabel hash tahan kunci mendukung penyisipan, pencarian, dan penghapusan secara bersamaan. Karena pasangan kunci-nilai selalu berada di salah satu dari empat keadaan dan kunci tidak berpindah, tabel menjamin kebenarannya bahkan ketika berbagai jenis operasi digunakan secara bersamaan.
Namun, jika kita memproses sekumpulan penyisipan dan penghapusan secara paralel, dan jika array masukan dari pasangan berisi kunci duplikat, maka kita tidak akan dapat memprediksi pasangan mana yang akan βmenangββakan ditulis ke tabel hash terakhir. Katakanlah kita memanggil kode penyisipan dengan array masukan yang berpasangan A/0 B/1 A/2 C/3 A/4. Saat kode selesai, pasangkan B/1 ΠΈ C/3 dijamin ada di tabel, tetapi pada saat yang sama pasangan mana pun akan muncul di dalamnya A/0, A/2 ΠΈΠ»ΠΈ A/4. Ini mungkin menjadi masalah atau tidak - semuanya tergantung pada aplikasinya. Anda mungkin mengetahui sebelumnya bahwa tidak ada kunci duplikat dalam larik masukan, atau Anda mungkin tidak peduli nilai mana yang ditulis terakhir.
Jika ini merupakan masalah bagi Anda, maka Anda perlu memisahkan pasangan duplikat menjadi panggilan sistem CUDA yang berbeda. Di CUDA, setiap operasi yang memanggil kernel selalu selesai sebelum panggilan kernel berikutnya (setidaknya dalam satu thread. Di thread yang berbeda, kernel dieksekusi secara paralel). Pada contoh di atas, jika Anda memanggil satu kernel dengan A/0 B/1 A/2 C/3, dan yang lainnya dengan A/4, lalu kuncinya A akan mendapatkan nilainya 4.
Sekarang mari kita bicara tentang apakah fungsi seharusnya lookup() ΠΈ delete() gunakan pointer biasa atau mudah menguap ke array pasangan di tabel hash. Menyatakan bahwa:
Kompiler dapat memilih untuk mengoptimalkan pembacaan dan penulisan ke memori global atau bersama... Pengoptimalan ini dapat dinonaktifkan menggunakan kata kunci
volatile: ... referensi apa pun ke variabel ini dikompilasi ke dalam instruksi baca atau tulis memori nyata.
Pertimbangan kebenaran tidak memerlukan penerapan volatile. Jika thread eksekusi menggunakan nilai cache dari operasi pembacaan sebelumnya, maka thread tersebut akan menggunakan informasi yang sedikit ketinggalan jaman. Namun tetap saja, ini adalah informasi dari keadaan tabel hash yang benar pada saat tertentu dari panggilan kernel. Jika Anda perlu menggunakan informasi terbaru, Anda dapat menggunakan indeks volatile, namun performanya akan sedikit menurun: menurut pengujian saya, saat menghapus 32 juta elemen, kecepatannya menurun dari 500 juta penghapusan/detik menjadi 450 juta penghapusan/detik.
Performa
Dalam pengujian untuk memasukkan 64 juta elemen dan menghapus 32 juta di antaranya, persaingan antar std::unordered_map dan hampir tidak ada tabel hash untuk GPU:
std::unordered_map menghabiskan 70 ms untuk memasukkan dan menghapus elemen lalu membebaskannya unordered_map (menyingkirkan jutaan elemen membutuhkan banyak waktu, karena di dalam unordered_map beberapa alokasi memori dibuat). Sejujurnya, std:unordered_map pembatasan yang sangat berbeda. Ini adalah thread eksekusi CPU tunggal, mendukung nilai kunci dengan ukuran berapa pun, berkinerja baik pada tingkat pemanfaatan tinggi, dan menunjukkan kinerja stabil setelah beberapa kali penghapusan.
Durasi tabel hash untuk GPU dan komunikasi antar program adalah 984 ms. Ini termasuk waktu yang dihabiskan untuk menempatkan tabel di memori dan menghapusnya (mengalokasikan 1 GB memori satu kali, yang memerlukan beberapa waktu di CUDA), memasukkan dan menghapus elemen, dan mengulanginya. Semua salinan ke dan dari memori kartu video juga diperhitungkan.
Tabel hash itu sendiri membutuhkan waktu 271 ms untuk diselesaikan. Ini termasuk waktu yang dihabiskan oleh kartu video untuk memasukkan dan menghapus elemen, dan tidak memperhitungkan waktu yang dihabiskan untuk menyalin ke memori dan mengulangi tabel yang dihasilkan. Jika tabel GPU bertahan lama, atau jika tabel hash seluruhnya terdapat di memori kartu video (misalnya, untuk membuat tabel hash yang akan digunakan oleh kode GPU lain dan bukan prosesor pusat), maka hasil tesnya relevan.
Tabel hash untuk kartu video menunjukkan kinerja tinggi karena throughput tinggi dan paralelisasi aktif.
Kekurangan:
Arsitektur tabel hash memiliki beberapa masalah yang perlu diperhatikan:
- Penyelidikan linier terhambat oleh pengelompokan, yang menyebabkan penempatan kunci dalam tabel kurang sempurna.
- Kunci tidak dilepas menggunakan fungsi ini
deletedan seiring waktu mereka mengacaukan meja.
Akibatnya, kinerja tabel hash dapat menurun secara bertahap, terutama jika tabel tersebut sudah ada dalam jangka waktu lama dan memiliki banyak sisipan dan penghapusan. Salah satu cara untuk mengurangi kelemahan ini adalah dengan mengulangi tabel baru dengan tingkat pemanfaatan yang cukup rendah dan memfilter kunci yang dihapus selama pengulangan.
Untuk mengilustrasikan masalah yang dijelaskan, saya akan menggunakan kode di atas untuk membuat tabel dengan 128 juta elemen dan mengulang 4 juta elemen hingga saya mengisi 124 juta slot (tingkat pemanfaatan sekitar 0,96). Berikut adalah tabel hasil, setiap baris merupakan panggilan kernel CUDA untuk memasukkan 4 juta elemen baru ke dalam satu tabel hash:
Tingkat penggunaan
Durasi penyisipan 4 elemen
0,00
11,608448 ms (361,314798 juta kunci/detik)
0,03
11,751424 ms (356,918799 juta kunci/detik)
0,06
11,942592 ms (351,205515 juta kunci/detik)
0,09
12,081120 ms (347,178429 juta kunci/detik)
0,12
12,242560 ms (342,600233 juta kunci/detik)
0,16
12,396448 ms (338,347235 juta kunci/detik)
0,19
12,533024 ms (334,660176 juta kunci/detik)
0,22
12,703328 ms (330,173626 juta kunci/detik)
0,25
12,884512 ms (325,530693 juta kunci/detik)
0,28
13,033472 ms (321,810182 juta kunci/detik)
0,31
13,239296 ms (316,807174 juta kunci/detik)
0,34
13,392448 ms (313,184256 juta kunci/detik)
0,37
13,624000 ms (307,861434 juta kunci/detik)
0,41
13,875520 ms (302,280855 juta kunci/detik)
0,44
14,126528 ms (296,909756 juta kunci/detik)
0,47
14,399328 ms (291,284699 juta kunci/detik)
0,50
14,690304 ms (285,515123 juta kunci/detik)
0,53
15,039136 ms (278,892623 juta kunci/detik)
0,56
15,478656 ms (270,973402 juta kunci/detik)
0,59
15,985664 ms (262,379092 juta kunci/detik)
0,62
16,668673 ms (251,627968 juta kunci/detik)
0,66
17,587200 ms (238,486174 juta kunci/detik)
0,69
18,690048 ms (224,413765 juta kunci/detik)
0,72
20,278816 ms (206,831789 juta kunci/detik)
0,75
22,545408 ms (186,038058 juta kunci/detik)
0,78
26,053312 ms (160,989275 juta kunci/detik)
0,81
31,895008 ms (131,503463 juta kunci/detik)
0,84
42,103294 ms (99,619378 juta kunci/detik)
0,87
61,849056 ms (67,815164 juta kunci/detik)
0,90
105,695999 ms (39,682713 juta kunci/detik)
0,94
240,204636 ms (17,461378 juta kunci/detik)
Ketika pemanfaatan meningkat, kinerja menurun. Hal ini tidak diinginkan dalam banyak kasus. Jika suatu aplikasi menyisipkan elemen ke dalam tabel dan kemudian membuangnya (misalnya, saat menghitung kata dalam buku), maka hal ini tidak menjadi masalah. Namun jika aplikasi menggunakan tabel hash yang berumur panjang (misalnya, dalam editor grafis untuk menyimpan bagian gambar yang tidak kosong di mana pengguna sering menyisipkan dan menghapus informasi), maka perilaku ini dapat menimbulkan masalah.
Dan mengukur kedalaman penyelidikan tabel hash setelah 64 juta penyisipan (faktor pemanfaatan 0,5). Kedalaman rata-rata adalah 0,4774, sehingga sebagian besar kunci berada di slot terbaik atau satu slot jauhnya dari posisi terbaik. Kedalaman bunyi maksimum adalah 60.
Saya kemudian mengukur kedalaman probing di atas meja dengan 124 juta sisipan (faktor pemanfaatan 0,97). Kedalaman rata-rata sudah 10,1757, dan maksimum - 6474 (!!). Performa penginderaan linier turun secara signifikan pada tingkat pemanfaatan yang tinggi.
Yang terbaik adalah menjaga tingkat pemanfaatan tabel hash ini tetap rendah. Namun kemudian kami meningkatkan kinerja dengan mengorbankan konsumsi memori. Untungnya, dalam kasus kunci dan nilai 32-bit, hal ini dapat dibenarkan. Jika dalam contoh di atas, dalam tabel dengan 128 juta elemen, kita mempertahankan faktor pemanfaatan 0,25, maka kita dapat menempatkan tidak lebih dari 32 juta elemen di dalamnya, dan 96 juta slot sisanya akan hilang - 8 byte untuk setiap pasangan , 768 MB memori hilang.
Harap dicatat bahwa kita berbicara tentang hilangnya memori kartu video, yang merupakan sumber daya yang lebih berharga daripada memori sistem. Meskipun sebagian besar kartu grafis desktop modern yang mendukung CUDA memiliki memori minimal 4 GB (pada saat penulisan, NVIDIA 2080 Ti memiliki 11 GB), kehilangan jumlah tersebut bukanlah keputusan yang paling bijaksana.
Nanti saya akan menulis lebih banyak tentang membuat tabel hash untuk kartu video yang tidak memiliki masalah kedalaman probing, serta cara menggunakan kembali slot yang dihapus.
Pengukuran kedalaman bunyi
Untuk menentukan kedalaman probing suatu kunci, kita dapat mengekstrak hash kunci tersebut (indeks tabel idealnya) dari indeks tabel sebenarnya:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Karena keajaiban bilangan biner komplemen dua dan fakta bahwa kapasitas tabel hash adalah dua pangkat dua, pendekatan ini akan berhasil bahkan ketika indeks kunci dipindahkan ke awal tabel. Mari kita ambil kunci yang di-hash ke 1, tetapi dimasukkan ke slot 3. Kemudian untuk tabel dengan kapasitas 4 kita dapatkan (3 β 1) & 3, yang setara dengan 2.
Kesimpulan
Jika Anda memiliki pertanyaan atau komentar, silakan kirim email kepada saya di atau buka topik baru di.
Kode ini ditulis berdasarkan inspirasi dari artikel bagus:
Di masa depan, saya akan terus menulis tentang implementasi tabel hash untuk kartu video dan menganalisis kinerjanya. Rencana saya mencakup chaining, hashing Robin Hood, dan hashing cuckoo menggunakan operasi atom dalam struktur data yang ramah GPU.
Sumber: www.habr.com