Chuir mé ar Github é .
Is tábla simplí GPU hash é atá in ann na céadta milliún ionsá in aghaidh an tsoicind a phróiseáil. Ar mo ríomhaire glúine NVIDIA GTX 1060, cuireann an cód isteach 64 milliún péirí eochair-luacha a ghintear go randamach i thart ar 210 ms agus baintear 32 milliún péirí i thart ar 64 ms.
Is é sin, is é an luas ar ríomhaire glúine ná thart ar 300 milliún ionsáite/soicind agus 500 milliún scriosta/soicind.
Tá an tábla scríofa i CUDA, cé gur féidir an teicníocht chéanna a chur i bhfeidhm ar HLSL nó GLSL. Tá roinnt teorainneacha ag baint leis an gcur i bhfeidhm chun ardfheidhmíocht ar chárta físeáin a chinntiú:
- Ní dhéantar ach eochracha 32-giotán agus na luachanna céanna a phróiseáil.
- Tá méid seasta ag an tábla hash.
- Agus ní mór an méid seo a bheith comhionann le dhá chuig an chumhacht.
Le haghaidh eochracha agus luachanna, ní mór duit marcóir teorannóir simplí a chur in áirithe (is é seo an cód thuas 0xffffffff).
Tábla hash gan glas
Úsáideann an tábla hash seoltaí oscailte le , is é sin, níl ann ach sraith de phéirí eochairluacha a stóráiltear sa chuimhne agus a bhfuil feidhmíocht taisce níos fearr acu. Ní féidir an rud céanna a rá maidir le slabhraiú, rud a bhaineann le pointeoir a chuardach i liosta nasctha. Is éard atá i tábla hash ná eagar simplí chun eilimintí a stóráil KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Cumhacht dhá cheann, ní príomhuimhir, é méid an tábla, toisc gur leor teagasc tapa amháin chun an masc pow2/AND a chur i bhfeidhm, ach tá an t-oibreoir modulus i bhfad níos moille. Tá sé seo tábhachtach i gcás iniúchadh líneach, toisc go gcaithfidh an t-innéacs sliotán a bheith fillte i ngach sliotán i gclár líneach cuardaigh. Agus mar thoradh air sin, cuirtear modulo le costas na hoibríochta i ngach sliotán.
Ní stórálann an tábla ach eochair agus luach gach eiliminte, ní hais den eochair. Ós rud é nach stórálann an tábla ach eochracha 32-giotán, ríomhtar an hash go han-tapa. Úsáideann an cód thuas hash Murmur3, a dhéanann ach cúpla shifts, XORs agus iolraithe.
Úsáideann an tábla hash teicnící cosanta glasála atá neamhspleách ar ord cuimhne. Fiú má chuireann roinnt oibríochtaí scríbhneoireachta isteach ar ord oibríochtaí eile dá leithéid, coimeádfaidh an tábla hash an staid cheart fós. Déanfaimid labhairt faoi seo thíos. Oibríonn an teicníc go hiontach le cártaí físeáin a ritheann na mílte snáithe i gcomhthráth.
Tá na heochracha agus na luachanna sa tábla hash tosaithe a fholmhú.
Is féidir an cód a mhodhnú chun eochracha agus luachanna 64-giotán a láimhseáil freisin. Teastaíonn oibríochtaí adamhach léite, scríobh, agus comparáide agus babhtála le haghaidh eochracha. Agus éilíonn luachanna oibríochtaí adamhach léamh agus scríobh. Go fortunately, i CUDA, tá oibríochtaí léamh-scríofa le haghaidh luachanna 32- agus 64-giotán adamhach chomh fada agus a bhíonn siad ailínithe go nádúrtha (féach thíos). ), agus tacaíonn cártaí físe nua-aimseartha le hoibríochtaí comparáide agus malairte adamhach 64-giotán. Ar ndóigh, nuair a bhogann tú go 64 giotán, laghdóidh an fheidhmíocht beagán.
Staid tábla hash
Féadfaidh ceann amháin de cheithre stát a bheith ag gach péire eochairluacha i dtábla hash:
- Tá eochair agus luach folamh. Sa stát seo, cuirtear tús leis an tábla hash.
- Tá an eochair scríofa síos, ach níl an luach scríofa fós. Má tá snáithe eile ag léamh sonraí faoi láthair, filleann sé folamh ansin. Is gnách é seo, bheadh an rud céanna tarlaithe dá n-oibreodh snáithe forghníomhaithe eile beagán níos luaithe, agus táimid ag caint faoi struchtúr sonraí comhthráthach.
- Taifeadtar an eochair agus an luach araon.
- Tá an luach ar fáil do snáitheanna forghníomhaithe eile, ach níl an eochair fós. Féadfaidh sé seo tarlú toisc go bhfuil múnla cuimhne atá ordaithe go scaoilte ag samhail ríomhchláraithe CUDA. Is gnách é seo; ar aon nós, tá an eochair fós folamh, fiú mura bhfuil an luach amhlaidh a thuilleadh.
Is nuance tábhachtach é, nuair a bheidh an eochair scríofa chuig an sliotán, ní ghluaiseann sé a thuilleadh - fiú má scriostar an eochair, labhairfimid faoi seo thíos.
Oibríonn an cód tábla hash fiú le samhlacha cuimhne scaoilte ordaithe ina bhfuil an t-ord ina léitear agus ina scríobhtar an chuimhne anaithnid. Agus muid ag féachaint ar ionsá, cuardach, agus scriosadh i dtábla hash, cuimhnigh go bhfuil gach péire eochairluacha i gceann de na ceithre stát a gcuirtear síos orthu thuas.
Á chur isteach i tábla hash
Breathnaíonn an fheidhm CUDA a chuireann péirí eochairluacha isteach i dtábla hash mar seo:
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);
}
}
Chun eochair a chur isteach, atriálann an cód tríd an eagar tábla hash ag tosú le hash na heochrach a cuireadh isteach. Déanann gach sliotán san eagar oibríocht adamhach comparáide agus babhtála a dhéanann comparáid idir an eochair sa sliotán sin agus folmhú. Má aimsítear neamhréir, déantar an eochair sa sliotán a nuashonrú leis an eochair ionsáite, agus ansin cuirtear an eochair sliotán bunaidh ar ais. Má bhí an eochair bhunaidh seo folamh nó má mheaitseáil leis an eochair a cuireadh isteach, ansin fuair an cód sliotán oiriúnach le cur isteach agus chuir sé an luach ionsáite isteach sa sliotán.
Más rud é i nglao eithne amháin gpu_hashtable_insert() tá gnéithe iolracha leis an eochair chéanna, ansin is féidir aon cheann dá luachanna a scríobh chuig an sliotán eochair. Meastar gur gnáthrud é seo: éireoidh le ceann de na heochairluachanna a scríobhtar le linn an ghlao, ach toisc go dtarlaíonn sé seo go léir go comhthreomhar laistigh de roinnt snáitheanna forghníomhaithe, ní féidir linn a thuar cén scríobh cuimhne a bheidh ar an gceann deireanach.
Cuardach tábla hash
Cód le haghaidh eochracha cuardaigh:
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);
}
}
Chun luach eochair atá stóráilte i dtábla a fháil, déanaimid athrá tríd an eagar ag tosú le hash na heochrach atá á lorg againn. I ngach sliotán, déanaimid seiceáil an bhfuil an eochair an ceann atá á lorg againn, agus má tá, cuirimid a luach ar ais. Déanaimid seiceáil freisin an bhfuil an eochair folamh, agus má tá, déanaimid deireadh leis an gcuardach.
Mura féidir linn an eochair a aimsiú, filleann an cód luach folamh.
Is féidir na hoibríochtaí cuardaigh seo go léir a dhéanamh i gcomhthráth trí ionsá agus scriosadh. Beidh ceann amháin de na ceithre staid a thuairiscítear thuas don sreabhadh ag gach péire sa tábla.
Scriosadh i tábla hash
Cód chun eochracha a scriosadh:
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);
}
}
Déantar eochair a scriosadh ar bhealach neamhghnách: fágaimid an eochair sa tábla agus marcáilimid a luach (ní an eochair féin) mar folamh. Tá an cód seo an-chosúil le lookup(), ach amháin nuair a aimsítear meaits ar eochair, déanann sé a luach folamh.
Mar a luadh thuas, a luaithe a scríobhtar eochair chuig sliotán, ní bhogtar a thuilleadh í. Fiú nuair a scriostar eilimint ón tábla, fanann an eochair i bhfeidhm, ní bhíonn ach a luach folamh. Ciallaíonn sé seo nach gá dúinn oibríocht scríofa adamhach a úsáid don luach sliotán, mar is cuma an bhfuil an luach reatha folamh nó nach bhfuil - beidh sé fós folamh.
Tábla hash a athrú méid
Is féidir leat méid tábla hash a athrú trí tábla níos mó a chruthú agus eilimintí neamhfholamh ón sean tábla a chur isteach ann. Níor chuir mé an fheidhm seo i bhfeidhm mar bhí mé ag iarraidh an cód samplach a choinneáil simplí. Thairis sin, i gcláir CUDA, is minic a dhéantar leithdháileadh cuimhne sa chód óstach seachas san eithne CUDA.
An t-alt cur síos ar conas struchtúr sonraí atá faoi chosaint ghlas a mhodhnú.
Iomaíochas
Sna míreanna cód feidhme thuas gpu_hashtable_insert(), _lookup() и _delete() péire eochairluacha amháin a phróiseáil ag an am. Agus níos ísle gpu_hashtable_insert(), _lookup() и _delete() sraith péirí a phróiseáil go comhthreomhar, gach péire i snáithe forghníomhaithe GPU ar leith:
// 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);
}
}
Tacaíonn an tábla hash-resistant glas le ionsáigh chomhthráthacha, cuardaigh agus scriosta. Toisc go bhfuil péirí eochairluacha i gceann de cheithre stát i gcónaí agus nach bogann na heochracha, ráthaíonn an tábla cruinneas fiú nuair a úsáidtear cineálacha éagsúla oibríochtaí ag an am céanna.
Mar sin féin, má dhéanaimid próiseáil ar bhaisc ionsáite agus scriosta ag an am céanna, agus má tá eochracha dúblacha san raon ionchuir de phéirí, ní bheidh muid in ann a thuar cé na péirí a “bhuann”—a scríobhfar chuig an tábla hash deireanach. Ligean le rá thugamar an cód isteach le sraith ionchuir de péirí A/0 B/1 A/2 C/3 A/4. Nuair a chríochnaíonn an cód, péirí B/1 и C/3 ráthaithe a bheith i láthair sa tábla, ach ag an am céanna beidh aon cheann de na péirí le feiceáil ann A/0, A/2 nó A/4. D'fhéadfadh sé seo nó nach féidir a bheith ina fhadhb - braitheann sé ar fad ar an iarratas. B'fhéidir go mbeadh a fhios agat roimh ré nach bhfuil aon eochracha dúblacha san eagar ionchuir, nó b'fhéidir nach bhfuil suim agat cén luach a scríobhadh an uair dheireanach.
Más fadhb é seo duit, ní mór duit na péirí dúblacha a scaradh i nglaonna córais CUDA éagsúla. I CUDA, críochnaítear aon oibríocht a ghlaonn an eithne i gcónaí roimh an gcéad ghlao eithne eile (ar a laghad laistigh de snáithe amháin. I snáitheanna éagsúla, déantar eithne go comhthreomhar). Sa sampla thuas, má ghlaonn tú eithne amháin le A/0 B/1 A/2 C/3, agus an ceann eile le A/4, ansin an eochair A gheobhaidh an luach 4.
Anois, déanaimis labhairt faoi cé acu ba chóir feidhmeanna lookup() и delete() úsáid a bhaint as pointeoir simplí nó so-ghalaithe le sraith péirí sa tábla hais. Deirtear go:
B'fhéidir go roghnódh an tiomsaitheoir léann agus scríobhann a bharrfheabhsú chuig cuimhne dhomhanda nó go cuimhne chomhroinnte... Is féidir na huasmhéaduithe seo a dhíchumasú ag baint úsáide as an eochairfhocal
volatile: ... déantar aon tagairt don athróg seo a thiomsú i bhfíor-theagasc chun cuimhne a léamh nó a scríobh.
Ní gá iarratas a dhéanamh ar bhreithnithe ceartais volatile. Má úsáideann an snáithe forghníomhaithe luach i dtaisce ó oibríocht a léadh níos luaithe, beidh sé ag baint úsáide as faisnéis atá beagán as dáta. Ach fós, is faisnéis é seo ó staid cheart an tábla hash ag tráth áirithe den ghlao eithne. Más gá duit an fhaisnéis is déanaí a úsáid, is féidir leat an t-innéacs a úsáid volatile, ach ansin laghdóidh an fheidhmíocht beagán: de réir mo thástálacha, nuair a bhí 32 milliún eilimint á scriosadh, laghdaigh an luas ó 500 milliún scriosadh / soic go 450 milliún scriosadh / soic.
Táirgiúlacht
Sa tástáil chun 64 milliún eilimint a chur isteach agus 32 milliún díobh a scriosadh, tá iomaíocht idir std::unordered_map agus níl mórán tábla hash ann don GPU:
std::unordered_map chaith sé 70 ms ag cur isteach agus ag baint eilimintí agus ansin iad a shaoradh unordered_map (Bíonn go leor ama fáil réidh leis na milliúin eilimintí, mar gheall ar an taobh istigh unordered_map déantar leithdháiltí cuimhne iolracha). Ag labhairt go hionraic, std:unordered_map srianta go hiomlán difriúil. Is snáithe forghníomhaithe amháin LAP é, tacaíonn sé le heochairluachanna ar aon mhéid, feidhmíonn sé go maith ag rátaí arda úsáide, agus taispeánann sé feidhmíocht chobhsaí tar éis scriosadh iolrach.
Ba é ré an tábla hash don GPU agus don chumarsáid idirchláir ná 984 ms. Áirítear leis seo an t-am a chaitear ag cur an tábla i gcuimhne agus é a scriosadh (1 GB de chuimhne a leithdháileadh uair amháin, a thógann roinnt ama i CUDA), eilimintí a chur isteach agus a scriosadh, agus atriall os a gcionn. Cuirtear gach cóip go dtí agus ó chuimhne an chárta físeáin san áireamh freisin.
Thóg an tábla hash féin 271 ms le críochnú. Áiríonn sé seo an t-am a chaitheann an cárta físeáin ag cur isteach agus ag scriosadh eilimintí, agus ní chuireann sé san áireamh an t-am a chaitear ag cóipeáil sa chuimhne agus ag atriall thar an tábla mar thoradh air. Má tá an tábla GPU ina chónaí ar feadh i bhfad, nó má tá an tábla hash go hiomlán i gcuimhne an chárta físeán (mar shampla, a chruthú tábla hash a bheidh in úsáid ag cód GPU eile agus nach bhfuil an próiseálaí lárnach), ansin tá toradh na tástála ábhartha.
Léiríonn an tábla hash le haghaidh cárta físeán ardfheidhmíocht mar gheall ar tréchur ard agus comhthreomharú gníomhach.
Teorainneacha
Tá roinnt saincheisteanna le cur ar an eolas faoi ailtireacht tábla hash:
- Cuirtear bac ar thaiscéalaíocht líneach trí bhraisliú, rud a fhágann nach gcuirtear na heochracha sa tábla níos lú ná go foirfe.
- Ní bhaintear eochracha leis an bhfeidhm
deleteagus le himeacht ama déanann siad bearradh ar an mbord.
Mar thoradh air sin, is féidir le feidhmíocht tábla hash díghrádú de réir a chéile, go háirithe má tá sé ann ar feadh i bhfad agus go bhfuil go leor cuir isteach agus scriosadh ann. Bealach amháin leis na míbhuntáistí seo a mhaolú ná athshnaidhmeadh isteach i dtábla nua le ráta úsáide measartha íseal agus na heochracha a baineadh amach a scagadh le linn an athsheasaithe.
Chun na saincheisteanna a gcuirtear síos orthu a léiriú, bainfidh mé úsáid as an gcód thuas chun tábla a chruthú le 128 milliún eilimint agus lúb trí 4 mhilliún eilimint go dtí go mbeidh 124 milliún sliotán líonta agam (ráta úsáide de thart ar 0,96). Seo é an tábla torthaí, is glao eithne CUDA é gach sraith chun 4 mhilliún eilimint nua a chur isteach i dtábla hash amháin:
Ráta úsáide
Fad an ionchuir 4 eilimint
0,00
11,608448 ms (361,314798 milliún eochair / soicind)
0,03
11,751424 ms (356,918799 milliún eochair / soicind)
0,06
11,942592 ms (351,205515 milliún eochair / soicind)
0,09
12,081120 ms (347,178429 milliún eochair / soicind)
0,12
12,242560 ms (342,600233 milliún eochair / soicind)
0,16
12,396448 ms (338,347235 milliún eochair / soicind)
0,19
12,533024 ms (334,660176 milliún eochair / soicind)
0,22
12,703328 ms (330,173626 milliún eochair / soicind)
0,25
12,884512 ms (325,530693 milliún eochair / soicind)
0,28
13,033472 ms (321,810182 milliún eochair / soicind)
0,31
13,239296 ms (316,807174 milliún eochair / soicind)
0,34
13,392448 ms (313,184256 milliún eochair / soicind)
0,37
13,624000 ms (307,861434 milliún eochair / soicind)
0,41
13,875520 ms (302,280855 milliún eochair / soicind)
0,44
14,126528 ms (296,909756 milliún eochair / soicind)
0,47
14,399328 ms (291,284699 milliún eochair / soicind)
0,50
14,690304 ms (285,515123 milliún eochair / soicind)
0,53
15,039136 ms (278,892623 milliún eochair / soicind)
0,56
15,478656 ms (270,973402 milliún eochair / soicind)
0,59
15,985664 ms (262,379092 milliún eochair / soicind)
0,62
16,668673 ms (251,627968 milliún eochair / soicind)
0,66
17,587200 ms (238,486174 milliún eochair / soicind)
0,69
18,690048 ms (224,413765 milliún eochair / soicind)
0,72
20,278816 ms (206,831789 milliún eochair / soicind)
0,75
22,545408 ms (186,038058 milliún eochair / soicind)
0,78
26,053312 ms (160,989275 milliún eochair / soicind)
0,81
31,895008 ms (131,503463 milliún eochair / soicind)
0,84
42,103294 ms (99,619378 milliún eochair / soicind)
0,87
61,849056 ms (67,815164 milliún eochair / soicind)
0,90
105,695999 ms (39,682713 milliún eochair / soicind)
0,94
240,204636 ms (17,461378 milliún eochair / soicind)
De réir mar a mhéadaíonn úsáid, laghdaítear feidhmíocht. Níl sé seo inmhianaithe i bhformhór na gcásanna. Má chuireann feidhmchlár gnéithe isteach i dtábla agus go gcaitheann sé iad (mar shampla, nuair a bhíonn focail á gcomhaireamh i leabhar), ní fadhb í sin. Ach má úsáideann an feidhmchlár tábla hash fadsaoil (mar shampla, in eagarthóir grafaice chun codanna neamhfholamh d'íomhánna a stóráil ina gcuireann an t-úsáideoir isteach agus go scriosann sé faisnéis go minic), is féidir leis an iompar seo a bheith ina fhadhb.
Agus thomhas an tábla hash doimhneacht probing tar éis 64 milliún ionsá (fachtóir úsáide 0,5). Ba é 0,4774 an meándoimhneacht, agus mar sin bhí an chuid is mó de na heochracha sa sliotán is fearr is féidir nó i sliotán amháin ar shiúl ón suíomh is fearr. Ba é 60 an doimhneacht fuaime uasta.
Thomhais mé ansin an doimhneacht taiscéalaíochta ar bhord le 124 milliún ionsá (fachtóir úsáide 0,97). Ba é 10,1757 an meándoimhneacht cheana féin, agus an t-uasmhéid - 6474 (!!). Titeann feidhmíocht braite líneach go suntasach ag rátaí arda úsáide.
Is fearr ráta úsáide an tábla hash seo a choinneáil íseal. Ach ansin déanaimid méadú ar fheidhmíocht ar chostas tomhaltas cuimhne. Go fortunately, i gcás eochracha agus luachanna 32-giotán, is féidir é seo a chosaint. Más rud é sa sampla thuas, i dtábla le 128 milliún eilimint, coinnímid an fachtóir úsáide de 0,25, ansin ní féidir linn níos mó ná 32 milliún eilimint a chur ann, agus caillfear na 96 milliún sliotán atá fágtha - 8 bytes do gach péire. , 768 MB de chuimhne caillte.
Tabhair faoi deara le do thoil go bhfuil muid ag caint faoi chailliúint cuimhne cárta físeán, atá ina acmhainn níos luachmhaire ná cuimhne córais. Cé go bhfuil ar a laghad 4 GB de chuimhne ag an chuid is mó de na cártaí grafaicí deisce nua-aimseartha a thacaíonn le CUDA (ag an am seo scríofa, tá 2080 GB ag an NVIDIA 11 Ti), ní bheadh sé fós ar an gcinneadh is ciallmhar méideanna den sórt sin a chailleadh.
Níos déanaí beidh mé ag scríobh níos mó faoi chruthú táblaí hash le haghaidh cártaí físeán nach bhfuil fadhbanna le doimhneacht probing, chomh maith le bealaí chun athúsáid sliotán scriosta.
Tomhas doimhneacht fuaime
Chun doimhneacht iniúchta eochair a chinneadh, is féidir linn hash na heochrach (an t-innéacs tábla idéalach) a bhaint as a innéacs tábla iarbhír:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Mar gheall ar an draíocht a bhaineann le dhá uimhir dhénártha chomhlánú dhá agus ar an bhfíric go bhfuil cumas an tábla hash dhá chumhacht, oibreoidh an cur chuige seo fiú nuair a bhogtar an t-innéacs eochair go dtí tús an tábla. Glacaimid eochair a hashed go 1, ach a chuirtear isteach i sliotán 3. Ansin le haghaidh tábla le toilleadh 4 a fháil againn (3 — 1) & 3, atá comhionann le 2 .
Conclúid
Má tá ceisteanna nó tuairimí agat, seol ríomhphost chugam ag nó oscail topaic nua i.
Scríobhadh an cód seo faoi inspioráid ó earraí den scoth:
Sa todhchaí, leanfaidh mé ag scríobh faoi fheidhmiú tábla hash le haghaidh cártaí físeán agus anailís a dhéanamh ar a bhfeidhmíocht. I measc na bpleananna atá agam tá slabhraiú, hashing Robin Hood, agus hashing cuach ag baint úsáide as oibríochtaí adamhach i struchtúir sonraí atá cairdiúil don GPU.
Foinse: will.com