ខ្ញុំបានបង្ហោះវានៅលើ Github គម្រោងថ្មី A Simple GPU Hash Table.

វាជាតារាង GPU hash ដ៏សាមញ្ញដែលមានសមត្ថភាពដំណើរការរាប់រយលានបញ្ចូលក្នុងមួយវិនាទី។ នៅលើកុំព្យូទ័រយួរដៃ NVIDIA GTX 1060 របស់ខ្ញុំ លេខកូដបញ្ចូល 64 លានគូដែលបង្កើតដោយចៃដន្យក្នុងរយៈពេលប្រហែល 210 ms និងដកចេញ 32 លានគូក្នុងរយៈពេលប្រហែល 64 ms ។

នោះគឺល្បឿននៅលើកុំព្យូទ័រយួរដៃគឺប្រហែល 300 លានបញ្ចូល / វិនាទី និង 500 លានលុប / វិនាទី។

តារាងត្រូវបានសរសេរក្នុង CUDA ទោះបីជាបច្ចេកទេសដូចគ្នាអាចត្រូវបានអនុវត្តចំពោះ HLSL ឬ GLSL ក៏ដោយ។ ការអនុវត្តមានដែនកំណត់ជាច្រើនដើម្បីធានាបាននូវដំណើរការខ្ពស់នៅលើកាតវីដេអូ៖

• មានតែគ្រាប់ចុច 32 ប៊ីត និងតម្លៃដូចគ្នាត្រូវបានដំណើរការ។
• តារាង hash មានទំហំថេរ។
• ហើយទំហំនេះត្រូវតែស្មើនឹងពីរទៅថាមពល។

សម្រាប់កូនសោ និងតម្លៃ អ្នកត្រូវកក់សញ្ញាកំណត់ព្រំដែនសាមញ្ញ (ក្នុងកូដខាងលើនេះគឺ 0xffffffff)។

តារាង Hash ដោយគ្មានសោ

តារាង hash ប្រើអាសយដ្ឋានបើកចំហជាមួយ ការស៊ើបអង្កេតលីនេអ៊ែរនោះ​គឺ​វា​គ្រាន់​តែ​ជា​អារេ​នៃ​គូ​តម្លៃ​គន្លឹះ​ដែល​ត្រូវ​បាន​រក្សា​ទុក​ក្នុង​អង្គ​ចងចាំ​ និង​មាន​ដំណើរការ​ឃ្លាំង​សម្ងាត់​ល្អ​ជាង។ មិន​អាច​និយាយ​ដូច​គ្នា​នេះ​បាន​សម្រាប់​ការ​ដាក់​ច្រវាក់ ដែល​ពាក់ព័ន្ធ​នឹង​ការ​ស្វែងរក​ទ្រនិច​ក្នុង​បញ្ជី​ដែល​បាន​ភ្ជាប់។ តារាង hash គឺជាអារេសាមញ្ញសម្រាប់ផ្ទុកធាតុ KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

ទំហំនៃតារាងគឺជាថាមពលនៃពីរ មិនមែនជាលេខបឋមទេ ពីព្រោះការណែនាំលឿនមួយគឺគ្រប់គ្រាន់ដើម្បីអនុវត្តរបាំង pow2/AND ប៉ុន្តែប្រតិបត្តិករម៉ូឌុលគឺយឺតជាង។ នេះមានសារៈសំខាន់នៅក្នុងករណីនៃការស៊ើបអង្កេតលីនេអ៊ែរ ចាប់តាំងពីក្នុងការរកមើលតារាងលីនេអ៊ែរ សន្ទស្សន៍រន្ធដោតត្រូវតែរុំនៅក្នុងរន្ធនីមួយៗ។ ហើយជាលទ្ធផលការចំណាយនៃប្រតិបត្តិការត្រូវបានបន្ថែមម៉ូឌុលនៅក្នុងរន្ធនីមួយៗ។

តារាងរក្សាទុកតែកូនសោ និងតម្លៃសម្រាប់ធាតុនីមួយៗ មិនមែនជាសញ្ញានៃសោនោះទេ។ ដោយសារតារាងរក្សាទុកតែគ្រាប់ចុច 32 ប៊ីត លេខហាសត្រូវបានគណនាយ៉ាងរហ័ស។ កូដខាងលើប្រើ Murmur3 hash ដែលដំណើរការតែការផ្លាស់ប្តូរមួយចំនួន XOR និងគុណ។

