Fe'i postiais ar Github .
Mae'n dabl hash GPU syml sy'n gallu prosesu cannoedd o filiynau o fewnosodiadau yr eiliad. Ar fy ngliniadur NVIDIA GTX 1060, mae'r cod yn mewnosod 64 miliwn o barau gwerth allweddol a gynhyrchir ar hap mewn tua 210 ms ac yn dileu 32 miliwn o barau mewn tua 64 ms.
Hynny yw, y cyflymder ar liniadur yw tua 300 miliwn o fewnosodiadau/eiliad a 500 miliwn o ddileadau/eiliad.
Mae'r tabl wedi'i ysgrifennu mewn CUDA, er y gellir cymhwyso'r un dechneg i HLSL neu GLSL. Mae gan y gweithrediad nifer o gyfyngiadau i sicrhau perfformiad uchel ar gerdyn fideo:
- Dim ond allweddi 32-did a'r un gwerthoedd sy'n cael eu prosesu.
- Mae gan y bwrdd hash faint sefydlog.
- Ac mae'n rhaid i'r maint hwn fod yn hafal i ddau i'r pŵer.
Ar gyfer allweddi a gwerthoedd, mae angen i chi gadw marciwr amffinydd syml (yn y cod uchod dyma 0xffffffff).
Bwrdd hash heb gloeon
Mae'r tabl hash yn defnyddio cyfeiriadau agored gyda , hynny yw, yn syml, mae'n amrywiaeth o barau gwerth allweddol sy'n cael eu storio yn y cof ac sydd â pherfformiad cache uwch. Ni ellir dweud yr un peth am gadwyno, sy'n golygu chwilio am bwyntydd mewn rhestr gysylltiedig. Mae tabl stwnsh yn arae syml sy'n storio elfennau KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Mae maint y bwrdd yn bŵer o ddau, nid rhif cysefin, oherwydd mae un cyfarwyddyd cyflym yn ddigon i gymhwyso'r mwgwd pow2 / AND, ond mae gweithredwr y modwlws yn llawer arafach. Mae hyn yn bwysig yn achos stilio llinol, oherwydd mewn chwiliad tabl llinol rhaid lapio'r mynegai slot ym mhob slot. Ac o ganlyniad, mae cost y llawdriniaeth yn cael ei ychwanegu modwlo ym mhob slot.
Mae'r tabl yn storio'r allwedd a'r gwerth ar gyfer pob elfen yn unig, nid stwnsh o'r allwedd. Gan fod y tabl yn storio allweddi 32-did yn unig, mae'r hash yn cael ei gyfrifo'n gyflym iawn. Mae'r cod uchod yn defnyddio'r hash Murmur3, sydd ond yn perfformio ychydig o sifftiau, XORs a lluosi.
Mae'r tabl hash yn defnyddio technegau amddiffyn cloi sy'n annibynnol ar drefn cof. Hyd yn oed os bydd rhai gweithrediadau ysgrifennu yn amharu ar drefn gweithrediadau eraill o'r fath, bydd y tabl hash yn dal i gynnal y cyflwr cywir. Byddwn yn siarad am hyn isod. Mae'r dechneg yn gweithio'n wych gyda chardiau fideo sy'n rhedeg miloedd o edafedd ar yr un pryd.
Mae'r allweddi a'r gwerthoedd yn y tabl hash yn cael eu cychwyn i wagio.
Gellir addasu'r cod i drin allweddi a gwerthoedd 64-bit hefyd. Mae allweddi yn gofyn am weithrediadau darllen, ysgrifennu, a chymharu-a-cyfnewid atomig. Ac mae gwerthoedd yn gofyn am weithrediadau darllen ac ysgrifennu atomig. Yn ffodus, yn CUDA, mae gweithrediadau darllen-ysgrifennu ar gyfer gwerthoedd 32- a 64-bit yn atomig cyn belled â'u bod wedi'u halinio'n naturiol (gweler isod). ), ac mae cardiau fideo modern yn cefnogi gweithrediadau cymharu a chyfnewid atomig 64-did. Wrth gwrs, wrth symud i 64 did, bydd perfformiad yn gostwng ychydig.
Cyflwr tabl hash
Gall pob pâr gwerth allweddol mewn tabl stwnsh fod ag un o bedwar cyflwr:
- Mae allwedd a gwerth yn wag. Yn y cyflwr hwn, mae'r tabl hash yn cael ei gychwyn.
- Mae'r allwedd wedi'i ysgrifennu, ond nid yw'r gwerth wedi'i ysgrifennu eto. Os yw edefyn arall yn darllen data ar hyn o bryd, mae'n dychwelyd yn wag. Mae hyn yn normal, byddai'r un peth wedi digwydd pe bai llinyn gweithredu arall wedi gweithio ychydig yn gynharach, ac rydym yn sôn am strwythur data cydamserol.
- Mae'r allwedd a'r gwerth yn cael eu cofnodi.
- Mae'r gwerth ar gael i edafedd cyflawni eraill, ond nid yw'r allwedd eto. Gall hyn ddigwydd oherwydd bod gan fodel rhaglennu CUDA fodel cof sydd wedi'i drefnu'n llac. Mae hyn yn normal; beth bynnag, mae'r allwedd yn dal yn wag, hyd yn oed os nad yw'r gwerth bellach.
Naws bwysig yw, unwaith y bydd yr allwedd wedi'i ysgrifennu i'r slot, nid yw'n symud mwyach - hyd yn oed os caiff yr allwedd ei ddileu, byddwn yn siarad am hyn isod.
Mae'r cod tabl hash hyd yn oed yn gweithio gyda modelau cof sydd wedi'u trefnu'n llac lle nad yw'r drefn y caiff cof ei ddarllen a'i ysgrifennu yn hysbys. Wrth i ni edrych ar fewnosod, chwilio, a dileu mewn tabl hash, cofiwch fod pob pâr gwerth allweddol yn un o'r pedwar cyflwr a ddisgrifir uchod.
Mewnosod mewn bwrdd hash
Mae'r swyddogaeth CUDA sy'n mewnosod parau gwerth allweddol mewn tabl stwnsh yn edrych fel hyn:
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);
}
}
I fewnosod allwedd, mae'r cod yn ailadrodd trwy'r gyfres bwrdd stwnsh gan ddechrau gyda stwnsh yr allwedd a fewnosodwyd. Mae pob slot yn yr arae yn perfformio gweithrediad cymharu-a-cyfnewid atomig sy'n cymharu'r allwedd yn y slot hwnnw â gwag. Os canfyddir diffyg cyfatebiaeth, caiff yr allwedd yn y slot ei diweddaru gyda'r allwedd a fewnosodwyd, ac yna dychwelir yr allwedd slot wreiddiol. Os oedd yr allwedd wreiddiol hon yn wag neu'n cyfateb i'r allwedd a fewnosodwyd, yna daeth y cod o hyd i slot addas i'w fewnosod a gosododd y gwerth a fewnosodwyd yn y slot.
Os mewn un galwad cnewyllyn gpu_hashtable_insert() mae yna elfennau lluosog gyda'r un allwedd, yna gellir ysgrifennu unrhyw un o'u gwerthoedd i'r slot allweddol. Ystyrir hyn yn normal: bydd un o'r ysgrifeniadau gwerth allweddol yn ystod yr alwad yn llwyddo, ond gan fod hyn i gyd yn digwydd ochr yn ochr â sawl llinyn gweithredu, ni allwn ragweld pa ysgrifen cof fydd yr un olaf.
Chwilio tabl hash
Cod ar gyfer chwilio bysellau:
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);
}
}
I ddarganfod gwerth allwedd sydd wedi'i storio mewn tabl, rydyn ni'n ailadrodd trwy'r arae gan ddechrau gyda stwnsh yr allwedd rydyn ni'n edrych amdano. Ym mhob slot, rydyn ni'n gwirio ai'r allwedd yw'r un rydyn ni'n edrych amdano, ac os felly, rydyn ni'n dychwelyd ei werth. Rydym hefyd yn gwirio a yw'r allwedd yn wag, ac os felly, rydym yn rhoi'r gorau i'r chwiliad.
Os na allwn ddod o hyd i'r allwedd, mae'r cod yn dychwelyd gwerth gwag.
Gellir cyflawni'r holl weithrediadau chwilio hyn ar yr un pryd trwy fewnosod a dileu. Bydd gan bob pâr yn y tabl un o'r pedwar cyflwr a ddisgrifir uchod ar gyfer y llif.
Dileu mewn tabl hash
Cod ar gyfer dileu allweddi:
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);
}
}
Mae dileu allwedd yn cael ei wneud mewn ffordd anarferol: rydyn ni'n gadael yr allwedd yn y tabl ac yn nodi ei werth (nid yr allwedd ei hun) yn wag. Mae'r cod hwn yn debyg iawn i lookup(), ac eithrio pan ddarganfyddir cyfatebiaeth ar allwedd, mae'n gwneud ei werth yn wag.
Fel y soniwyd uchod, unwaith y bydd allwedd wedi'i ysgrifennu i slot, ni chaiff ei symud mwyach. Hyd yn oed pan fydd elfen yn cael ei dileu o'r tabl, mae'r allwedd yn parhau yn ei le, mae ei werth yn dod yn wag. Mae hyn yn golygu nad oes angen i ni ddefnyddio gweithrediad ysgrifennu atomig ar gyfer y gwerth slot, oherwydd nid oes ots a yw'r gwerth presennol yn wag ai peidio - bydd yn dal yn wag.
Newid maint bwrdd stwnsh
Gallwch newid maint bwrdd stwnsh trwy greu tabl mwy a mewnosod elfennau nad ydynt yn wag o'r hen fwrdd ynddo. Wnes i ddim gweithredu'r swyddogaeth hon oherwydd roeddwn i eisiau cadw'r cod sampl yn syml. Ar ben hynny, mewn rhaglenni CUDA, mae dyraniad cof yn aml yn cael ei wneud yn y cod gwesteiwr yn hytrach nag yn y cnewyllyn CUDA.
Yn yr erthygl yn disgrifio sut i addasu strwythur data o'r fath a ddiogelir gan glo.
Cystadleurwydd
Yn y pytiau cod swyddogaeth uchod gpu_hashtable_insert(), _lookup() и _delete() prosesu un pâr gwerth allweddol ar y tro. Ac yn is gpu_hashtable_insert(), _lookup() и _delete() prosesu amrywiaeth o barau yn gyfochrog, pob pâr mewn llinyn gweithredu GPU ar wahân:
// 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);
}
}
Mae'r bwrdd stwnsh sy'n gwrthsefyll clo yn cefnogi mewnosodiadau, edrychiadau a dileadau cydamserol. Oherwydd bod parau gwerth allweddol bob amser mewn un o bedwar cyflwr ac nad yw'r allweddi'n symud, mae'r tabl yn gwarantu cywirdeb hyd yn oed pan ddefnyddir gwahanol fathau o weithrediadau ar yr un pryd.
Fodd bynnag, os byddwn yn prosesu swp o fewnosodiadau a dileadau yn gyfochrog, ac os yw'r arae mewnbwn o barau yn cynnwys allweddi dyblyg, yna ni fyddwn yn gallu rhagweld pa barau fydd yn “ennill” - a fydd yn cael ei ysgrifennu i'r tabl hash olaf. Gadewch i ni ddweud ein bod wedi galw'r cod mewnosod gydag amrywiaeth mewnbwn o barau A/0 B/1 A/2 C/3 A/4. Pan fydd y cod wedi'i gwblhau, parau B/1 и C/3 yn sicr o fod yn bresennol yn y tabl, ond ar yr un pryd bydd unrhyw un o'r parau yn ymddangos ynddo A/0, A/2 neu A/4. Gall hyn fod yn broblem neu beidio - mae'r cyfan yn dibynnu ar y cais. Efallai eich bod yn gwybod ymlaen llaw nad oes allweddi dyblyg yn yr arae mewnbwn, neu efallai nad oes ots gennych pa werth a ysgrifennwyd ddiwethaf.
Os yw hyn yn broblem i chi, yna mae angen i chi wahanu'r parau dyblyg i wahanol alwadau system CUDA. Yn CUDA, mae unrhyw weithrediad sy'n galw'r cnewyllyn bob amser yn dod i ben cyn yr alwad cnewyllyn nesaf (o leiaf o fewn un edefyn. Mewn edafedd gwahanol, gweithredir cnewyllyn yn gyfochrog). Yn yr enghraifft uchod, os byddwch yn ffonio un cnewyllyn gyda A/0 B/1 A/2 C/3, a'r llall gyda A/4, yna yr allwedd A bydd yn cael y gwerth 4.
Nawr gadewch i ni siarad a ddylai swyddogaethau lookup() и delete() defnyddio pwyntydd plaen neu gyfnewidiol i amrywiaeth o barau yn y tabl stwnsh. Yn datgan bod:
Gall y casglwr ddewis optimeiddio darlleniadau ac ysgrifennu i gof byd-eang neu gof a rennir... Gellir analluogi'r optimeiddiadau hyn gan ddefnyddio'r allweddair
volatile: ... mae unrhyw gyfeiriad at y newidyn hwn yn cael ei grynhoi mewn cyfarwyddyd darllen neu ysgrifennu cof go iawn.
Nid oes angen cymhwyso ystyriaethau cywirdeb volatile. Os yw'r edefyn gweithredu yn defnyddio gwerth wedi'i storio o weithrediad a ddarllenwyd yn gynharach, yna bydd yn defnyddio gwybodaeth sydd ychydig yn hen ffasiwn. Ond o hyd, mae hyn yn wybodaeth o gyflwr cywir y tabl hash ar adeg benodol o'r alwad cnewyllyn. Os oes angen i chi ddefnyddio'r wybodaeth ddiweddaraf, gallwch ddefnyddio'r mynegai volatile, ond yna bydd y perfformiad yn gostwng ychydig: yn ôl fy mhrofion, wrth ddileu 32 miliwn o elfennau, gostyngodd y cyflymder o 500 miliwn o ddileadau / eiliad i 450 miliwn o ddileadau / eiliad.
Cynhyrchiant
Yn y prawf ar gyfer mewnosod 64 miliwn o elfennau a dileu 32 miliwn ohonynt, cystadleuaeth rhwng std::unordered_map ac nid oes fawr ddim bwrdd hash ar gyfer y GPU:
std::unordered_map gwario 70 ms yn mewnosod a thynnu elfennau ac yna'n eu rhyddhau unordered_map (Mae cael gwared ar filiynau o elfennau yn cymryd llawer o amser, oherwydd y tu mewn unordered_map gwneir dyraniadau cof lluosog). A siarad yn onest, std:unordered_map cyfyngiadau hollol wahanol. Mae'n edefyn gweithredu CPU sengl, yn cefnogi gwerthoedd allweddol o unrhyw faint, yn perfformio'n dda ar gyfraddau defnyddio uchel, ac yn dangos perfformiad sefydlog ar ôl dileu lluosog.
Hyd y tabl hash ar gyfer y GPU a chyfathrebu rhwng rhaglenni oedd 984 ms. Mae hyn yn cynnwys yr amser a dreulir yn gosod y bwrdd yn y cof a'i ddileu (dyrannu 1 GB o gof un tro, sy'n cymryd peth amser yn CUDA), mewnosod a dileu elfennau, ac iteru drostynt. Mae pob copi i ac o gof y cerdyn fideo hefyd yn cael eu hystyried.
Cymerodd y tabl hash ei hun 271 ms i'w gwblhau. Mae hyn yn cynnwys yr amser a dreulir gan y cerdyn fideo yn mewnosod a dileu elfennau, ac nid yw'n ystyried yr amser a dreulir yn copïo i'r cof ac yn ailadrodd dros y tabl canlyniadol. Os yw'r tabl GPU yn byw am amser hir, neu os yw'r tabl hash wedi'i gynnwys yn gyfan gwbl yng nghof y cerdyn fideo (er enghraifft, i greu tabl hash a fydd yn cael ei ddefnyddio gan god GPU arall ac nid y prosesydd canolog), yna mae canlyniad y prawf yn berthnasol.
Mae'r tabl hash ar gyfer cerdyn fideo yn dangos perfformiad uchel oherwydd trwybwn uchel a chyfochrogrwydd gweithredol.
Cyfyngiadau
Mae gan bensaernïaeth bwrdd hash rai problemau i fod yn ymwybodol ohonynt:
- Mae stilio llinellol yn cael ei rwystro gan glystyru, sy'n achosi i'r allweddi yn y tabl gael eu gosod yn llai na pherffaith.
- Nid yw allweddi yn cael eu tynnu gan ddefnyddio'r ffwythiant
deletea thros amser maent yn annibendod y bwrdd.
O ganlyniad, gall perfformiad tabl hash ddiraddio'n raddol, yn enwedig os yw'n bodoli am amser hir a bod ganddo nifer o fewnosodiadau a dileu. Un ffordd o liniaru'r anfanteision hyn yw ail-wampio i mewn i dabl newydd gyda chyfradd defnyddio eithaf isel a hidlo'r allweddi a dynnwyd allan yn ystod yr ail-wampio.
I ddarlunio'r materion a ddisgrifiwyd, byddaf yn defnyddio'r cod uchod i greu tabl gyda 128 miliwn o elfennau a dolen trwy 4 miliwn o elfennau nes i mi lenwi 124 miliwn o slotiau (cyfradd defnyddio o tua 0,96). Dyma'r tabl canlyniadau, mae pob rhes yn alwad cnewyllyn CUDA i fewnosod 4 miliwn o elfennau newydd mewn un tabl stwnsh:
Cyfradd defnydd
Hyd y mewnosodiad 4 o elfennau
0,00
11,608448 ms (361,314798 miliwn o allweddi/eiliad)
0,03
11,751424 ms (356,918799 miliwn o allweddi/eiliad)
0,06
11,942592 ms (351,205515 miliwn o allweddi/eiliad)
0,09
12,081120 ms (347,178429 miliwn o allweddi/eiliad)
0,12
12,242560 ms (342,600233 miliwn o allweddi/eiliad)
0,16
12,396448 ms (338,347235 miliwn o allweddi/eiliad)
0,19
12,533024 ms (334,660176 miliwn o allweddi/eiliad)
0,22
12,703328 ms (330,173626 miliwn o allweddi/eiliad)
0,25
12,884512 ms (325,530693 miliwn o allweddi/eiliad)
0,28
13,033472 ms (321,810182 miliwn o allweddi/eiliad)
0,31
13,239296 ms (316,807174 miliwn o allweddi/eiliad)
0,34
13,392448 ms (313,184256 miliwn o allweddi/eiliad)
0,37
13,624000 ms (307,861434 miliwn o allweddi/eiliad)
0,41
13,875520 ms (302,280855 miliwn o allweddi/eiliad)
0,44
14,126528 ms (296,909756 miliwn o allweddi/eiliad)
0,47
14,399328 ms (291,284699 miliwn o allweddi/eiliad)
0,50
14,690304 ms (285,515123 miliwn o allweddi/eiliad)
0,53
15,039136 ms (278,892623 miliwn o allweddi/eiliad)
0,56
15,478656 ms (270,973402 miliwn o allweddi/eiliad)
0,59
15,985664 ms (262,379092 miliwn o allweddi/eiliad)
0,62
16,668673 ms (251,627968 miliwn o allweddi/eiliad)
0,66
17,587200 ms (238,486174 miliwn o allweddi/eiliad)
0,69
18,690048 ms (224,413765 miliwn o allweddi/eiliad)
0,72
20,278816 ms (206,831789 miliwn o allweddi/eiliad)
0,75
22,545408 ms (186,038058 miliwn o allweddi/eiliad)
0,78
26,053312 ms (160,989275 miliwn o allweddi/eiliad)
0,81
31,895008 ms (131,503463 miliwn o allweddi/eiliad)
0,84
42,103294 ms (99,619378 miliwn o allweddi/eiliad)
0,87
61,849056 ms (67,815164 miliwn o allweddi/eiliad)
0,90
105,695999 ms (39,682713 miliwn o allweddi/eiliad)
0,94
240,204636 ms (17,461378 miliwn o allweddi/eiliad)
Wrth i'r defnydd gynyddu, mae perfformiad yn lleihau. Nid yw hyn yn ddymunol yn y rhan fwyaf o achosion. Os yw cymhwysiad yn mewnosod elfennau mewn tabl ac yna'n eu taflu (er enghraifft, wrth gyfrif geiriau mewn llyfr), yna nid yw hyn yn broblem. Ond os yw'r rhaglen yn defnyddio tabl stwnsh hirhoedlog (er enghraifft, mewn golygydd graffeg i storio rhannau o ddelweddau nad ydynt yn wag lle mae'r defnyddiwr yn aml yn mewnosod ac yn dileu gwybodaeth), yna gall yr ymddygiad hwn fod yn broblemus.
A mesur dyfnder stilio tabl hash ar ôl 64 miliwn o fewnosodiadau (ffactor defnydd 0,5). Y dyfnder cyfartalog oedd 0,4774, felly roedd y rhan fwyaf o allweddi naill ai yn y slot gorau posibl neu un slot i ffwrdd o'r safle gorau. Y dyfnder sain uchaf oedd 60.
Yna mesurais y dyfnder stilio ar fwrdd gyda 124 miliwn o fewnosodiadau (ffactor defnydd 0,97). Roedd y dyfnder cyfartalog eisoes yn 10,1757, a'r uchafswm - 6474 (!!). Mae perfformiad synhwyro llinol yn gostwng yn sylweddol ar gyfraddau defnyddio uchel.
Mae'n well cadw cyfradd defnyddio'r tabl hash hwn yn isel. Ond yna rydym yn cynyddu perfformiad ar draul defnydd cof. Yn ffodus, yn achos allweddi a gwerthoedd 32-did, gellir cyfiawnhau hyn. Os yn yr enghraifft uchod, mewn tabl gyda 128 miliwn o elfennau, rydym yn cadw'r ffactor defnydd o 0,25, yna ni allwn osod mwy na 32 miliwn o elfennau ynddo, a bydd y 96 miliwn o slotiau sy'n weddill yn cael eu colli - 8 bytes ar gyfer pob pâr. , 768 MB o gof coll.
Sylwch ein bod yn sôn am golli cof cerdyn fideo, sy'n adnodd mwy gwerthfawr na chof system. Er bod gan y mwyafrif o gardiau graffeg bwrdd gwaith modern sy'n cefnogi CUDA o leiaf 4 GB o gof (ar adeg ysgrifennu, mae gan yr NVIDIA 2080 Ti 11 GB), nid dyma'r penderfyniad doethaf o hyd i golli symiau o'r fath.
Yn ddiweddarach byddaf yn ysgrifennu mwy am greu tablau hash ar gyfer cardiau fideo nad ydynt yn cael problemau gyda dyfnder stilio, yn ogystal â ffyrdd o ailddefnyddio slotiau wedi'u dileu.
Mesur dyfnder swnio
Er mwyn pennu dyfnder treiddgar allwedd, gallwn dynnu hash yr allwedd (ei fynegai tabl delfrydol) o'i fynegai tabl gwirioneddol:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Oherwydd hud dau o rifau deuaidd ategu dau a'r ffaith bod cynhwysedd y tabl hash yn ddau i bŵer dau, bydd y dull hwn yn gweithio hyd yn oed pan fydd y mynegai allweddol yn cael ei symud i ddechrau'r tabl. Gadewch i ni gymryd allwedd sy'n stwnsio i 1, ond yn cael ei fewnosod yn slot 3. Yna ar gyfer bwrdd gyda chynhwysedd 4 rydym yn cael (3 — 1) & 3, sy'n cyfateb i 2.
Casgliad
Os oes gennych gwestiynau neu sylwadau, anfonwch e-bost ataf yn neu agor pwnc newydd i mewn.
Ysgrifennwyd y cod hwn dan ysbrydoliaeth o erthyglau rhagorol:
Yn y dyfodol, byddaf yn parhau i ysgrifennu am weithrediadau tabl hash ar gyfer cardiau fideo a dadansoddi eu perfformiad. Mae fy nghynlluniau'n cynnwys cadwyno, stwnsio Robin Hood, a stwnsio gog gan ddefnyddio gweithrediadau atomig mewn strwythurau data sy'n gyfeillgar i GPU.
Ffynhonnell: hab.com