Github'da yayınladım .
Saniyede yüz milyonlarca eklemeyi işleyebilen basit bir GPU karma tablosudur. NVIDIA GTX 1060 dizüstü bilgisayarımda kod, rastgele oluşturulmuş 64 milyon anahtar/değer çiftini yaklaşık 210 ms'de ekler ve 32 milyon çifti yaklaşık 64 ms'de kaldırır.
Yani bir dizüstü bilgisayardaki hız yaklaşık olarak 300 milyon ekleme/sn ve 500 milyon silme/sn'dir.
Tablo CUDA'da yazılmıştır, ancak aynı teknik HLSL veya GLSL'ye de uygulanabilir. Uygulamanın, video kartında yüksek performans sağlamak için çeşitli sınırlamaları vardır:
- Yalnızca 32 bitlik anahtarlar ve aynı değerler işlenir.
- Hash tablosunun sabit bir boyutu vardır.
- Ve bu büyüklüğün iki üssüne eşit olması gerekir.
Anahtarlar ve değerler için basit bir sınırlayıcı işaretleyici ayırmanız gerekir (yukarıdaki kodda bu 0xffffffff'dir).
Kilitsiz karma tablosu
Karma tablosu açık adreslemeyi kullanır yani, bellekte depolanan ve üstün önbellek performansına sahip olan bir anahtar/değer çiftleri dizisidir. Aynı şey bağlantılı bir listede bir işaretçinin aranmasını içeren zincirleme için söylenemez. Karma tablosu, öğeleri depolayan basit bir dizidir KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Tablonun boyutu asal bir sayı değil, ikinin katıdır, çünkü pow2/AND maskesini uygulamak için hızlı bir komut yeterlidir, ancak modül operatörü çok daha yavaştır. Bu, doğrusal problama durumunda önemlidir, çünkü doğrusal bir tablo aramasında yuva indeksinin her yuvaya sarılması gerekir. Ve sonuç olarak, operasyonun maliyeti her slotta modülo olarak eklenir.
Tablo, anahtarın karmasını değil, yalnızca her öğenin anahtarını ve değerini saklar. Tabloda yalnızca 32 bitlik anahtarlar saklandığından hash çok hızlı hesaplanır. Yukarıdaki kod, yalnızca birkaç kaydırma, XOR ve çarpma işlemi gerçekleştiren Murmur3 karmasını kullanır.
Karma tablosu, bellek sırasından bağımsız kilitleme koruma teknikleri kullanır. Bazı yazma işlemleri bu tür diğer işlemlerin sırasını bozsa bile karma tablo yine de doğru durumu koruyacaktır. Aşağıda bunun hakkında konuşacağız. Bu teknik, binlerce iş parçacığını aynı anda çalıştıran video kartlarıyla harika çalışıyor.
Hash tablosundaki anahtarlar ve değerler boş olarak başlatılmıştır.
Kod, 64 bit anahtarları ve değerleri de işleyecek şekilde değiştirilebilir. Anahtarlar atomik okuma, yazma ve karşılaştırma ve değiştirme işlemlerini gerektirir. Ve değerler atomik okuma ve yazma işlemlerini gerektirir. Neyse ki CUDA'da 32 ve 64 bit değerler için okuma-yazma işlemleri doğal olarak hizalandıkları sürece atomiktir (aşağıya bakın). ) ve modern video kartları 64 bit atomik karşılaştırma ve değişim işlemlerini destekler. Elbette 64 bit'e geçtiğinizde performans biraz düşecektir.
Karma tablo durumu
Bir karma tablosundaki her anahtar/değer çifti dört durumdan birine sahip olabilir:
- Anahtar ve değer boş. Bu durumda karma tablosu başlatılır.
- Anahtar yazıldı ancak değeri henüz yazılmadı. Başka bir iş parçacığı şu anda veri okuyorsa boş döndürür. Bu normaldir, başka bir yürütme iş parçacığı biraz daha erken çalışsaydı aynı şey olurdu ve eşzamanlı bir veri yapısından bahsediyoruz.
- Hem anahtar hem de değer kaydedilir.
- Değer diğer yürütme iş parçacıkları tarafından kullanılabilir, ancak anahtar henüz mevcut değildir. Bunun nedeni CUDA programlama modelinin gevşek sıralı bir bellek modeline sahip olmasıdır. Bu normaldir; her durumda, değer artık öyle olmasa bile anahtar hala boştur.
Önemli bir nüans, anahtar yuvaya yazıldıktan sonra artık hareket etmemesidir - anahtar silinse bile, bunun hakkında aşağıda konuşacağız.
Karma tablo kodu, belleğin okunma ve yazılma sırasının bilinmediği gevşek sıralı bellek modelleriyle bile çalışır. Bir karma tablosunda ekleme, arama ve silme işlemlerine baktığımızda her anahtar/değer çiftinin yukarıda açıklanan dört durumdan birinde olduğunu unutmayın.
Karma tabloya ekleme
Anahtar/değer çiftlerini karma tablosuna ekleyen CUDA işlevi şuna benzer:
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);
}
}
Bir anahtar eklemek için kod, eklenen anahtarın karmasından başlayarak karma tablosu dizisi boyunca yinelenir. Dizideki her yuva, o yuvadaki anahtarı boşla karşılaştıran atomik bir karşılaştırma ve takas işlemi gerçekleştirir. Uyumsuzluk tespit edilirse yuvadaki anahtar, takılan anahtarla güncellenir ve ardından orijinal yuva anahtarı iade edilir. Bu orijinal anahtar boşsa veya eklenen anahtarla eşleşiyorsa, kod, yerleştirme için uygun bir yuva buldu ve girilen değeri yuvaya yerleştirdi.
Bir çekirdek çağrısında ise gpu_hashtable_insert() Aynı anahtara sahip birden fazla eleman varsa, bunların değerlerinden herhangi biri anahtar yuvasına yazılabilir. Bu normal kabul edilir: çağrı sırasında anahtar/değer yazmalarından biri başarılı olur, ancak tüm bunlar birkaç yürütme iş parçacığı içinde paralel olarak gerçekleştiğinden, hangi bellek yazma işleminin sonuncusu olacağını tahmin edemeyiz.
Karma tablo araması
Anahtarları aramak için kod:
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);
}
}
Tabloda saklanan bir anahtarın değerini bulmak için aradığımız anahtarın karmasından başlayarak diziyi yineleriz. Her yuvada aradığımız anahtarın olup olmadığını kontrol ediyoruz ve eğer öyleyse değerini döndürüyoruz. Ayrıca anahtarın boş olup olmadığını da kontrol ediyoruz ve boşsa aramayı iptal ediyoruz.
Anahtarı bulamazsak kod boş bir değer döndürür.
Bu arama işlemlerinin tümü ekleme ve silme işlemleriyle eş zamanlı olarak gerçekleştirilebilir. Tablodaki her çift, akış için yukarıda açıklanan dört durumdan birine sahip olacaktır.
Karma tablosunda silme
Anahtarları silme kodu:
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);
}
}
Bir anahtarın silinmesi alışılmadık bir şekilde yapılır: anahtarı tabloda bırakırız ve değerini (anahtarın kendisini değil) boş olarak işaretleriz. Bu kod şuna çok benzer: lookup()ancak bir anahtarda bir eşleşme bulunduğunda değerini boş bırakır.
Yukarıda belirtildiği gibi, bir anahtar bir yuvaya yazıldığında artık hareket etmez. Bir öğe tablodan silindiğinde bile anahtar yerinde kalır, değeri boşalır. Bu, yuva değeri için atomik yazma işlemi kullanmamıza gerek olmadığı anlamına gelir, çünkü mevcut değerin boş olup olmaması önemli değildir; yine de boş olacaktır.
Karma tablosunu yeniden boyutlandırma
Daha büyük bir tablo oluşturarak ve eski tablodan boş olmayan öğeleri bu tabloya ekleyerek karma tablosunun boyutunu değiştirebilirsiniz. Örnek kodu basit tutmak istediğim için bu işlevi uygulamadım. Ayrıca CUDA programlarında bellek tahsisi genellikle CUDA çekirdeği yerine ana bilgisayar kodunda yapılır.
makale böyle bir kilit korumalı veri yapısının nasıl değiştirileceğini açıklamaktadır.
Rekabet gücü
Yukarıdaki işlev kodu parçacıklarında gpu_hashtable_insert(), _lookup() и _delete() Aynı anda bir anahtar/değer çiftini işleyin. Ve daha aşağıda gpu_hashtable_insert(), _lookup() и _delete() her çift ayrı bir GPU yürütme iş parçacığında olacak şekilde bir dizi çifti paralel olarak işleyin:
// 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);
}
}
Kilitlenmeye dayanıklı karma tablosu eşzamanlı eklemeleri, aramaları ve silmeleri destekler. Anahtar/değer çiftleri her zaman dört durumdan birinde olduğundan ve anahtarlar hareket etmediğinden, farklı türdeki işlemler aynı anda kullanıldığında bile tablo doğruluğu garanti eder.
Bununla birlikte, bir dizi ekleme ve silme işlemini paralel olarak işlersek ve giriş çiftleri dizisi yinelenen anahtarlar içeriyorsa, o zaman hangi çiftlerin "kazanacağını" tahmin edemeyiz; karma tablosuna en son yazılacaktır. Diyelim ki ekleme kodunu çiftlerden oluşan bir giriş dizisiyle çağırdık A/0 B/1 A/2 C/3 A/4. Kod tamamlandığında çiftler B/1 и C/3 tabloda bulunması garanti edilir, ancak aynı zamanda çiftlerden herhangi biri de tabloda görünecektir A/0, A/2 veya A/4. Bu bir sorun olabilir veya olmayabilir; her şey uygulamaya bağlıdır. Giriş dizisinde yinelenen anahtarların olmadığını önceden biliyor olabilirsiniz veya en son hangi değerin yazıldığını önemsemiyor olabilirsiniz.
Bu sizin için bir sorunsa, yinelenen çiftleri farklı CUDA sistem çağrılarına ayırmanız gerekir. CUDA'da çekirdeği çağıran herhangi bir işlem her zaman bir sonraki çekirdek çağrısından önce tamamlanır (en azından bir iş parçacığı içinde. Farklı iş parçacıklarında çekirdekler paralel olarak yürütülür). Yukarıdaki örnekte, bir çekirdeği şununla çağırırsanız: A/0 B/1 A/2 C/3, ve diğeri ile A/4, ardından anahtar A değerini alacak 4.
Şimdi fonksiyonların gerekli olup olmadığı hakkında konuşalım. lookup() и delete() karma tablosundaki bir çift dizisine yönelik düz veya uçucu bir işaretçi kullanın. şunları belirtir:
Derleyici, okuma ve yazma işlemlerini genel veya paylaşılan belleğe optimize etmeyi seçebilir... Bu optimizasyonlar, anahtar kelime kullanılarak devre dışı bırakılabilir.
volatile: ... bu değişkene yapılan herhangi bir referans, gerçek bir hafıza okuma veya yazma talimatına derlenir.
Doğruluk hususları başvuru gerektirmez volatile. Yürütme iş parçacığı daha önceki bir okuma işleminden önbelleğe alınmış bir değer kullanıyorsa, bu durumda biraz güncel olmayan bilgileri kullanıyor olacaktır. Ancak yine de bu, çekirdek çağrısının belirli bir anında karma tablosunun doğru durumundan gelen bilgidir. En son bilgileri kullanmanız gerekiyorsa dizini kullanabilirsiniz. volatile, ancak bu durumda performans biraz düşecek: Testlerime göre 32 milyon öğeyi silerken hız 500 milyon silme/sn'den 450 milyon silme/sn'ye düştü.
Proizvoditelnost
64 milyon elementin yerleştirilip 32 milyonunun silindiği testte aralarında rekabet yaşandı. std::unordered_map ve GPU için neredeyse hiç karma tablosu yok:
std::unordered_map Öğeleri ekleyip çıkarmak ve ardından serbest bırakmak için 70 ms harcadı unordered_map (Milyonlarca elementten kurtulmak çok zaman alır çünkü içeride unordered_map birden fazla bellek tahsisi yapılır). Dürüstçe söylüyorum, std:unordered_map tamamen farklı kısıtlamalar. Tek bir CPU yürütme iş parçacığıdır, her boyuttaki anahtar/değer çiftlerini destekler, yüksek kullanım oranlarında iyi performans gösterir ve birden fazla silme işleminden sonra istikrarlı performans gösterir.
GPU ve programlar arası iletişim için hash tablosunun süresi 984 ms idi. Bu, tablonun belleğe yerleştirilmesi ve silinmesi (bir defada 1 GB bellek tahsis edilmesi, CUDA'da biraz zaman alır), öğelerin eklenmesi, silinmesi ve bunların üzerinde yineleme yapılması için harcanan zamanı içerir. Video kartı belleğine gelen ve gelen tüm kopyalar da dikkate alınır.
Hash tablosunun tamamlanması 271 ms sürdü. Bu, video kartının öğeleri takmak ve silmek için harcadığı süreyi içerir ve belleğe kopyalama ve sonuç tablosu üzerinde yineleme için harcanan süreyi hesaba katmaz. GPU tablosu uzun süre dayanıyorsa veya karma tablosu tamamen video kartının belleğinde bulunuyorsa (örneğin, merkezi işlemci tarafından değil başka GPU kodu tarafından kullanılacak bir karma tablosu oluşturmak için), o zaman test sonucu önemlidir.
Bir video kartının karma tablosu, yüksek verim ve aktif paralelleştirme nedeniyle yüksek performansı gösterir.
Dezavantajları:
Karma tablo mimarisinde dikkat edilmesi gereken birkaç sorun vardır:
- Doğrusal problama, kümeleme nedeniyle sekteye uğrar, bu da tablodaki anahtarların mükemmelden daha az yerleştirilmesine neden olur.
- İşlev kullanılarak anahtarlar kaldırılmaz
deleteve zamanla masayı doldururlar.
Sonuç olarak, bir karma tablosunun performansı, özellikle uzun süredir mevcutsa ve çok sayıda ekleme ve silme işlemi içeriyorsa, yavaş yavaş düşebilir. Bu dezavantajları azaltmanın bir yolu, oldukça düşük bir kullanım oranına sahip yeni bir tabloya yeniden karma yapmak ve yeniden karma sırasında kaldırılan anahtarları filtrelemektir.
Açıklanan sorunları göstermek için, 128 milyon öğeli bir tablo oluşturmak ve 4 milyon alanı doldurana kadar (kullanım oranı yaklaşık 124) 0,96 milyon öğe arasında döngü yapmak için yukarıdaki kodu kullanacağım. İşte sonuç tablosu, her satır, bir karma tablosuna 4 milyon yeni öğe eklemek için bir CUDA çekirdek çağrısıdır:
Kullanım oranı
Ekleme süresi 4 öğe
0,00
11,608448 ms (361,314798 milyon anahtar/sn.)
0,03
11,751424 ms (356,918799 milyon anahtar/sn.)
0,06
11,942592 ms (351,205515 milyon anahtar/sn.)
0,09
12,081120 ms (347,178429 milyon anahtar/sn.)
0,12
12,242560 ms (342,600233 milyon anahtar/sn.)
0,16
12,396448 ms (338,347235 milyon anahtar/sn.)
0,19
12,533024 ms (334,660176 milyon anahtar/sn.)
0,22
12,703328 ms (330,173626 milyon anahtar/sn.)
0,25
12,884512 ms (325,530693 milyon anahtar/sn.)
0,28
13,033472 ms (321,810182 milyon anahtar/sn.)
0,31
13,239296 ms (316,807174 milyon anahtar/sn.)
0,34
13,392448 ms (313,184256 milyon anahtar/sn.)
0,37
13,624000 ms (307,861434 milyon anahtar/sn.)
0,41
13,875520 ms (302,280855 milyon anahtar/sn.)
0,44
14,126528 ms (296,909756 milyon anahtar/sn.)
0,47
14,399328 ms (291,284699 milyon anahtar/sn.)
0,50
14,690304 ms (285,515123 milyon anahtar/sn.)
0,53
15,039136 ms (278,892623 milyon anahtar/sn.)
0,56
15,478656 ms (270,973402 milyon anahtar/sn.)
0,59
15,985664 ms (262,379092 milyon anahtar/sn.)
0,62
16,668673 ms (251,627968 milyon anahtar/sn.)
0,66
17,587200 ms (238,486174 milyon anahtar/sn.)
0,69
18,690048 ms (224,413765 milyon anahtar/sn.)
0,72
20,278816 ms (206,831789 milyon anahtar/sn.)
0,75
22,545408 ms (186,038058 milyon anahtar/sn.)
0,78
26,053312 ms (160,989275 milyon anahtar/sn.)
0,81
31,895008 ms (131,503463 milyon anahtar/sn.)
0,84
42,103294 ms (99,619378 milyon anahtar/sn.)
0,87
61,849056 ms (67,815164 milyon anahtar/sn.)
0,90
105,695999 ms (39,682713 milyon anahtar/sn.)
0,94
240,204636 ms (17,461378 milyon anahtar/sn.)
Kullanım arttıkça performans düşer. Bu çoğu durumda arzu edilen bir durum değildir. Bir uygulama bir tabloya öğeler ekleyip sonra bunları atarsa (örneğin, bir kitaptaki sözcükleri sayarken), bu bir sorun değildir. Ancak uygulama uzun ömürlü bir karma tablo kullanıyorsa (örneğin, kullanıcının sık sık bilgi ekleyip sildiği görüntülerin boş olmayan bölümlerini depolamak için bir grafik düzenleyicide), bu davranış sorunlu olabilir.
Ve 64 milyon eklemeden sonra karma tablosunun derinliklerini ölçtü (kullanım faktörü 0,5). Ortalama derinlik 0,4774'tü, dolayısıyla çoğu anahtar ya mümkün olan en iyi yuvadaydı ya da en iyi konumdan bir yuva uzaktaydı. Maksimum sondaj derinliği 60 idi.
Daha sonra prob derinliğini 124 milyon kesici uç içeren bir masa üzerinde ölçtüm (kullanım faktörü 0,97). Ortalama derinlik zaten 10,1757 idi ve maksimum - 6474 (!!). Doğrusal algılama performansı yüksek kullanım oranlarında önemli ölçüde düşer.
Bu karma tablonun kullanım oranını düşük tutmak en iyisidir. Ancak daha sonra bellek tüketimi pahasına performansı artırıyoruz. Neyse ki, 32 bitlik anahtarlar ve değerler söz konusu olduğunda bu haklı gösterilebilir. Yukarıdaki örnekte 128 milyon öğeli bir tabloda kullanım faktörünü 0,25 tutarsak, o zaman içine 32 milyondan fazla öğe yerleştiremeyiz ve kalan 96 milyon yuva kaybolacaktır - her çift için 8 bayt 768 MB kayıp hafıza.
Sistem belleğinden daha değerli bir kaynak olan ekran kartı belleğinin kaybından bahsettiğimizi lütfen unutmayın. Her ne kadar CUDA'yı destekleyen modern masaüstü grafik kartlarının çoğu en az 4 GB belleğe sahip olsa da (bu yazının yazıldığı sırada NVIDIA 2080 Ti'de 11 GB vardı), bu miktarları kaybetmek yine de en akıllıca karar olmayacaktır.
Daha sonra derinlik sorunu olmayan video kartları için karma tabloları oluşturmanın yanı sıra silinen yuvaları yeniden kullanma yolları hakkında daha fazla yazacağım.
Sondaj derinliği ölçümü
Bir anahtarın araştırma derinliğini belirlemek için, anahtarın karmasını (ideal tablo dizini) gerçek tablo dizininden çıkarabiliriz:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
İkinin ikinin tümleyeni ikili sayıların büyüsü ve karma tablosunun kapasitesinin iki üzeri iki olması gerçeği nedeniyle, bu yaklaşım anahtar dizini tablonun başına taşındığında bile işe yarayacaktır. 1'e hashlenmiş ancak 3. yuvaya yerleştirilmiş bir anahtarı alalım. Daha sonra 4 kapasiteli bir tablo için şunu elde ederiz: (3 — 1) & 32'ye eşdeğerdir.
Sonuç
Sorularınız veya yorumlarınız varsa lütfen bana e-posta gönderin: veya yeni bir konu açın.
Bu kod mükemmel makalelerden ilham alınarak yazılmıştır:
İlerleyen zamanlarda ekran kartlarına hash table uygulamaları yazmaya ve performanslarını analiz etmeye devam edeceğim. Planlarım arasında GPU dostu veri yapılarında atomik işlemleri kullanarak zincirleme, Robin Hood hashing ve guguklu hashing yer alıyor.
Kaynak: habr.com