Posui eam Github .
Simplex est mensa GPU detrahenda apta ad expediendas centena milia decies centena millia insertorum secundo. In meo NVIDIA GTX 1060 laptop, codice interserit 64 decies centena millia passibus generata key-valorem paria circiter 210 ms, et aufert 32 decies centena fere 64 ms.
Hoc est, celeritas in laptop est circiter 300 decies centena milia interiicit/sec et 500 miliones delet/sec.
Tabula in CUDA scripta est, quamvis ars eadem applicari possit cum HLSL vel GLSL. Exsecutio plures limitationes habet ad invigilandum, magni effectus in cinematographico:
- Tantum 32-bit claves easdemque valores discursum sunt.
- Nullam mensam certam magnitudinem habet.
- Et haec magnitudo debet esse potentiae aequalis duobus.
Pro clavibus et valoribus, titulum simplicem delimitatorem (in superiore codice hoc 0xffffffff esse oportet).
Nullam mensam sine seras
Nullam in mensa utitur oratio aperta , id est, simpliciter ordinata praecipuorum valorum paria quae in memoria reposita sunt et praestantiores enucleatus effectus habet. Idem dici non potest de connexione, quae in investigatione monstratorem in nexu cum indice implicat. Nullam mensam simplex ordinata repono elementis KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Magnitudo mensae vis duorum est, non numerus primus, quia una instructio celeris sufficit ad larvam pow2/AND applicandam, sed modus operantis multo tardius est. Hoc interest in perscrutandis linearibus, cum in tabula lineari speculi socors indice singulis socors involvatur. Et per consequens modulo in unaquaque socors sumptum operationis additur.
Mensa solum clavem et valorem pro unoquoque elemento reponit, non Nullam clavis. Cum tabula tantum thesauros 32-bit claves, Nullam celerrime computatur. Codex supra Murmur3 Nullam utitur, qui paucas vices, XORs et multiplicationes facit.
In mensa Nullam densis artificiis tutelae utitur quae ordinis memoriae independentes sunt. Etiamsi aliqui scribunt operationes ordinem aliarum talium operationum perturbare, adhuc tabula Nullam conservabit statum rectam. De hoc infra loquemur. Ars magna cum cinematographicis video quae milia staminum currunt simul.
Claves et valores in mensa fasciculata initiales sunt vacuae.
Codex mutari potest ad claves ac valores etiam 64-frenas tractandas. Claves nuclei legunt, scribentes, comparant et operationes permuto postulant. Et bona atomica requirunt operationes legere et scribere. Fortunate, in CUDA, operationes scribes lege 32- et 64-bit valores atomici dum naturaliter alignantur (vide infra). ) et moderni cinematographici cinematographici 64-bit nuclei comparandi et operationes commutationis adiuvant. Scilicet, cum ad 64 frena moveatur, effectus leviter decrescet.
Nullam mensam civitatis
Quisque clavem pretii par in Nullam mensam habere potest unum ex quattuor civitatibus:
- Clavis et pretii vana sunt. In hoc statu, mensa Nullam initialized est.
- Clavis scripta est, sed valor nondum scriptus est. Si aliud linum nunc data est legens, tunc vacuus redit. Hoc normale, idem futurum esset si aliud stamen executionis paulo ante laborasset, et loquimur de instrumento concurrente.
- Tam clavis quam valor referuntur.
- Valor aliis stamina exsecutionis praesto est, sed clavis nondum est. Hoc fieri potest, quia exemplar programmandi CUDA exemplum memoriae laxe ordinatum habet. Hoc normale est, quoquo modo, clavis vacua adhuc est, etsi valor amplius non est.
Momenti momenti est quod semel clavis ad socors scripta est, eam non iam movet - etsi clavis deleta est, de hoc infra loquemur.
Nullam tabulae codicem etiam opera memoriae laxius ordinatae sunt in quibus ordo in quo memoria legatur et scriptus sit ignoratur. Sicut inspicimus insertionem, aspectum et deletionem in tabula Nullam, memini singulas par valoris in uno ex quattuor civitatibus supra descriptis.
Inserentes Nullam mensam
Munus CUDA quae paria clavem pretii inmittit in tabulam Nullam similis est:
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);
}
}
Ad inserere clavem, signum per tabulam detrahere ordinatam iterat initium cum clavis insertae Nullam. Quaelibet socors in ordine facit atomicam comparationem et operationem permuto quae clavem in socors vacuam comparat. Si mismatch detegitur, clavis in socors cum clavis insertis renovatur et tunc clavis rimae originalis redditur. Si haec clavis originalis vacua erat vel clavis inserta aequavit, tunc signum aptam socors inserendi invenit et valorem in rimam inseruit.
Si in nucleo vocatus gpu_hashtable_insert() cum eadem clavis elementa multa sint, tum quaelibet earum bona ad clavem foramen scribi possunt. Hoc normale habetur: unum praecipuum valorem scribens in vocationi succedet, sed cum haec omnia in pluribus sequelis exsecutionis parallela fiant, divinare non possumus quae memoria scribenda erit novissima.
Nullam mensam lookup
Codices ad claves quaerendas:
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);
}
}
Ad valorem clavis in mensa repositum invenire, iteramus per ordinatam initium a Nullam clavem quaerimus. In unaquaque socora, si clavem illam quam quaerimus est reprimimus, et si ita est, eius valorem reddimus. Etiam reprehendo si clavis vacua est et si ita investigationem abortum facimus.
Si clavem invenire non possumus, signum inane valorem reddit.
Omnes hae operationes inquisitionis simul per insertiones et deletiones perfici possunt. Quodlibet par in mensa habe- bat unum ex quatuor civitatibus suprascriptis pro fluxu.
Deletis in Nullam mensam
Code ad delendo claves:
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);
}
}
Clavis deleto modo insolito fit: clavem in mensa relinquimus et eius valorem notamus (non ipsam clavem) inanem. Hoc signum est simillimum lookup()nisi quod, cum compositus in clavibus reperitur, inanem facit valorem suum.
Ut supra, cum clavis ad foramen scriptum est, iam non movetur. Etiam cum elementum e mensa deletum est, clavis in loco remanet, eius valor simpliciter vacuus fit. Hoc significat quod non opus est utendo scribendo operationem atomicam propter valorem socors, quia nihil refert utrum valor praesens inanis sit vel non - adhuc inanis fiet.
Resizing a Nullam mensam
Magnitudinem mensae detrahentis mutare potes maiorem mensam creando et elementa non vacua ex veteri tabula in illam inserendo. Hanc functionem non feci, quia codicem sample simplex servare volui. Praeterea in programs CUDA, destinatio memoriae saepe fit in codice castrorum potius quam in nucleo CUDA.
articulus describitur quomodo modi seram munita notitia structura.
Competitiveness
Munus supra in codice excerpta gpu_hashtable_insert(), _lookup() ΠΈ _delete() aliquid unum clavem-valorem par ad tempus. Et inferiora gpu_hashtable_insert(), _lookup() ΠΈ _delete() ordinate emittunt paria in parallela, singula bina in separato GPU executioni filo;
// 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);
}
}
Sera resistens Nullam mensam subsidia concurrentia interserit, vultus, deletiones. Quia paria valoris semper in una quattuor civitatibus et clavibus non moventur, tabula rectitudinem praestat etiam cum diversae operationum rationes simul utuntur.
Sed si massam insertionum et deletionum in parallelis procedamus, et si input duplicata clavium series contineat, tunc praedicere non poterimus quae paria "vincent" - ad Nullam ultimam tabulam scribentur. Dicamus vocavimus codicem insertionem cum initus ordinata paria A/0 B/1 A/2 C/3 A/4. Cum codice complet, paria B/1 ΠΈ C/3 spondentur in mensa adesse, sed simul quaelibet paria in ea apparebunt A/0, A/2 aut A/4. Hoc forte vel non potest esse problema - totum pendet in applicatione. Scias ante non esse claves duplicatas in inordinatione ordinatarum, vel non curare quod valor proxime scriptus est.
Si hoc problema pro te est, necesse est paria duplicata in diversa ratio vocationis CUDA dividere. In CUDA, quaelibet operatio quae nucleum vocat, semper perficit ante nucleum proximum (saltem intra unum fila. In filis diversis, nuclei paralleli efficiuntur). In exemplo supra, si vocas unum nucleum cum A/0 B/1 A/2 C/3et alter cum A/4Ergo clavem A erit in valorem 4.
Nunc fama num munera lookup() ΠΈ delete() uti monstratorem campum vel volatile Nullam aciem in mensa paria. Asserit:
Compilator eligere potest ad optimize legentem et scribentem ad memoriam globalem vel communem ... Hae optimizationes uti keyword erre possunt.
volatile: ... quaelibet relatio huius variabilis compilata in veram memoriam legendi vel scribendi disciplinam.
Recta considerations applicationem non eget volatile. Si filum exsecutionis valorem conditivo ab operatione antecedente legeris utitur, tunc notitia aliquantulum outdlatata erit. Sed adhuc, hoc indicium est ex recto statu tabulae detrahendae certo momento nuclei vocati. Si notitia novissima uti debes, indice uti potes volatilesed tunc effectus leviter decrescet: iuxta meae probationes, cum deletis 32 miliones elementorum, celeritas minuitur ex 500 miliones deletionum/sec ad 450 miliones deletionum/sec.
productivity
In probatione ad inserendas elementa 64 decies centena et deleta 32 decies centena, inter certamen std::unordered_map et nulla fere mensa est GPU:
std::unordered_map MS 70 insertis et exemptis elementis et deinde eas liberans unordered_map (Expellere decies elementis multum temporis, quod intus unordered_map multae memoriae prouinciis fiunt). Honeste loqui, std:unordered_map omnino diversis angustiis. Una CPU est stamina executionis, cuiuslibet magnitudinis valores clavos sustinet, bene in rates utendo altas facit et post multas deletiones stabiles effectus ostendit.
Duratio tabulae detrahendae pro GPU et inter-programma communicationis 984 ms erat. Hoc includit tempus, ponendo mensam in memoriam et eam delendo (I GB memoriae collocans unum tempus, quod in CUDA aliquo tempore occupat), elementa inserendo et delendo, et in ea iterando. Omnes codices ad et ex memoria card video habiti sunt etiam in ratione.
Nullam mensa ipsa 271 ms complere cepit. Hoc includit tempus a card electronicis inserendis et delendis elementis confectum, nec rationem temporis in memoria describendi et iterandi super tabulam consequentem. Si tabula GPU diu vivit, aut si mensa detrahens in memoria card curriculi contineatur (exempli gratia, tabulam Nullam creare quae ab aliis GPU codice adhibebitur et non processus centralis), tunc test eventus pertinet.
Nullam mensam pro card in video ostendit altam observantiam propter altitudinem throughput et activam parallelizationem.
defectuum oeconomicorum
Architectura tabulae Nullam paucas quaestiones ut conscias habeat:
- Perscrutatio linearis ligaturae impedita est, quae facit claves in mensa minus quam perfecte ponendas.
- Claves non remota sunt per munus
deleteac tempus cluunt mensam.
Quam ob rem, functio mensae detrahendae gradatim degradare potest, praesertim si diu existit ac numerosos inserit ac delevit. Uno modo ad haec incommoda mitiganda est novas tabulas resilire cum rate satis depresso utendo et claves remotas in retrahendo eliquare.
Ad quaestiones descriptos illustrandas, codice superiore utar ut mensam cum 128 decies elementis conficiam et per 4 decies elementis ansam, donec 124 decies centena foramina implevi (utendo rate circiter 0,96). Hic est eventus mensae, cuiuslibet ordinis est CUDA nucleus vocatus ut 4 decies centena nova elementa in unam tabulam detraheret;
Usus rate
Insertio durationis 4 elementorum
0,00
11,608448 ms (361,314798 decies centena millia claves/sec.)
0,03
11,751424 ms (356,918799 decies centena millia claves/sec.)
0,06
11,942592 ms (351,205515 decies centena millia claves/sec.)
0,09
12,081120 ms (347,178429 decies centena millia claves/sec.)
0,12
12,242560 ms (342,600233 decies centena millia claves/sec.)
0,16
12,396448 ms (338,347235 decies centena millia claves/sec.)
0,19
12,533024 ms (334,660176 decies centena millia claves/sec.)
0,22
12,703328 ms (330,173626 decies centena millia claves/sec.)
0,25
12,884512 ms (325,530693 decies centena millia claves/sec.)
0,28
13,033472 ms (321,810182 decies centena millia claves/sec.)
0,31
13,239296 ms (316,807174 decies centena millia claves/sec.)
0,34
13,392448 ms (313,184256 decies centena millia claves/sec.)
0,37
13,624000 ms (307,861434 decies centena millia claves/sec.)
0,41
13,875520 ms (302,280855 decies centena millia claves/sec.)
0,44
14,126528 ms (296,909756 decies centena millia claves/sec.)
0,47
14,399328 ms (291,284699 decies centena millia claves/sec.)
0,50
14,690304 ms (285,515123 decies centena millia claves/sec.)
0,53
15,039136 ms (278,892623 decies centena millia claves/sec.)
0,56
15,478656 ms (270,973402 decies centena millia claves/sec.)
0,59
15,985664 ms (262,379092 decies centena millia claves/sec.)
0,62
16,668673 ms (251,627968 decies centena millia claves/sec.)
0,66
17,587200 ms (238,486174 decies centena millia claves/sec.)
0,69
18,690048 ms (224,413765 decies centena millia claves/sec.)
0,72
20,278816 ms (206,831789 decies centena millia claves/sec.)
0,75
22,545408 ms (186,038058 decies centena millia claves/sec.)
0,78
26,053312 ms (160,989275 decies centena millia claves/sec.)
0,81
31,895008 ms (131,503463 decies centena millia claves/sec.)
0,84
42,103294 ms (99,619378 decies centena millia claves/sec.)
0,87
61,849056 ms (67,815164 decies centena millia claves/sec.)
0,90
105,695999 ms (39,682713 decies centena millia claves/sec.)
0,94
240,204636 ms (17,461378 decies centena millia claves/sec.)
Utendo crescit, effectus decrescit. Hoc in pluribus non est optabile. Si applicatio elementa in tabulam immittit et ea reicit (exempli gratia cum verba numerandi in libro), haec quaestio non est. Sed si applicatio tabulam Nullam longaevam adhibet (exempli gratia, in graphice editore ut congregem partes imaginum non vacuas, ubi user notitias frequenter inserit et delet), tunc hic mores inconveniens esse potest.
Mensam Nullam et mensus est profunditatem perscrutandi post 64 decies centena millia inserta (utendo factor 0,5). Mediocris profunditas erat 0,4774, ergo claves maximae erant vel optime socors vel socors ab optimo loco. Maxima altitudo sonans LX erat.
Tunc metiebam profunditatem perscrutandi in mensa cum 124 decies insertis (utendo factor 0,97). Mediocris profunditas iam 10,1757 erat, et maxima - 6474 (!!). Lineae sentiendi effectus guttae significanter in altum utendo rates.
Optimum est ut hanc tabulam Nullam utendo rate humilis. Sed tum sumptu memoriae consummatio perficiendi augemus. Fortunate, in casu 32-clavium et bonorum, hoc iustificari potest. Si in exemplo superiore in tabula cum 128 miliones elementorum utendo factorem 0,25 tenemus, tunc non plus quam 32 decies centena milia in ea ponere possumus, reliquae 96 decies centena foramina amittentur - 8 bytes pro singulis binis. , 768 MB amissae memoriae.
Quaeso note nos loquimur de amissione schedularum memoria, quae pretiosior est resource quam ratio memoriae. Etsi recentissimi cinematographici graphicae cinematographicae maxime recentes quae CUDA sustentant memoriae saltem 4 GB habent (tempore scriptionis NVIDIA 2080 Ti 11 GB habet), tamen sapientissimum consilium non erit tantas pondera perdere.
Postea plura scribemus de tabulis celorum creandis pro schedulis video quae quaestiones profunditatem non habent, ac vias ad reuse foramina deleta.
Sonus altitudinem mensurae
Ut perscrutandi altitudinem clavis definias, possumus clavem detrahere (specialem tabulam indicem) ex ipsa tabula eius indicem extrahere;
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Propter magicam duorum numerorum binariorum complementum et quod capacitas mensae Nullam duo potestati duorum, aditus etiam operabitur cum index clavis ad principium tabulae movetur. Clavem sumamus ad hashed 1 , sed in socors inseritur 3. Deinde ad mensam cum capacitate 4 accipimus (3 β 1) & 3quod idem valet ad II.
conclusio,
Si quaestiones vel commentationes habes, electronicas ad me aut novum topic in aperire.
Codex hic ex praestantissimis articulis sub inspiratione conscriptus est;
In futurum, scribere pergam de instrumentis cinematographicis pro instrumentis cinematographicis ac faciendis analysi. Consilia mea includunt catenam, Robin Hood hashing, et cuculus hasting utens operationes atomicas in notitia structurae quae sunt GPU amicae.
Source: www.habr.com