តារាង hash ប្រើបច្ចេកទេសការពារចាក់សោដែលឯករាជ្យពីលំដាប់អង្គចងចាំ។ ទោះបីជាប្រតិបត្តិការសរសេរខ្លះរំខានដល់លំដាប់នៃប្រតិបត្តិការបែបនេះក៏ដោយ តារាង hash នឹងនៅតែរក្សាស្ថានភាពត្រឹមត្រូវ។ យើងនឹងនិយាយអំពីរឿងនេះខាងក្រោម។ បច្ចេកទេសនេះដំណើរការល្អជាមួយកាតវីដេអូដែលដំណើរការរាប់ពាន់ខ្សែក្នុងពេលដំណាលគ្នា។

គ្រាប់ចុច និងតម្លៃនៅក្នុងតារាង hash ត្រូវបានចាប់ផ្តើមទៅជាទទេ។

កូដអាចត្រូវបានកែប្រែដើម្បីគ្រប់គ្រងគ្រាប់ចុច 64 ប៊ីត និងតម្លៃផងដែរ។ គ្រាប់ចុចត្រូវការអាតូមអាន សរសេរ និងប្រតិបត្តិការប្រៀបធៀប និងប្តូរ។ ហើយតម្លៃទាមទារនូវប្រតិបត្តិការអាន និងសរសេរអាតូម។ ជាសំណាងល្អនៅក្នុង CUDA ប្រតិបត្តិការអាន-សរសេរសម្រាប់តម្លៃ 32- និង 64-bit គឺជាអាតូម ដរាបណាពួកវាត្រូវបានតម្រឹមតាមធម្មជាតិ (សូមមើលខាងក្រោម)។ នៅទីនេះ) និងកាតវីដេអូទំនើបគាំទ្រប្រតិបត្តិការប្រៀបធៀប និងផ្លាស់ប្តូរអាតូមិច 64 ប៊ីត។ ជាការពិតណាស់នៅពេលផ្លាស់ទីទៅ 64 ប៊ីត ដំណើរការនឹងថយចុះបន្តិច។

ស្ថានភាពតារាង Hash

គូតម្លៃគន្លឹះនីមួយៗនៅក្នុងតារាង hash អាចមានរដ្ឋមួយក្នុងចំណោមរដ្ឋចំនួនបួន៖

• កូនសោ និងតម្លៃគឺទទេ។ នៅក្នុងស្ថានភាពនេះ តារាង hash ត្រូវបានចាប់ផ្តើម។
• កូនសោ​ត្រូវ​បាន​សរសេរ​ចុះ ប៉ុន្តែ​តម្លៃ​មិន​ទាន់​ត្រូវ​បាន​សរសេរ​ចុះ​នៅ​ឡើយ​ទេ។ ប្រសិនបើខ្សែបណ្ដាញផ្សេងទៀតកំពុងអានទិន្នន័យ នោះវានឹងត្រឡប់ទទេ។ នេះគឺជារឿងធម្មតា រឿងដដែលនឹងកើតឡើងប្រសិនបើដំណើរការប្រតិបត្តិមួយផ្សេងទៀតបានដំណើរការមុននេះបន្តិច ហើយយើងកំពុងនិយាយអំពីរចនាសម្ព័ន្ធទិន្នន័យស្របគ្នា។
• ទាំងសោ និងតម្លៃត្រូវបានកត់ត្រា។
• តម្លៃ​គឺ​អាច​ប្រើ​បាន​សម្រាប់​ការ​ប្រតិបត្តិ​ខ្សែ​ស្រឡាយ​ផ្សេង​ទៀត ប៉ុន្តែ​គន្លឹះ​មិន​ទាន់​មាន​នៅ​ឡើយ​ទេ។ នេះអាចកើតឡើងដោយសារតែគំរូសរសេរកម្មវិធី CUDA មានគំរូអង្គចងចាំដែលបានបញ្ជារលុង នេះជារឿងធម្មតាទេ នៅក្នុងព្រឹត្តិការណ៍ណាមួយ សោនៅតែទទេ ទោះបីជាតម្លៃលែងដូច្នេះក៏ដោយ។

