Gi-post ko kini sa Github .
Kini usa ka yano nga lamesa sa hash sa GPU nga makahimo sa pagproseso sa gatusan ka milyon nga mga pagsal-ot matag segundo. Sa akong NVIDIA GTX 1060 laptop, ang code nagsal-ot sa 64 milyon nga random nga nahimo nga key-value pairs sa mga 210 ms ug nagtangtang sa 32 milyon nga mga pares sa mga 64 ms.
Sa ato pa, ang katulin sa usa ka laptop gibana-bana nga 300 milyon nga pagsulud / sec ug 500 milyon nga pagtangtang / sec.
Ang lamesa gisulat sa CUDA, bisan kung ang parehas nga teknik mahimong magamit sa HLSL o GLSL. Ang pagpatuman adunay daghang mga limitasyon aron masiguro ang taas nga pasundayag sa usa ka video card:
- Ang 32-bit nga mga yawe ug parehas nga kantidad ang giproseso.
- Ang hash nga lamesa adunay usa ka piho nga gidak-on.
- Ug kini nga gidak-on kinahanglan nga katumbas sa duha sa gahum.
Alang sa mga yawe ug mga kantidad, kinahanglan nimo nga magreserba ug usa ka yano nga marka sa delimiter (sa taas nga code kini 0xffffffff).
Hash nga lamesa nga walay mga kandado
Ang hash nga lamesa naggamit sa bukas nga pag-address sa , sa ato pa, usa lang kini ka han-ay sa key-value pairs nga gitipigan sa memorya ug adunay labaw nga performance sa cache. Ang sama nga dili masulti alang sa pagkadena, nga naglakip sa pagpangita alang sa usa ka pointer sa usa ka nalambigit nga lista. Ang hash table usa ka yano nga array nga nagtipig sa mga elemento KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Ang gidak-on sa lamesa usa ka gahum sa duha, dili usa ka punoan nga numero, tungod kay ang usa ka paspas nga panudlo igo na aron magamit ang maskara sa pow2 / AND, apan ang modulus operator labi ka hinay. Importante kini sa kaso sa linear probing, tungod kay sa linear table lookup ang slot index kinahanglang maputos sa matag slot. Ug ingon nga resulta, ang gasto sa operasyon gidugang modulo sa matag slot.
Ang lamesa nagtipig lamang sa yawe ug bili alang sa matag elemento, dili usa ka hash sa yawe. Tungod kay ang lamesa nagtipig lamang sa 32-bit nga mga yawe, ang hash dali nga kalkulado. Ang code sa ibabaw naggamit sa Murmur3 hash, nga naghimo lamang sa pipila ka mga pagbalhin, XOR ug pagpadaghan.
Ang hash table naggamit ug locking protection techniques nga independente sa memory order. Bisan kung ang pipila nga mga operasyon sa pagsulat makabalda sa han-ay sa ubang mga operasyon, ang lamesa sa hash magpadayon gihapon sa husto nga kahimtang. Atong hisgotan kini sa ubos. Ang teknik maayo kaayo sa mga video card nga nagdagan sa libu-libo nga mga hilo nga dungan.
Ang mga yawe ug mga kantidad sa hash nga lamesa gisugdan aron walay sulod.
Ang code mahimong usbon aron madumala ang 64-bit nga mga yawe ug mga kantidad usab. Ang mga yawe nagkinahanglan og atomic nga pagbasa, pagsulat, ug pagtandi-ug-pagbaylo nga mga operasyon. Ug ang mga kantidad nanginahanglan mga operasyon sa pagbasa ug pagsulat sa atomic. Maayo na lang, sa CUDA, ang read-write nga mga operasyon alang sa 32- ug 64-bit nga mga kantidad kay atomic basta kini natural nga nahiangay (tan-awa sa ubos). ), ug ang modernong mga video card nagsuporta sa 64-bit nga atomic compare-and-exchange nga mga operasyon. Siyempre, kung mobalhin sa 64 bits, gamay ra ang pasundayag.
Hash nga kahimtang sa lamesa
Ang matag key-value pares sa hash table mahimong adunay usa sa upat ka estado:
- Ang yawe ug bili walay sulod. Niini nga estado, ang hash table gisugdan.
- Ang yawe gisulat na, apan ang bili wala pa masulat. Kung ang laing thread karon nagbasa sa datos, kini mobalik nga walay sulod. Normal kini, ang sama nga butang mahitabo kung ang laing hilo sa pagpatay nagtrabaho sa sayo pa, ug naghisgot kami bahin sa usa ka dungan nga istruktura sa datos.
- Ang yawe ug ang bili natala.
- Ang kantidad magamit sa ubang mga hilo sa pagpatuman, apan ang yawe wala pa. Kini mahitabo tungod kay ang CUDA programming model adunay usa ka loosely ordered memory model. Normal kini; sa bisan unsa nga panghitabo, ang yawe walay sulod gihapon, bisan kung ang bili dili na.
Ang usa ka importante nga nuance mao nga sa higayon nga ang yawe gisulat sa slot, kini dili na molihok - bisan kon ang yawe matangtang, kita sa paghisgot mahitungod niini sa ubos.
Ang hash table code bisan sa pagtrabaho uban sa loosely ordered memory models diin ang han-ay diin ang memorya gibasa ug gisulat wala mahibaloi. Sa atong pagtan-aw sa pagsal-ot, pagpangita, ug pagtangtang sa hash table, hinumdumi nga ang matag key-value pair anaa sa usa sa upat ka estado nga gihulagway sa ibabaw.
Pagsulod sa hash table
Ang CUDA function nga nagsal-ot sa key-value pairs ngadto sa hash table ingon niini:
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);
}
}
Aron masulod ang usa ka yawe, ang kodigo mag-usab-usab pinaagi sa hash table array sugod sa hash sa gisulod nga yawe. Ang matag slot sa array nagpahigayon og atomic compare-and-swap operation nga nagtandi sa yawe sa maong slot nga walay sulod. Kung adunay makit-an nga mismatch, ang yawe sa slot i-update sa gisulud nga yawe, ug dayon ang orihinal nga yawe sa slot ibalik. Kung kining orihinal nga yawe walay sulod o gipares sa gisulod nga yawe, nan ang kodigo nakakaplag ug haom nga slot para sa pagsal-ot ug gisulod ang gisulod nga bili ngadto sa slot.
Kung sa usa ka tawag sa kernel gpu_hashtable_insert() adunay daghang mga elemento nga adunay parehas nga yawe, unya ang bisan unsa sa ilang mga kantidad mahimong isulat sa yawe nga slot. Giisip kini nga normal: ang usa sa mga key-value nga pagsulat sa panahon sa tawag magmalampuson, apan tungod kay kining tanan mahitabo nga managsama sulod sa daghang mga hilo sa pagpatay, dili nato matag-an kung unsang memorya ang pagsulat mao ang katapusan.
Hash nga pagpangita sa lamesa
Code para sa pagpangita sa mga yawe:
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);
}
}
Aron makit-an ang bili sa usa ka yawe nga gitipigan sa usa ka lamesa, kita mag-uli sa laray sugod sa hash sa yawe nga atong gipangita. Sa matag slot, among susihon kung ang yawe mao ang among gipangita, ug kung mao, among ibalik ang kantidad niini. Gisusi usab namo kung walay sulod ang yawe, ug kung mao, among gi-abort ang pagpangita.
Kung dili namo makit-an ang yawe, ang code nagbalik sa usa ka walay sulod nga kantidad.
Ang tanan niini nga mga operasyon sa pagpangita mahimo nga dungan nga himuon pinaagi sa mga pagsal-ot ug pagtangtang. Ang matag pares sa lamesa adunay usa sa upat ka estado nga gihulagway sa ibabaw alang sa dagan.
Pagtangtang sa usa ka hash table
Code para sa pagtangtang sa mga yawe:
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);
}
}
Ang pagtangtang sa usa ka yawe gihimo sa usa ka dili kasagaran nga paagi: gibiyaan namon ang yawe sa lamesa ug markahan ang kantidad niini (dili ang yawe mismo) nga walaβy sulod. Kini nga code susama kaayo sa lookup(), gawas nga kung ang usa ka posporo makit-an sa usa ka yawe, gihimo niini nga walay sulod ang kantidad niini.
Sama sa gihisgutan sa ibabaw, sa higayon nga ang usa ka yawe gisulat sa usa ka slot, kini dili na ibalhin. Bisan kung ang usa ka elemento gitangtang gikan sa lamesa, ang yawe nagpabilin sa lugar, ang kantidad niini mahimong walay sulod. Kini nagpasabot nga dili na nato kinahanglan nga mogamit ug atomic write operation alang sa slot value, tungod kay dili igsapayan kung ang kasamtangan nga bili walay sulod o wala - kini mahimo gihapon nga walay sulod.
Pag-usab sa gidak-on sa hash table
Mahimo nimong usbon ang gidak-on sa usa ka hash table pinaagi sa paghimo og mas dako nga lamesa ug pagsal-ot sa mga elemento nga walay sulod gikan sa daan nga lamesa ngadto niini. Wala nako gipatuman kini nga gamit tungod kay gusto nako nga huptan nga simple ang sample code. Dugang pa, sa mga programa sa CUDA, ang alokasyon sa memorya kanunay nga gihimo sa host code kaysa sa CUDA kernel.
Ang artikulo naghulagway kung giunsa pag-usab ang ingon nga istruktura sa datos nga giprotektahan sa kandado.
Pagkakompetisyon
Sa ibabaw nga function code snippet gpu_hashtable_insert(), _lookup() ΠΈ _delete() iproseso ang usa ka key-value pair matag higayon. Ug ubos gpu_hashtable_insert(), _lookup() ΠΈ _delete() iproseso ang usa ka han-ay sa mga pares nga magkaparehas, ang matag pares sa usa ka bulag nga execution thread sa 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);
}
}
Ang lock-resistant hash table nagsuporta sa dungan nga pagsal-ot, pagpangita, ug pagtangtang. Tungod kay ang key-value pairs kanunay anaa sa usa sa upat ka mga estado ug ang mga yawe dili molihok, ang lamesa naggarantiya sa pagkahusto bisan kung lain-laing mga matang sa mga operasyon ang gigamit nga dungan.
Bisan pa, kung magproseso kami usa ka hugpong sa mga pagsal-ot ug pagtangtang nga managsama, ug kung ang input array sa mga pares adunay mga doble nga yawe, nan dili kami makatagna kung unsang mga pares ang "modaog" -isulat sa lamesa sa hash nga katapusan. Ingnon ta nga gitawag nato ang insertion code nga adunay input array of pairs A/0 B/1 A/2 C/3 A/4. Kung makompleto na ang code, ipares B/1 ΠΈ C/3 gigarantiyahan nga naa sa lamesa, apan sa parehas nga oras bisan unsang mga pares ang makita niini A/0, A/2 o A/4. Mahimo kini o dili usa ka problema - kini tanan nagdepende sa aplikasyon. Mahimong nahibal-an nimo nga daan nga walaβy doble nga mga yawe sa input array, o dili nimo igsapayan kung unsang kantidad ang gisulat sa katapusan.
Kung kini usa ka problema alang kanimo, nan kinahanglan nimo nga ibulag ang mga doble nga pares sa lainlaing mga tawag sa sistema sa CUDA. Sa CUDA, ang bisan unsang operasyon nga nagtawag sa kernel kanunay nga makompleto sa dili pa ang sunod nga kernel call (labing menos sulod sa usa ka thread. Sa lain-laing mga thread, ang mga kernel gipatuman sa parallel). Sa pananglitan sa ibabaw, kung tawgon nimo ang usa ka kernel nga adunay A/0 B/1 A/2 C/3, ug ang uban nga adunay A/4, unya ang yawe A makuha ang bili 4.
Karon hisgutan naton kung kinahanglan ba ang mga function lookup() ΠΈ delete() gamita ang usa ka yano o dali nga tudlo sa usa ka han-ay sa mga pares sa hash table. Gipahayag nga:
Mahimong pilion sa compiler nga i-optimize ang mga pagbasa ug pagsulat sa global o gipaambit nga memorya... Kini nga mga pag-optimize mahimong ma-disable gamit ang keyword
volatile: ... bisan unsa nga reperensiya niini nga baryable gihugpong ngadto sa tinuod nga memorya sa pagbasa o pagsulat sa instruksiyon.
Ang mga konsiderasyon sa pagkatul-id wala magkinahanglan og aplikasyon volatile. Kung ang execution thread naggamit sa usa ka cached value gikan sa usa ka naunang read operation, nan kini mogamit og gamay nga outdated nga impormasyon. Apan bisan pa, kini ang kasayuran gikan sa husto nga kahimtang sa lamesa sa hash sa usa ka piho nga oras sa tawag sa kernel. Kung kinahanglan nimo gamiton ang pinakabag-o nga impormasyon, mahimo nimong gamiton ang indeks volatile, apan ang pasundayag mokunhod og gamay: sumala sa akong mga pagsulay, kung gitangtang ang 32 milyon nga mga elemento, ang katulin mikunhod gikan sa 500 milyon nga mga pagtangtang / sec ngadto sa 450 milyon nga mga pagtangtang / sec.
Pag-uswag
Sa pagsulay alang sa pagsal-ot sa 64 milyon nga mga elemento ug pagtangtang sa 32 milyon niini, kompetisyon tali sa std::unordered_map ug halos walay hash table para sa GPU:
std::unordered_map migasto og 70 ms sa pagsal-ot ug pagtangtang sa mga elemento ug dayon pagpalingkawas niini unordered_map (Ang pagtangtang sa milyon-milyon nga mga elemento nanginahanglan daghang oras, tungod kay sa sulod unordered_map daghang mga alokasyon sa memorya ang gihimo). Sa tinuod lang, std:unordered_map hingpit nga lahi nga mga pagdili. Kini usa ka us aka CPU nga hilo sa pagpatuman, nagsuporta sa mga yawe nga kantidad sa bisan unsang gidak-on, maayo nga nahimo sa taas nga rate sa paggamit, ug nagpakita nga lig-on nga pasundayag pagkahuman sa daghang mga pagtangtang.
Ang gidugayon sa hash table alang sa GPU ug inter-program nga komunikasyon mao ang 984 ms. Naglakip kini sa oras nga gigugol sa pagbutang sa lamesa sa panumduman ug pagtangtang niini (paggahin sa 1 GB nga panumduman usa ka higayon, nga magdugay sa CUDA), pagsal-ot ug pagtangtang sa mga elemento, ug pag-uli sa ibabaw niini. Ang tanan nga mga kopya ngadto ug gikan sa memorya sa video card gikonsiderar usab.
Ang hash table mismo mikuha ug 271 ms aron makompleto. Naglakip kini sa oras nga gigugol sa video card sa pagsulod ug pagtangtang sa mga elemento, ug wala gikonsiderar ang oras nga gigugol sa pagkopya sa memorya ug pag-uli sa resulta nga lamesa. Kung ang lamesa sa GPU mabuhi sa dugay nga panahon, o kung ang lamesa sa hash naa sa tibuuk nga panumduman sa video card (pananglitan, paghimo usa ka lamesa nga hash nga gamiton sa ubang GPU code ug dili ang sentral nga processor), nan ang resulta sa pagsulay may kalabutan.
Ang hash table alang sa usa ka video card nagpakita sa taas nga performance tungod sa taas nga throughput ug aktibong parallelization.
mga kakulangan
Ang arkitektura sa hash table adunay pipila ka mga isyu nga kinahanglan mahibal-an:
- Ang linear probing gibabagan sa clustering, nga maoy hinungdan nga ang mga yawe sa lamesa gibutang nga dili kaayo hingpit.
- Ang mga yawe dili makuha gamit ang function
deleteug sa paglabay sa panahon sila nagkalat sa lamesa.
Ingon usa ka sangputanan, ang paghimo sa usa ka lamesa sa hash mahimong hinayhinay nga madaot, labi na kung kini naglungtad sa dugay nga panahon ug adunay daghang mga pagsal-ot ug pagtangtang. Usa ka paagi aron maminusan kini nga mga kakulangan mao ang pag-rehash sa usa ka bag-ong lamesa nga adunay gamay nga rate sa paggamit ug pagsala sa mga gikuha nga yawe sa panahon sa pag-rehashing.
Sa pag-ilustrar sa mga isyu nga gihulagway, akong gamiton ang kodigo sa ibabaw sa paghimo og usa ka lamesa nga adunay 128 ka milyon nga mga elemento ug pag-loop sa 4 ka milyon nga mga elemento hangtud nga napuno ko ang 124 ka milyon nga mga slots (paggamit rate sa mga 0,96). Ania ang lamesa sa resulta, ang matag laray usa ka tawag sa kernel sa CUDA aron isulud ang 4 milyon nga bag-ong mga elemento sa usa ka lamesa sa hash:
Rate sa paggamit
Gidugayon sa pagsal-ot 4 ka elemento
0,00
11,608448 ms (361,314798 milyon nga yawe/seg.)
0,03
11,751424 ms (356,918799 milyon nga yawe/seg.)
0,06
11,942592 ms (351,205515 milyon nga yawe/seg.)
0,09
12,081120 ms (347,178429 milyon nga yawe/seg.)
0,12
12,242560 ms (342,600233 milyon nga yawe/seg.)
0,16
12,396448 ms (338,347235 milyon nga yawe/seg.)
0,19
12,533024 ms (334,660176 milyon nga yawe/seg.)
0,22
12,703328 ms (330,173626 milyon nga yawe/seg.)
0,25
12,884512 ms (325,530693 milyon nga yawe/seg.)
0,28
13,033472 ms (321,810182 milyon nga yawe/seg.)
0,31
13,239296 ms (316,807174 milyon nga yawe/seg.)
0,34
13,392448 ms (313,184256 milyon nga yawe/seg.)
0,37
13,624000 ms (307,861434 milyon nga yawe/seg.)
0,41
13,875520 ms (302,280855 milyon nga yawe/seg.)
0,44
14,126528 ms (296,909756 milyon nga yawe/seg.)
0,47
14,399328 ms (291,284699 milyon nga yawe/seg.)
0,50
14,690304 ms (285,515123 milyon nga yawe/seg.)
0,53
15,039136 ms (278,892623 milyon nga yawe/seg.)
0,56
15,478656 ms (270,973402 milyon nga yawe/seg.)
0,59
15,985664 ms (262,379092 milyon nga yawe/seg.)
0,62
16,668673 ms (251,627968 milyon nga yawe/seg.)
0,66
17,587200 ms (238,486174 milyon nga yawe/seg.)
0,69
18,690048 ms (224,413765 milyon nga yawe/seg.)
0,72
20,278816 ms (206,831789 milyon nga yawe/seg.)
0,75
22,545408 ms (186,038058 milyon nga yawe/seg.)
0,78
26,053312 ms (160,989275 milyon nga yawe/seg.)
0,81
31,895008 ms (131,503463 milyon nga yawe/seg.)
0,84
42,103294 ms (99,619378 milyon nga yawe/seg.)
0,87
61,849056 ms (67,815164 milyon nga yawe/seg.)
0,90
105,695999 ms (39,682713 milyon nga yawe/seg.)
0,94
240,204636 ms (17,461378 milyon nga yawe/seg.)
Samtang nagkadaghan ang paggamit, mikunhod ang performance. Dili kini gusto sa kadaghanan nga mga kaso. Kung ang usa ka aplikasyon magsal-ot sa mga elemento sa usa ka lamesa ug dayon isalikway kini (pananglitan, kung mag-ihap sa mga pulong sa usa ka libro), nan kini dili usa ka problema. Apan kung ang aplikasyon naggamit sa usa ka taas nga kinabuhi nga hash nga lamesa (pananglitan, sa usa ka editor sa graphic aron tipigan ang mga walaβy sulod nga mga bahin sa mga imahe diin ang tiggamit kanunay nga nagsal-ot ug nagtangtang sa kasayuran), nan kini nga pamatasan mahimong adunay problema.
Ug gisukod ang giladmon sa pagsusi sa hash table pagkahuman sa 64 milyon nga mga pagsal-ot (gamit ang hinungdan 0,5). Ang kasagaran nga giladmon mao ang 0,4774, mao nga kadaghanan sa mga yawe anaa sa labing maayo nga posible nga slot o usa ka slot gikan sa pinakamaayo nga posisyon. Ang labing kataas nga giladmon sa tunog mao ang 60.
Gisukod dayon nako ang giladmon sa pagsusi sa usa ka lamesa nga adunay 124 milyon nga mga pagsal-ot (gamit ang hinungdan 0,97). Ang kasagaran nga giladmon mao na ang 10,1757, ug ang pinakataas nga - 6474 (!!). Ang pasundayag sa linear sensing mikunhod pag-ayo sa taas nga rate sa paggamit.
Labing maayo nga huptan nga ubos ang rate sa paggamit sa hash table. Apan unya gipataas namon ang pasundayag sa gasto sa pagkonsumo sa memorya. Maayo na lang, sa kaso sa 32-bit nga mga yawe ug mga kantidad, kini mahimong makatarunganon. Kung sa panig-ingnan sa ibabaw, sa usa ka lamesa nga adunay 128 milyon nga mga elemento, gitipigan namon ang hinungdan sa paggamit nga 0,25, nan mahimo namon nga ibutang ang dili molapas sa 32 milyon nga mga elemento niini, ug ang nahabilin nga 96 milyon nga mga slot mawala - 8 bytes alang sa matag pares. , 768 MB nga nawala nga memorya.
Palihug timan-i nga naghisgot kami mahitungod sa pagkawala sa memorya sa video card, nga usa ka mas bililhon nga kapanguhaan kay sa memorya sa sistema. Bisan kung kadaghanan sa mga modernong desktop graphics card nga nagsuporta sa CUDA adunay labing menos 4 GB nga panumduman (sa panahon sa pagsulat, ang NVIDIA 2080 Ti adunay 11 GB), dili gihapon kini ang labing maalamon nga desisyon nga mawala ang ingon nga mga kantidad.
Sa ulahi ako mosulat og dugang mahitungod sa paghimo og hash nga mga lamesa alang sa mga video card nga walay mga problema sa pagsusi sa giladmon, ingon man usab sa mga paagi sa paggamit pag-usab sa mga natangtang nga mga slot.
Pagsukod sa giladmon sa tingog
Aron mahibal-an ang giladmon sa pagsusi sa usa ka yawe, mahimo naton makuha ang hash sa yawe (ang labing maayo nga indeks sa lamesa) gikan sa aktwal nga indeks sa lamesa:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Tungod sa salamangka sa duha ka komplemento nga binary nga mga numero ug ang kamatuoran nga ang kapasidad sa hash nga lamesa mao ang duha sa gahum sa duha, kini nga pamaagi molihok bisan kung ang yawe nga indeks gibalhin sa sinugdanan sa lamesa. Atong kuhaon ang usa ka yawe nga gihash ngadto sa 1, apan gisal-ut ngadto sa slot 3. Unya alang sa usa ka lamesa nga adunay kapasidad 4 atong makuha (3 β 1) & 3, nga katumbas sa 2.
konklusyon
Kung naa kay pangutana o komento, palihog email nako sa o pagbukas ug bag-ong topiko sa.
Kini nga code gisulat ubos sa inspirasyon gikan sa maayo kaayo nga mga artikulo:
Sa umaabot, magpadayon ako sa pagsulat bahin sa mga pagpatuman sa hash table alang sa mga video card ug pag-analisar sa ilang pasundayag. Ang akong mga plano naglakip sa chaining, Robin Hood hashing, ug cuckoo hashing gamit ang atomic operations sa data structures nga GPU friendly.
Source: www.habr.com