Би үүнийг Github дээр нийтэлсэн .
Энэ нь секундэд хэдэн зуун сая оруулгыг боловсруулах чадвартай энгийн GPU хэш хүснэгт юм. Миний NVIDIA GTX 1060 зөөврийн компьютер дээр код нь санамсаргүй байдлаар үүсгэгдсэн 64 сая түлхүүр утгыг 210 мс-д оруулж, 32 сая хосыг 64 мс-д устгадаг.
Өөрөөр хэлбэл, зөөврийн компьютер дээрх хурд нь ойролцоогоор 300 сая оруулга / сек, 500 сая устгах / сек байна.
Хүснэгтийг CUDA хэл дээр бичсэн боловч HLSL эсвэл GLSL-д ижил аргыг хэрэглэж болно. Энэхүү хэрэгжилт нь видео картын өндөр гүйцэтгэлийг хангах хэд хэдэн хязгаарлалттай байдаг:
- Зөвхөн 32 битийн түлхүүрүүд болон ижил утгуудыг боловсруулдаг.
- Хэш хүснэгт нь тогтмол хэмжээтэй байна.
- Мөн энэ хэмжээ нь чадалтай хоёр тэнцүү байх ёстой.
Түлхүүр болон утгуудын хувьд та энгийн хязгаарлагч тэмдэглэгээг нөөцлөх хэрэгтэй (дээрх кодонд энэ нь 0xffffffff байна).
Түгжээгүй хэш хүснэгт
Хэш хүснэгт нь нээлттэй хаягийг ашиглана , өөрөөр хэлбэл, энэ нь санах ойд хадгалагдсан, дээд зэргийн кэшийн гүйцэтгэлтэй түлхүүр-утга хосын массив юм. Холбогдсон жагсаалтаас заагч хайхтай холбоотой гинжний хувьд ижил зүйлийг хэлж болохгүй. Хэш хүснэгт нь элементүүдийг хадгалах энгийн массив юм KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Хүснэгтийн хэмжээ нь анхны тоо биш хоёрын зэрэгтэй, учир нь нэг хурдан заавар нь pow2/AND маск хэрэглэхэд хангалттай боловч модулийн оператор нь хамаагүй удаан байдаг. Шугаман хүснэгт хайхад үүрний индексийг үүр бүрт ороосон байх ёстой тул энэ нь шугаман шалгалтын хувьд чухал юм. Үүний үр дүнд үйл ажиллагааны өртөгийг үүр бүрт модуль нэмдэг.
Хүснэгт нь зөвхөн элемент бүрийн түлхүүр ба утгыг хадгалдаг бөгөөд түлхүүрийн хэш биш. Хүснэгтэнд зөвхөн 32 битийн түлхүүр хадгалагддаг тул хэшийг маш хурдан тооцдог. Дээрх код нь Murmur3 хэшийг ашигладаг бөгөөд энэ нь хэдхэн ээлж, XOR болон үржүүлэлтийг гүйцэтгэдэг.
Хэш хүснэгт нь санах ойн дарааллаас хамааралгүй түгжих хамгаалалтын аргуудыг ашигладаг. Зарим бичих үйлдлүүд нь бусад үйлдлийн дарааллыг тасалдуулж байсан ч хэш хүснэгт нь зөв төлөвийг хадгалсаар байх болно. Энэ талаар бид доор ярих болно. Энэ техник нь олон мянган хэлхээг нэгэн зэрэг ажиллуулдаг видео картуудтай маш сайн ажилладаг.
Хэш хүснэгт дэх түлхүүрүүд болон утгуудыг хоосон болгохын тулд эхлүүлсэн.
Кодыг 64 битийн түлхүүр, утгыг зохицуулахын тулд өөрчилж болно. Түлхүүрүүд нь атомын унших, бичих, харьцуулах, солих үйлдлүүдийг шаарддаг. Мөн утгууд нь атомын унших, бичих үйлдлийг шаарддаг. Аз болоход, CUDA-д 32 ба 64 битийн утгуудын унших, бичих үйлдлүүд нь байгалийн жамаар тохируулагдсан тохиолдолд атомын шинж чанартай байдаг (доороос харна уу). ), орчин үеийн видео картууд нь 64 битийн атомын харьцуулах, солилцох үйлдлийг дэмждэг. Мэдээжийн хэрэг, 64 бит рүү шилжих үед гүйцэтгэл бага зэрэг буурах болно.
Хэш хүснэгтийн төлөв
Хэш хүснэгт дэх түлхүүр-утга хос бүр дөрвөн төлөвийн аль нэгтэй байж болно:
- Түлхүүр болон утга хоосон байна. Энэ төлөвт хэш хүснэгтийг эхлүүлнэ.
- Түлхүүр нь бичигдсэн боловч үнэ цэнэ нь хараахан бичигдээгүй байна. Хэрэв өөр хэлхээ өгөгдөл уншиж байгаа бол энэ нь хоосон буцаж ирнэ. Энэ бол хэвийн зүйл, хэрэв өөр нэг гүйцэтгэлийн хэлхээ арай эрт ажиллаж байсан бол ижил зүйл тохиолдох байсан бөгөөд бид зэрэгцээ өгөгдлийн бүтцийн тухай ярьж байна.
- Түлхүүр болон утгыг хоёуланг нь бүртгэдэг.
- Энэ утгыг гүйцэтгэх бусад хэлхээнд ашиглах боломжтой боловч түлхүүр хараахан болоогүй байна. CUDA програмчлалын загвар нь санах ойн загвар нь сул захиалгатай байдаг тул ийм зүйл тохиолдож болно. Энэ бол хэвийн зүйл; ямар ч тохиолдолд утга нь байхгүй болсон ч гэсэн түлхүүр хоосон хэвээр байна.
Нэг чухал нюанс бол түлхүүрийг үүрэнд бичсэний дараа тэр нь хөдлөхөө больсон - түлхүүр устгагдсан байсан ч бид энэ тухай доор ярих болно.
Хэш хүснэгтийн код нь санах ой унших, бичих дараалал нь тодорхойгүй сул эрэмбэлэгдсэн санах ойн загваруудтай ч ажилладаг. Хэш хүснэгтэд оруулах, хайх, устгах үйлдлийг харахад түлхүүр-утга хос бүр дээр дурдсан дөрвөн төлөвийн аль нэгэнд байгааг санаарай.
Хэш хүснэгтэд оруулах
Түлхүүр-утга хосыг хэш хүснэгтэд оруулах CUDA функц дараах байдалтай байна.
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);
}
}
Түлхүүр оруулахын тулд код нь оруулсан түлхүүрийн хэшээс эхлэн хэш хүснэгтийн массиваар давтагдана. Массив дахь оролт бүр нь атомын харьцуулах, солих үйлдлийг гүйцэтгэдэг бөгөөд энэ үүрэнд байгаа түлхүүрийг хоосон болгон харьцуулдаг. Хэрэв тохиромжгүй байдал илэрсэн бол үүрэнд байгаа түлхүүрийг оруулсан түлхүүрээр шинэчилж, дараа нь анхны үүрний түлхүүрийг буцаана. Хэрэв энэ анхны түлхүүр хоосон байсан эсвэл оруулсан түлхүүртэй таарч байвал код нь оруулах тохиромжтой үүрийг олж, оруулсан утгыг үүрэнд оруулсан болно.
Хэрэв нэг цөмийн дуудлагад байвал gpu_hashtable_insert() ижил түлхүүртэй олон элемент байгаа бол тэдгээрийн аль нэг утгыг түлхүүрийн үүрэнд бичиж болно. Энэ нь хэвийн гэж тооцогддог: дуудлагын явцад нэг түлхүүр-утга бичих нь амжилттай болох боловч энэ бүхэн гүйцэтгэлийн хэд хэдэн хэлхээнд зэрэгцэн явагддаг тул санах ойн аль бичвэр сүүлчийнх болохыг бид таамаглаж чадахгүй.
Хэш хүснэгт хайх
Түлхүүр хайх код:
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);
}
}
Хүснэгтэд хадгалагдсан түлхүүрийн утгыг олохын тулд хайж буй түлхүүрийнхээ хэшээс эхлэн массивыг давтдаг. Слот бүрт бид түлхүүр нь бидний хайж буй түлхүүр мөн эсэхийг шалгаж, хэрэв тийм бол бид түүний утгыг буцаана. Мөн бид түлхүүр хоосон эсэхийг шалгаж, хэрэв байгаа бол хайлтыг зогсооно.
Хэрэв бид түлхүүрийг олж чадахгүй бол код нь хоосон утгыг буцаана.
Эдгээр бүх хайлтын үйлдлүүдийг оруулах, устгах замаар нэгэн зэрэг гүйцэтгэж болно. Хүснэгтийн хос бүр дээр урсгалын хувьд дээр дурдсан дөрвөн төлөвийн аль нэгтэй байх болно.
Хэш хүснэгтээс устгаж байна
Түлхүүрүүдийг устгах код:
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);
}
}
Түлхүүрийг устгах нь ер бусын аргаар хийгддэг: бид түлхүүрийг хүснэгтэд үлдээж, түүний утгыг (түлхүүр өөрөө биш) хоосон гэж тэмдэглэнэ. Энэ код нь маш төстэй юм lookup(), түлхүүр дээр тохирох зүйл олдвол түүний утгыг хоосон болгодог.
Дээр дурьдсанчлан, түлхүүрийг үүрэнд бичсэний дараа түүнийг шилжүүлэхээ болино. Хүснэгтээс элементийг устгасан ч гэсэн түлхүүр нь байрандаа үлдэж, утга нь хоосон болно. Энэ нь бид слот утгын хувьд атом бичих үйлдлийг ашиглах шаардлагагүй гэсэн үг, учир нь одоогийн утга хоосон эсэх нь хамаагүй - энэ нь хоосон хэвээр байх болно.
Хэш хүснэгтийн хэмжээг өөрчлөх
Та илүү том хүснэгт үүсгэж, хуучин хүснэгтээс хоосон бус элементүүдийг оруулах замаар хэш хүснэгтийн хэмжээг өөрчилж болно. Би жишээ кодыг энгийн байлгахыг хүссэн учраас энэ функцийг хэрэгжүүлээгүй. Түүнчлэн CUDA программуудад санах ойн хуваарилалтыг ихэвчлэн CUDA цөмд бус хост кодоор хийдэг.
Нийтлэл Ийм түгжээгээр хамгаалагдсан өгөгдлийн бүтцийг хэрхэн өөрчлөх талаар тайлбарласан.
Өрсөлдөх чадвар
Дээрх функцийн кодын хэсгүүдэд gpu_hashtable_insert(), _lookup() и _delete() нэг удаад нэг түлхүүр-утга хос боловсруулах. Мөн доогуур gpu_hashtable_insert(), _lookup() и _delete() хос бүрийг тусдаа 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);
}
}
Түгжихэд тэсвэртэй хэш хүснэгт нь нэгэн зэрэг оруулах, хайх, устгах үйлдлийг дэмждэг. Түлхүүр-утгийн хосууд үргэлж дөрвөн төлөвийн аль нэгэнд байдаг бөгөөд товчлуурууд нь хөдөлдөггүй тул хүснэгт нь янз бүрийн төрлийн үйлдлүүдийг нэгэн зэрэг ашигласан ч зөв байдлыг баталгаажуулдаг.
Гэсэн хэдий ч, хэрэв бид оруулах, устгах багцыг зэрэгцүүлэн боловсруулж, хосуудын оролтын массив нь давхардсан түлхүүрүүдийг агуулж байвал аль хосууд "ялах" талаар таамаглах боломжгүй болно - хамгийн сүүлд хэш хүснэгтэд бичигдэх болно. Бид оруулах кодыг хос оролтын массиваар дуудсан гэж бодъё A/0 B/1 A/2 C/3 A/4. Код дууссаны дараа хослоно B/1 и C/3 Хүснэгтэнд байх нь баталгаатай боловч үүнтэй зэрэгцэн аль ч хос нь гарч ирнэ A/0, A/2 буюу A/4. Энэ нь асуудал байж болно, үгүй ч байж магадгүй - энэ бүхэн програмаас хамаарна. Та оролтын массив дээр давхардсан түлхүүр байхгүй гэдгийг урьдчилан мэдэж болно, эсвэл аль утгыг хамгийн сүүлд бичсэн нь танд хамаагүй байж болно.
Хэрэв энэ нь танд асуудал байгаа бол та давхардсан хосуудыг өөр өөр CUDA системийн дуудлага болгон салгах хэрэгтэй. CUDA-д цөмийг дууддаг аливаа үйлдэл нь дараагийн цөмийн дуудлагаас өмнө үргэлж дуусдаг (ядаж нэг thread дотор. Өөр өөр урсгалуудад цөмүүд зэрэгцээ гүйцэтгэгддэг). Дээрх жишээн дээр хэрэв та нэг цөмийг дуудвал A/0 B/1 A/2 C/3, нөгөө нь хамт A/4, дараа нь түлхүүр A үнэ цэнийг авах болно 4.
Одоо функцууд шаардлагатай эсэх талаар ярилцъя lookup() и delete() хэш хүснэгт дэх хосуудын массив руу энгийн эсвэл тогтворгүй заагч ашиглана уу. Үүнд:
Хөрвүүлэгч нь глобал эсвэл хуваалцсан санах ойд унших, бичих үйлдлийг оновчтой болгохоор сонгож болно... Эдгээр оновчлолыг түлхүүр үг ашиглан идэвхгүй болгож болно.
volatile: ... энэ хувьсагчийн дурын лавлагааг бодит санах ойн унших, бичих заавар болгон хөрвүүлдэг.
Зөв байдлыг харгалзан үзэх нь өргөдөл гаргах шаардлагагүй volatile. Хэрэв гүйцэтгэлийн хэлхээ нь өмнөх уншсан үйлдлээс хадгалагдсан утгыг ашигладаг бол энэ нь бага зэрэг хуучирсан мэдээллийг ашиглах болно. Гэсэн хэдий ч энэ нь цөмийн дуудлагын тодорхой агшинд хэш хүснэгтийн зөв төлөвийн мэдээлэл юм. Хэрэв та хамгийн сүүлийн үеийн мэдээллийг ашиглах шаардлагатай бол индексийг ашиглаж болно volatile, гэхдээ дараа нь гүйцэтгэл бага зэрэг буурах болно: миний туршилтын дагуу 32 сая элементийг устгах үед хурд нь 500 сая устгал / сек-ээс 450 сая устгалт / сек болж буурсан.
Бүтээмж
Туршилтанд 64 сая элемент оруулж, 32 саяыг нь устгаж, хоорондоо өрсөлдөж байна std::unordered_map GPU-д зориулсан хэш хүснэгт бараг байхгүй:
std::unordered_map Элементүүдийг оруулах, арилгахад 70 мс зарцуулсан unordered_map (Сая сая элементүүдээс салах нь маш их цаг хугацаа шаарддаг, учир нь дотор нь unordered_map олон санах ойн хуваарилалт хийгдсэн). Үнэнийг хэлэхэд, std:unordered_map огт өөр хязгаарлалт. Энэ нь гүйцэтгэлийн нэг CPU урсгал бөгөөд ямар ч хэмжээтэй түлхүүр утгыг дэмждэг, ашиглалтын өндөр түвшинд сайн ажилладаг, олон устгасны дараа тогтвортой гүйцэтгэлийг харуулдаг.
GPU болон програм хоорондын харилцааны хэш хүснэгтийн үргэлжлэх хугацаа 984 мс байв. Үүнд хүснэгтийг санах ойд байршуулах, устгах (1 ГБ санах ойг нэг удаа хуваарилах, энэ нь CUDA-д бага зэрэг хугацаа шаардагдах), элементүүдийг оруулах, устгах, тэдгээрийг давталт хийхэд зарцуулсан хугацаа орно. Видео картын санах ой руу болон түүнээс гарах бүх хуулбарыг мөн харгалзан үзнэ.
Хэш хүснэгтийг дуусгахад 271 мс зарцуулсан. Үүнд видео картын элементүүдийг оруулах, устгахад зарцуулсан цагийг багтаасан бөгөөд санах ойд хуулж, үр дүнгийн хүснэгтийг давтахад зарцуулсан хугацааг тооцохгүй. Хэрэв GPU хүснэгт удаан хугацаагаар амьдардаг эсвэл хэш хүснэгт нь видео картын санах ойд бүрэн агуулагдсан бол (жишээлбэл, төв процессор биш харин өөр GPU код ашиглах хэш хүснэгт үүсгэх) туршилтын үр дүн хамааралтай.
Видео картын хэш хүснэгт нь өндөр дамжуулах чадвар, идэвхтэй параллелжуулалтаас шалтгаалан өндөр гүйцэтгэлийг харуулж байна.
дутагдлууд
Хэш хүснэгтийн архитектур нь анхаарах ёстой хэд хэдэн асуудалтай:
- Шугаман шалгалтыг кластер хийх замаар саатуулдаг бөгөөд энэ нь хүснэгтэд байгаа түлхүүрүүдийг төгс бус байрлуулахад хүргэдэг.
- Функцийг ашиглан товчлууруудыг хасдаггүй
deleteмөн цаг хугацаа өнгөрөхөд тэд ширээгээ эмх замбараагүй болгодог.
Үүний үр дүнд хэш хүснэгтийн гүйцэтгэл аажмаар доройтож болно, ялангуяа энэ нь удаан хугацаанд оршин тогтнож, олон тооны оруулга, устгалттай бол. Эдгээр сул талуудыг багасгах нэг арга бол ашиглалтын хурд багатай шинэ хүснэгтийг дахин тохируулах, дахин тохируулах явцад устгасан түлхүүрүүдийг шүүж авах явдал юм.
Тайлбарласан асуудлуудыг тайлбарлахын тулд би дээрх кодыг ашиглан 128 сая элемент бүхий хүснэгт үүсгэж, 4 сая слотыг дүүргэх хүртэл 124 сая элементээр давталт хийнэ (ашиглалтын хувь ойролцоогоор 0,96). Энд үр дүнгийн хүснэгт байна, мөр бүр нь нэг хэш хүснэгтэд 4 сая шинэ элемент оруулах CUDA цөмийн дуудлага юм.
Ашиглалтын түвшин
Оруулах хугацаа 4 элемент
0,00
11,608448 мс (361,314798 сая түлхүүр/сек)
0,03
11,751424 мс (356,918799 сая түлхүүр/сек)
0,06
11,942592 мс (351,205515 сая түлхүүр/сек)
0,09
12,081120 мс (347,178429 сая түлхүүр/сек)
0,12
12,242560 мс (342,600233 сая түлхүүр/сек)
0,16
12,396448 мс (338,347235 сая түлхүүр/сек)
0,19
12,533024 мс (334,660176 сая түлхүүр/сек)
0,22
12,703328 мс (330,173626 сая түлхүүр/сек)
0,25
12,884512 мс (325,530693 сая түлхүүр/сек)
0,28
13,033472 мс (321,810182 сая түлхүүр/сек)
0,31
13,239296 мс (316,807174 сая түлхүүр/сек)
0,34
13,392448 мс (313,184256 сая түлхүүр/сек)
0,37
13,624000 мс (307,861434 сая түлхүүр/сек)
0,41
13,875520 мс (302,280855 сая түлхүүр/сек)
0,44
14,126528 мс (296,909756 сая түлхүүр/сек)
0,47
14,399328 мс (291,284699 сая түлхүүр/сек)
0,50
14,690304 мс (285,515123 сая түлхүүр/сек)
0,53
15,039136 мс (278,892623 сая түлхүүр/сек)
0,56
15,478656 мс (270,973402 сая түлхүүр/сек)
0,59
15,985664 мс (262,379092 сая түлхүүр/сек)
0,62
16,668673 мс (251,627968 сая түлхүүр/сек)
0,66
17,587200 мс (238,486174 сая түлхүүр/сек)
0,69
18,690048 мс (224,413765 сая түлхүүр/сек)
0,72
20,278816 мс (206,831789 сая түлхүүр/сек)
0,75
22,545408 мс (186,038058 сая түлхүүр/сек)
0,78
26,053312 мс (160,989275 сая түлхүүр/сек)
0,81
31,895008 мс (131,503463 сая түлхүүр/сек)
0,84
42,103294 мс (99,619378 сая түлхүүр/сек)
0,87
61,849056 мс (67,815164 сая түлхүүр/сек)
0,90
105,695999 мс (39,682713 сая түлхүүр/сек)
0,94
240,204636 мс (17,461378 сая түлхүүр/сек)
Ашиглалт нэмэгдэхийн хэрээр гүйцэтгэл буурдаг. Энэ нь ихэнх тохиолдолд хүсээгүй юм. Хэрэв програм нь хүснэгтэд элементүүдийг оруулаад дараа нь тэдгээрийг хаявал (жишээлбэл, номонд үг тоолох үед) энэ нь асуудал биш юм. Гэхдээ хэрэв програм нь урт хугацааны хэш хүснэгтийг ашигладаг бол (жишээлбэл, хэрэглэгч байнга мэдээлэл оруулж, устгадаг зургийн хоосон бус хэсгийг хадгалахын тулд график засварлагч) ашигладаг бол энэ үйлдэл нь асуудалтай байж болно.
Мөн 64 сая оруулга хийсний дараа хэш хүснэгтийг шалгах гүнийг хэмжсэн (ашиглалтын коэффициент 0,5). Дундаж гүн нь 0,4774 байсан тул ихэнх түлхүүрүүд аль болох хамгийн сайн үүрэнд эсвэл хамгийн сайн байрлалаас нэг үүрний зайд байсан. Дууны хамгийн их гүн нь 60 байв.
Дараа нь би 124 сая оруулга бүхий ширээн дээр шалгах гүнийг хэмжсэн (ашиглалтын коэффициент 0,97). Дундаж гүн нь аль хэдийн 10,1757, хамгийн их нь - 6474 (!!). Шугаман мэдрэгчийн гүйцэтгэл өндөр ашиглалтын үед мэдэгдэхүйц буурдаг.
Энэ хэш хүснэгтийн ашиглалтын түвшинг бага байлгах нь хамгийн сайн арга юм. Гэхдээ дараа нь бид санах ойн хэрэглээний зардлаар гүйцэтгэлийг нэмэгдүүлдэг. Аз болоход, 32 битийн түлхүүрүүд болон утгуудын хувьд үүнийг зөвтгөж болно. Дээрх жишээн дээр 128 сая элемент бүхий хүснэгтэд ашиглалтын коэффициентийг 0,25-д үлдээвэл бид 32 саяас илүүгүй элементийг байрлуулж болох бөгөөд үлдсэн 96 сая слот алдагдах болно - хос бүрт 8 байт. , 768 MB санах ой алдагдсан.
Бид системийн санах ойноос илүү үнэ цэнэтэй нөөц болох видео картын санах ойг алдах тухай ярьж байгааг анхаарна уу. Хэдийгээр CUDA-г дэмждэг орчин үеийн ихэнх ширээний график картууд дор хаяж 4 ГБ санах ойтой (бичиж байх үед NVIDIA 2080 Ti нь 11 ГБ санах ойтой байсан) ийм хэмжээгээр алдах нь хамгийн ухаалаг шийдвэр биш хэвээр байна.
Хожим нь би шалгах гүнтэй холбоотой асуудалгүй видео картуудад зориулсан хэш хүснэгтүүд, мөн устгасан слотуудыг дахин ашиглах аргуудын талаар дэлгэрэнгүй бичих болно.
Дууны гүнийг хэмжих
Түлхүүрийг шалгах гүнийг тодорхойлохын тулд бид түлхүүрийн хэшийг (түүний хамгийн тохиромжтой хүснэгтийн индекс) бодит хүснэгтийн индексээс гаргаж авч болно:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Хоёрын хоёрын нөхөх хоёртын тоонуудын ид шидтэй, хэш хүснэгтийн багтаамж хоёр хоёрын зэрэгтэй байдаг тул энэ арга нь түлхүүр индексийг хүснэгтийн эхэнд шилжүүлсэн ч ажиллах болно. 1-д хэшлэгдсэн боловч 3-р үүрэнд оруулсан түлхүүрийг авцгаая. Дараа нь 4-ийн багтаамжтай хүснэгтийн хувьд бид авна. (3 — 1) & 3, энэ нь 2-той тэнцэнэ.
дүгнэлт
Хэрэв танд асуулт эсвэл сэтгэгдэл байвал над руу имэйлээр илгээнэ үү эсвэл шинэ сэдэв нээх.
Энэхүү кодыг маш сайн нийтлэлүүдээс сэдэвлэн бичсэн болно:
Ирээдүйд би видео картуудад зориулсан хэш хүснэгтийн хэрэгжилтийн талаар үргэлжлүүлэн бичиж, тэдгээрийн гүйцэтгэлд дүн шинжилгээ хийх болно. Миний төлөвлөгөөнд GPU-д ээлтэй өгөгдлийн бүтцэд атомын үйлдлийг ашиглан гинжлэх, Робин Гуд хэшлэх, хөхөө хэшлэх зэрэг багтана.
Эх сурвалж: www.habr.com