ចំនុចសំខាន់មួយគឺថានៅពេលដែលសោត្រូវបានសរសេរទៅរន្ធដោត វាលែងផ្លាស់ទីទៀតហើយ - ទោះបីជាសោត្រូវបានលុបក៏ដោយ យើងនឹងនិយាយអំពីរឿងនេះខាងក្រោម។

កូដតារាង hash ថែមទាំងដំណើរការជាមួយគំរូអង្គចងចាំរលុង ដែលលំដាប់ដែលអង្គចងចាំត្រូវបានអាន និងសរសេរមិនស្គាល់។ នៅពេលយើងពិនិត្យមើលការបញ្ចូល ការរកមើល និងការលុបនៅក្នុងតារាងសញ្ញា សូមចាំថាគូតម្លៃសោនីមួយៗស្ថិតនៅក្នុងរដ្ឋមួយក្នុងចំណោមរដ្ឋទាំងបួនដែលបានពិពណ៌នាខាងលើ។

ការបញ្ចូលទៅក្នុងតារាងសញ្ញា

អនុគមន៍ 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);
    }
}

ដើម្បីបញ្ចូលកូនសោ លេខកូដធ្វើម្តងទៀតតាមរយៈអារេតារាង hash ដោយចាប់ផ្តើមដោយសញ្ញានៃគ្រាប់ចុចដែលបានបញ្ចូល។ រន្ធដោតនីមួយៗនៅក្នុងអារេធ្វើប្រតិបត្តិការប្រៀបធៀប និងស្វបអាតូម ដែលប្រៀបធៀបសោនៅក្នុងរន្ធនោះទៅទទេ។ ប្រសិនបើរកឃើញភាពមិនស៊ីសង្វាក់គ្នា សោនៅក្នុងរន្ធដោតត្រូវបានធ្វើបច្ចុប្បន្នភាពជាមួយនឹងសោដែលបានបញ្ចូល ហើយបន្ទាប់មកគ្រាប់ចុចរន្ធដោតដើមត្រូវបានត្រឡប់មកវិញ។ ប្រសិនបើសោដើមនេះទទេ ឬត្រូវគ្នានឹងសោដែលបានបញ្ចូល នោះលេខកូដបានរកឃើញរន្ធដោតសមរម្យសម្រាប់ការបញ្ចូល និងបញ្ចូលតម្លៃដែលបានបញ្ចូលទៅក្នុងរន្ធដោត។

ប្រសិនបើនៅក្នុងការហៅខឺណែលមួយ។ gpu_hashtable_insert() មានធាតុជាច្រើនដែលមានគ្រាប់ចុចដូចគ្នា បន្ទាប់មកតម្លៃណាមួយរបស់ពួកគេអាចត្រូវបានសរសេរទៅកាន់រន្ធដោតគ្រាប់ចុច។ នេះត្រូវបានចាត់ទុកថាជារឿងធម្មតា៖ មួយក្នុងចំនោម key-value សរសេរក្នុងអំឡុងពេលហៅទូរសព្ទនឹងជោគជ័យ ប៉ុន្តែដោយសារអ្វីៗទាំងអស់នេះកើតឡើងស្របគ្នាក្នុងដំណើរការប្រតិបត្តិ យើងមិនអាចទាយទុកជាមុនបានថាតើការសរសេរ memory មួយណានឹងក្លាយជាការសរសេរចុងក្រោយ។

រកមើលតារាង Hash

កូដសម្រាប់សោស្វែងរក៖

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()លើកលែងតែនៅពេលដែលការផ្គូផ្គងត្រូវបានរកឃើញនៅលើសោ វាធ្វើឱ្យតម្លៃរបស់វាទទេ។

ដូចដែលបានរៀបរាប់ខាងលើ នៅពេលដែលសោមួយត្រូវបានសរសេរទៅកាន់រន្ធដោត នោះវានឹងលែងផ្លាស់ទីទៀតហើយ។ សូម្បីតែនៅពេលដែលធាតុមួយត្រូវបានលុបចេញពីតារាងក៏ដោយ សោនៅតែនៅដដែល តម្លៃរបស់វានឹងក្លាយជាទទេ។ នេះមានន័យថាយើងមិនចាំបាច់ប្រើប្រតិបត្តិការសរសេរអាតូមិកសម្រាប់តម្លៃរន្ធនោះទេព្រោះវាមិនមានបញ្ហាថាតើតម្លៃបច្ចុប្បន្នទទេឬអត់ - វានឹងនៅតែទទេ។

ផ្លាស់ប្តូរទំហំតារាង hash

អ្នកអាចផ្លាស់ប្តូរទំហំនៃតារាង hash ដោយបង្កើតតារាងធំមួយ ហើយបញ្ចូលធាតុមិនទទេពីតារាងចាស់ទៅក្នុងវា។ ខ្ញុំ​មិន​បាន​អនុវត្ត​មុខងារ​នេះ​ទេ ព្រោះ​ខ្ញុំ​ចង់​រក្សា​កូដ​គំរូ​ឲ្យ​សាមញ្ញ។ លើសពីនេះទៅទៀត នៅក្នុងកម្មវិធី CUDA ការបែងចែកអង្គចងចាំជាញឹកញាប់ត្រូវបានធ្វើនៅក្នុងកូដម៉ាស៊ីនជាជាងនៅក្នុងខឺណែល CUDA ។

នៅក្នុងអត្ថបទ តារាង Hash Free Wait-Free ចាក់សោ ពិពណ៌នាអំពីរបៀបកែប្រែរចនាសម្ព័ន្ធទិន្នន័យដែលការពារដោយចាក់សោបែបនេះ។

ភាពប្រកួតប្រជែង

នៅក្នុងអត្ថបទកូដមុខងារខាងលើ 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);
    }
}

តារាង hash ធន់នឹងការចាក់សោ គាំទ្រការបញ្ចូល ការរកមើល និងការលុបក្នុងពេលដំណាលគ្នា។ ដោយសារតែគូតម្លៃសោតែងតែស្ថិតនៅក្នុងរដ្ឋមួយក្នុងចំណោមរដ្ឋទាំងបួន ហើយសោមិនផ្លាស់ទី តារាងធានានូវភាពត្រឹមត្រូវ ទោះបីជាប្រតិបត្តិការប្រភេទផ្សេងគ្នាត្រូវបានប្រើក្នុងពេលដំណាលគ្នាក៏ដោយ។

ទោះជាយ៉ាងណាក៏ដោយ ប្រសិនបើយើងដំណើរការបណ្តុំនៃការបញ្ចូល និងការលុបស្របគ្នា ហើយប្រសិនបើអារេបញ្ចូលនៃគូមានសោស្ទួន នោះយើងនឹងមិនអាចទស្សន៍ទាយថាតើគូណានឹង "ឈ្នះ" ទេ នឹងត្រូវបានសរសេរទៅកាន់តារាងសញ្ញាចុងក្រោយ។ ឧបមាថា យើងបានហៅកូដបញ្ចូលជាមួយអារេបញ្ចូលនៃគូ A/0 B/1 A/2 C/3 A/4. នៅពេលដែលលេខកូដបញ្ចប់ គូ B/1 и C/3 ត្រូវ​បាន​ធានា​ថា​នឹង​មាន​វត្តមាន​ក្នុង​តារាង ប៉ុន្តែ​នៅ​ពេល​ជាមួយ​គ្នា​នោះ​គូ​ណា​មួយ​នឹង​បង្ហាញ​នៅ​ក្នុង​នោះ។ A/0, A/2 ឬ A/4. នេះអាចឬមិនមែនជាបញ្ហា - វាទាំងអស់គឺអាស្រ័យលើកម្មវិធី។ អ្នកអាចដឹងជាមុនថាមិនមានគ្រាប់ចុចស្ទួននៅក្នុងអារេបញ្ចូលទេ ឬអ្នកប្រហែលជាមិនខ្វល់ថាតម្លៃមួយណាត្រូវបានសរសេរចុងក្រោយ។

ប្រសិនបើនេះជាបញ្ហាសម្រាប់អ្នក នោះអ្នកត្រូវបំបែកគូស្ទួនទៅជាការហៅតាមប្រព័ន្ធ CUDA ផ្សេងៗ។ នៅ​ក្នុង CUDA ប្រតិបត្តិការ​ណា​មួយ​ដែល​ហៅ​ខឺណែល​តែង​តែ​បញ្ចប់​មុន​ការ​ហៅ​ខឺណែល​បន្ទាប់ (យ៉ាង​ហោច​ណាស់​ក្នុង​ខ្សែ​ស្រឡាយ​មួយ។ ក្នុង​ខ្សែ​ស្រឡាយ​ផ្សេង​គ្នា ខឺណែល​ត្រូវ​បាន​ប្រតិបត្តិ​ស្រប​គ្នា)។ នៅក្នុងឧទាហរណ៍ខាងលើប្រសិនបើអ្នកហៅខឺណែលមួយជាមួយ A/0 B/1 A/2 C/3, និងផ្សេងទៀតជាមួយ A/4បន្ទាប់មកគន្លឹះ A នឹងទទួលបានតម្លៃ 4.

ឥឡូវនេះសូមនិយាយអំពីថាតើមុខងារគួរ lookup() и delete() ប្រើទ្រនិចធម្មតា ឬងាយនឹងបង្កជាហេតុទៅអារេនៃគូក្នុងតារាងសញ្ញា។ ឯកសារ CUDA បញ្ជាក់ថា៖

កម្មវិធីចងក្រងអាចជ្រើសរើសដើម្បីបង្កើនប្រសិទ្ធភាពការអាន និងសរសេរទៅកាន់សតិសកល ឬចែករំលែក... ការបង្កើនប្រសិទ្ធភាពទាំងនេះអាចត្រូវបានបិទដោយប្រើពាក្យគន្លឹះ volatile: ... សេចក្តីយោងណាមួយចំពោះអថេរនេះត្រូវបានចងក្រងជាមេរៀនអាន ឬសរសេរសតិពិត។

ការពិចារណាលើភាពត្រឹមត្រូវមិនតម្រូវឱ្យមានការអនុវត្តទេ។ volatile. ប្រសិនបើខ្សែស្រឡាយប្រតិបត្តិប្រើតម្លៃឃ្លាំងសម្ងាត់ពីប្រតិបត្តិការដែលបានអានពីមុន នោះវានឹងប្រើព័ត៌មានហួសសម័យបន្តិច។ ប៉ុន្តែនៅតែ នេះជាព័ត៌មានពីស្ថានភាពត្រឹមត្រូវនៃតារាង hash នៅគ្រាជាក់លាក់នៃការហៅ kernel ។ ប្រសិនបើអ្នកត្រូវការប្រើព័ត៌មានចុងក្រោយ អ្នកអាចប្រើលិបិក្រម volatileប៉ុន្តែបន្ទាប់មកការអនុវត្តនឹងថយចុះបន្តិច៖ យោងតាមការធ្វើតេស្តរបស់ខ្ញុំ នៅពេលលុបធាតុ 32 លាន ល្បឿនបានថយចុះពី 500 លានលុប/វិនាទី ទៅ 450 លានលុប/វិនាទី។

ផលិតភាព

នៅក្នុងការធ្វើតេស្តសម្រាប់ការបញ្ចូលធាតុ 64 លាននិងការលុប 32 លាននៃពួកគេការប្រកួតប្រជែងរវាង std::unordered_map ហើយស្ទើរតែមិនមានតារាង hash សម្រាប់ GPU៖

std::unordered_map ចំណាយពេល 70 ms បញ្ចូល និងដកធាតុចេញ ហើយបន្ទាប់មកដោះលែងពួកគេ។ unordered_map (ការកម្ចាត់ធាតុរាប់លានត្រូវចំណាយពេលច្រើន ព្រោះនៅខាងក្នុង unordered_map ការបែងចែកអង្គចងចាំច្រើនត្រូវបានធ្វើឡើង) ។ និយាយដោយស្មោះត្រង់, std:unordered_map ការដាក់កម្រិតខុសគ្នាទាំងស្រុង។ វាគឺជាដំណើរការស៊ីភីយូតែមួយ គាំទ្រតម្លៃគន្លឹះនៃទំហំណាមួយ ដំណើរការបានល្អក្នុងអត្រាប្រើប្រាស់ខ្ពស់ និងបង្ហាញពីដំណើរការមានស្ថេរភាពបន្ទាប់ពីការលុបច្រើន។

រយៈពេលនៃតារាង hash សម្រាប់ GPU និងការទំនាក់ទំនងអន្តរកម្មវិធីគឺ 984 ms ។ នេះរាប់បញ្ចូលទាំងពេលវេលាដែលបានចំណាយក្នុងការដាក់តារាងក្នុងអង្គចងចាំ និងការលុបវា (ការបែងចែកអង្គចងចាំ 1 GB ម្តង ដែលត្រូវចំណាយពេលខ្លះក្នុង CUDA) ការបញ្ចូល និងលុបធាតុ និងការធ្វើម្តងទៀតលើពួកវា។ ច្បាប់ចម្លងទាំងអស់ទៅ និងពីអង្គចងចាំកាតវីដេអូក៏ត្រូវយកមកពិចារណាផងដែរ។

តារាង hash ខ្លួនវាចំណាយពេល 271 ms ដើម្បីបញ្ចប់។ នេះរាប់បញ្ចូលទាំងពេលវេលាដែលបានចំណាយដោយការបញ្ចូល និងលុបធាតុកាតវីដេអូ និងមិនគិតពីពេលវេលាដែលបានចំណាយក្នុងការចម្លងទៅក្នុងអង្គចងចាំ និងការធ្វើឡើងវិញនៅលើតារាងលទ្ធផលនោះទេ។ ប្រសិនបើតារាង GPU រស់នៅបានយូរ ឬប្រសិនបើតារាង hash មានទាំងស្រុងនៅក្នុងអង្គចងចាំរបស់កាតវីដេអូ (ឧទាហរណ៍ ដើម្បីបង្កើតតារាង hash ដែលនឹងត្រូវបានប្រើដោយកូដ GPU ផ្សេងទៀត និងមិនមែនខួរក្បាលកណ្តាល) បន្ទាប់មក លទ្ធផលតេស្តគឺពាក់ព័ន្ធ។

តារាង hash សម្រាប់កាតវីដេអូបង្ហាញពីដំណើរការខ្ពស់ ដោយសារការបញ្ជូនខ្ពស់ និងការប៉ារ៉ាឡែលសកម្ម។

កំហុស

ស្ថាបត្យកម្មតារាង hash មានបញ្ហាមួយចំនួនដែលត្រូវដឹង៖

• ការស៊ើបអង្កេតលីនេអ៊ែរត្រូវបានរារាំងដោយការចង្កោម ដែលបណ្តាលឱ្យគ្រាប់ចុចនៅក្នុងតារាងត្រូវបានដាក់តិចជាងឥតខ្ចោះ។
• សោមិនត្រូវបានយកចេញដោយប្រើមុខងារទេ។ delete ហើយយូរ ៗ ទៅពួកគេបានពង្រាយតុ។

ជាលទ្ធផល ដំណើរការនៃតារាង hash អាចបន្ថយបន្តិចម្តងៗ ជាពិសេសប្រសិនបើវាមានរយៈពេលយូរ ហើយមានការបញ្ចូល និងលុបច្រើន។ មធ្យោបាយមួយដើម្បីកាត់បន្ថយគុណវិបត្តិទាំងនេះគឺត្រូវបញ្ចូលឡើងវិញទៅក្នុងតារាងថ្មីមួយជាមួយនឹងអត្រាប្រើប្រាស់ទាបគួរសម ហើយត្រងចេញនូវសោរដែលបានដកចេញកំឡុងពេល rehashing ។

ដើម្បីបង្ហាញពីបញ្ហាដែលបានពិពណ៌នា ខ្ញុំនឹងប្រើកូដខាងលើដើម្បីបង្កើតតារាងមួយដែលមានធាតុ 128 លាន ហើយរង្វិលជុំតាមរយៈធាតុ 4 លានរហូតដល់ខ្ញុំបានបំពេញ 124 លានរន្ធ (អត្រាប្រើប្រាស់ប្រហែល 0,96)។ នេះគឺជាតារាងលទ្ធផល ជួរនីមួយៗគឺជាការហៅខឺណែល CUDA ដើម្បីបញ្ចូលធាតុថ្មីចំនួន 4 លានទៅក្នុងតារាងសញ្ញាមួយ៖

អត្រាប្រើប្រាស់
រយៈពេលនៃការបញ្ចូល 4 ធាតុ

0,00
11,608448 ms (361,314798 លានសោ/វិនាទី)

0,03
11,751424 ms (356,918799 លានសោ/វិនាទី)

0,06
11,942592 ms (351,205515 លានសោ/វិនាទី)

0,09
12,081120 ms (347,178429 លានសោ/វិនាទី)

0,12
12,242560 ms (342,600233 លានសោ/វិនាទី)

0,16
12,396448 ms (338,347235 លានសោ/វិនាទី)

0,19
12,533024 ms (334,660176 លានសោ/វិនាទី)

0,22
12,703328 ms (330,173626 លានសោ/វិនាទី)

0,25
12,884512 ms (325,530693 លានសោ/វិនាទី)

0,28
13,033472 ms (321,810182 លានសោ/វិនាទី)

0,31
13,239296 ms (316,807174 លានសោ/វិនាទី)

0,34
13,392448 ms (313,184256 លានសោ/វិនាទី)

0,37
13,624000 ms (307,861434 លានសោ/វិនាទី)

0,41
13,875520 ms (302,280855 លានសោ/វិនាទី)

0,44
14,126528 ms (296,909756 លានសោ/វិនាទី)

0,47
14,399328 ms (291,284699 លានសោ/វិនាទី)

0,50
14,690304 ms (285,515123 លានសោ/វិនាទី)

0,53
15,039136 ms (278,892623 លានសោ/វិនាទី)

0,56
15,478656 ms (270,973402 លានសោ/វិនាទី)

0,59
15,985664 ms (262,379092 លានសោ/វិនាទី)

0,62
16,668673 ms (251,627968 លានសោ/វិនាទី)

0,66
17,587200 ms (238,486174 លានសោ/វិនាទី)

0,69
18,690048 ms (224,413765 លានសោ/វិនាទី)

0,72
20,278816 ms (206,831789 លានសោ/វិនាទី)

0,75
22,545408 ms (186,038058 លានសោ/វិនាទី)

0,78
26,053312 ms (160,989275 លានសោ/វិនាទី)

0,81
31,895008 ms (131,503463 លានសោ/វិនាទី)

0,84
42,103294 ms (99,619378 លានសោ/វិនាទី)

0,87
61,849056 ms (67,815164 លានសោ/វិនាទី)

0,90
105,695999 ms (39,682713 លានសោ/វិនាទី)

0,94
240,204636 ms (17,461378 លានសោ/វិនាទី)

នៅពេលដែលការប្រើប្រាស់កើនឡើង ដំណើរការថយចុះ។ នេះមិនគួរឱ្យចង់បានទេក្នុងករណីភាគច្រើន។ ប្រសិនបើកម្មវិធីបញ្ចូលធាតុទៅក្នុងតារាង ហើយបន្ទាប់មកបោះបង់វា (ឧទាហរណ៍ ពេលរាប់ពាក្យក្នុងសៀវភៅ) នោះមិនមែនជាបញ្ហាទេ។ ប៉ុន្តែប្រសិនបើកម្មវិធីប្រើតារាង hash ដែលមានអាយុកាលយូរ (ឧទាហរណ៍នៅក្នុងកម្មវិធីកែក្រាហ្វិកដើម្បីរក្សាទុកផ្នែកមិនទទេនៃរូបភាពដែលអ្នកប្រើប្រាស់បញ្ចូល និងលុបព័ត៌មានញឹកញាប់) នោះឥរិយាបថនេះអាចមានបញ្ហា។

ហើយបានវាស់ជម្រៅនៃការស៊ើបអង្កេតតារាង hash បន្ទាប់ពីការបញ្ចូលចំនួន 64 លាន (កត្តាប្រើប្រាស់ 0,5)។ ជម្រៅជាមធ្យមគឺ 0,4774 ដូច្នេះសោភាគច្រើនស្ថិតនៅក្នុងរន្ធដោតល្អបំផុត ឬរន្ធមួយនៅឆ្ងាយពីទីតាំងល្អបំផុត។ ជម្រៅសំឡេងអតិបរមាគឺ 60 ។

បន្ទាប់មកខ្ញុំបានវាស់ជម្រៅនៃការស៊ើបអង្កេតលើតារាងដែលមាន 124 លានបញ្ចូល (កត្តាប្រើប្រាស់ 0,97) ។ ជម្រៅជាមធ្យមគឺ 10,1757 ហើយអតិបរមា - 6474 (!!). ការអនុវត្តការចាប់សញ្ញាលីនេអ៊ែរធ្លាក់ចុះយ៉ាងខ្លាំងក្នុងអត្រាប្រើប្រាស់ខ្ពស់។

វាជាការល្អបំផុតក្នុងការរក្សាអត្រាប្រើប្រាស់តារាង hash នេះឱ្យទាប។ ប៉ុន្តែបន្ទាប់មកយើងបង្កើនការអនុវត្តដោយចំណាយនៃការប្រើប្រាស់អង្គចងចាំ។ ជាសំណាងល្អ នៅក្នុងករណីនៃគ្រាប់ចុច និងតម្លៃ 32 ប៊ីត នេះអាចត្រូវបានរាប់ជាសុចរិត។ ប្រសិនបើក្នុងឧទាហរណ៍ខាងលើក្នុងតារាងដែលមានធាតុ 128 លាន យើងរក្សាកត្តាប្រើប្រាស់ 0,25 បន្ទាប់មកយើងអាចដាក់មិនលើសពី 32 លានធាតុនៅក្នុងវា ហើយ 96 លានរន្ធដែលនៅសល់នឹងត្រូវបាត់បង់ - 8 បៃសម្រាប់គូនីមួយៗ។ , 768 មេកាបៃនៃការចងចាំដែលបាត់បង់។

សូមចំណាំថាយើងកំពុងនិយាយអំពីការបាត់បង់អង្គចងចាំកាតវីដេអូដែលជាធនធានដ៏មានតម្លៃជាងអង្គចងចាំប្រព័ន្ធ។ ទោះបីជាកាតក្រាហ្វិកកុំព្យូទ័រទំនើបភាគច្រើនដែលគាំទ្រ CUDA មានអង្គចងចាំយ៉ាងហោចណាស់ 4 GB (នៅពេលសរសេរ NVIDIA 2080 Ti មាន 11 GB) វានៅតែមិនមែនជាការសម្រេចចិត្តដ៏ឆ្លាតវៃបំផុតក្នុងការបាត់បង់បរិមាណបែបនេះ។

ក្រោយមកខ្ញុំនឹងសរសេរបន្ថែមអំពីការបង្កើតតារាង hash សម្រាប់កាតវីដេអូដែលមិនមានបញ្ហាជាមួយនឹងការស៊ើបអង្កេតជម្រៅ ក៏ដូចជាវិធីប្រើឡើងវិញនូវរន្ធដោតដែលបានលុប។

ការវាស់វែងជម្រៅសំឡេង

ដើម្បី​កំណត់​ជម្រៅ​នៃ​ការ​សាកល្បង​នៃ​សោ ​យើង​អាច​ទាញ​សញ្ញា​របស់​សោ (សន្ទស្សន៍​តារាង​ដ៏​ល្អ​របស់​វា) ពី​សន្ទស្សន៍​តារាង​ពិត​របស់​វា​បាន៖

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

ដោយសារតែវេទមន្តនៃលេខគោលពីរបំពេញបន្ថែមរបស់ពីរ និងការពិតដែលថាសមត្ថភាពនៃតារាង hash គឺពីរទៅថាមពលនៃពីរ វិធីសាស្រ្តនេះនឹងដំណើរការសូម្បីតែនៅពេលដែលសន្ទស្សន៍គន្លឹះត្រូវបានផ្លាស់ទីទៅដើមតារាង។ ចូរយើងយកកូនសោដែលដាក់ទៅលេខ 1 ប៉ុន្តែត្រូវបានបញ្ចូលទៅក្នុងរន្ធលេខ 3។ បន្ទាប់មកសម្រាប់តារាងដែលមានទំហំ 4 យើងទទួលបាន (3 — 1) & 3ដែលស្មើនឹង ២.

សេចក្តីសន្និដ្ឋាន

ប្រសិនបើអ្នកមានសំណួរឬមតិយោបល់សូមផ្ញើអ៊ីមែលមកខ្ញុំ Twitter ឬបើកប្រធានបទថ្មី។ ឃ្លាំង.

កូដនេះត្រូវបានសរសេរក្រោមការបំផុសគំនិតពីអត្ថបទល្អៗ៖

• តារាង Hash គ្មានសោរសាមញ្ញបំផុតរបស់ពិភពលោក
• តារាង Hash Free Wait-Free ចាក់សោ

នៅពេលអនាគត ខ្ញុំនឹងបន្តសរសេរអំពីការអនុវត្តតារាង hash សម្រាប់កាតវីដេអូ និងវិភាគដំណើរការរបស់ពួកគេ។ ផែនការរបស់ខ្ញុំរួមមានការដាក់ខ្សែសង្វាក់ Robin Hood hashing និង cuckoo hashing ដោយប្រើប្រតិបត្តិការអាតូមិចនៅក្នុងរចនាសម្ព័ន្ធទិន្នន័យដែលងាយស្រួលប្រើ GPU ។

ប្រភព: www.habr.com