मैले Github मा पोस्ट गरें नयाँ परियोजना एक साधारण GPU ह्यास तालिका.

यो एक साधारण GPU ह्यास तालिका हो जुन प्रति सेकेन्ड लाखौं इन्सर्टहरू प्रशोधन गर्न सक्षम छ। मेरो NVIDIA GTX 1060 ल्यापटपमा, कोडले लगभग 64 ms मा 210 मिलियन अनियमित रूपमा उत्पन्न कुञ्जी-मान जोडीहरू घुसाउँछ र लगभग 32 ms मा 64 मिलियन जोडीहरू हटाउँछ।

अर्थात्, ल्यापटपको गति लगभग ३० करोड इन्सर्ट/सेकेन्ड र ५० करोड डिलिट/सेकेन्ड हुन्छ।

तालिका CUDA मा लेखिएको छ, यद्यपि उही प्रविधि HLSL वा GLSL मा लागू गर्न सकिन्छ। भिडियो कार्डमा उच्च प्रदर्शन सुनिश्चित गर्न कार्यान्वयनमा धेरै सीमाहरू छन्:

• केवल 32-बिट कुञ्जीहरू र समान मानहरू प्रशोधन गरिएका छन्।
• ह्यास तालिकाको निश्चित आकार छ।
• र यो साइज पावरमा दुई बराबर हुनुपर्छ।

कुञ्जीहरू र मानहरूको लागि, तपाईंले एक साधारण डेलिमिटर मार्कर आरक्षित गर्न आवश्यक छ (माथिको कोडमा यो 0xffffffff हो)।

ताला बिना ह्यास टेबल

ह्यास तालिकाले खुला ठेगाना प्रयोग गर्दछ रैखिक जाँच, त्यो हो, यो केवल कुञ्जी-मान जोडीहरूको एरे हो जुन मेमोरीमा भण्डार गरिएको छ र उच्च क्यास प्रदर्शन छ। चेनिङको लागि पनि यही भन्न सकिँदैन, जसमा लिङ्क गरिएको सूचीमा सूचक खोज्नु समावेश छ। ह्यास तालिका एक साधारण सरणी भण्डारण तत्व हो KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

तालिकाको आकार दुईको पावर हो, प्राइम नम्बर होइन, किनभने pow2/AND मास्क लागू गर्नको लागि एउटा द्रुत निर्देशन पर्याप्त छ, तर मोडुलस अपरेटर धेरै ढिलो छ। यो रैखिक जाँचको मामलामा महत्त्वपूर्ण छ, किनकि रेखीय तालिका लुकअपमा प्रत्येक स्लटमा स्लट अनुक्रमणिका बेरिएको हुनुपर्छ। र नतिजाको रूपमा, सञ्चालनको लागत प्रत्येक स्लटमा मोड्युलो थपिएको छ।

तालिकाले प्रत्येक तत्वको लागि कुञ्जी र मान मात्र भण्डारण गर्छ, कुञ्जीको ह्यास होइन। तालिकाले 32-बिट कुञ्जीहरू मात्र भण्डारण गर्ने भएकोले, ह्यास धेरै छिटो गणना गरिन्छ। माथिको कोडले Murmur3 ह्यास प्रयोग गर्दछ, जसले केही शिफ्टहरू, XORs र गुणन मात्र गर्छ।

ह्यास तालिकाले मेमोरी अर्डरबाट स्वतन्त्र लक सुरक्षा प्रविधिहरू प्रयोग गर्दछ। यदि केहि लेखन अपरेसनहरूले अन्य त्यस्ता अपरेसनहरूको क्रमलाई बाधा पुर्‍याउँछ भने, ह्यास तालिकाले अझै पनि सही स्थिति कायम राख्नेछ। हामी तल यस बारे कुरा गर्नेछौं। यो प्रविधिले भिडियो कार्डहरूसँग राम्रो काम गर्छ जुन एकैसाथ हजारौं थ्रेडहरू चलाउँछ।

ह्यास तालिकामा कुञ्जीहरू र मानहरू खाली गर्न प्रारम्भिक छन्।

कोड 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 मा, कर्नेललाई कल गर्ने कुनै पनि अपरेसनले सधैं अर्को कर्नेल कल अघि पूरा हुन्छ (कम्तिमा एउटा थ्रेड भित्र। विभिन्न थ्रेडहरूमा, कर्नेलहरू समानान्तरमा कार्यान्वयन गरिन्छ)। माथिको उदाहरणमा, यदि तपाइँ एक कर्नेलसँग कल गर्नुहुन्छ A/0 B/1 A/2 C/3, र अर्को संग A/4, त्यसपछि कुञ्जी A मूल्य प्राप्त हुनेछ 4.

अब कुरा गरौं कि कार्यहरू गर्नुपर्छ lookup() и delete() ह्यास तालिकामा जोडीहरूको एर्रेमा सादा वा अस्थिर सूचक प्रयोग गर्नुहोस्। CUDA दस्तावेजीकरण भनिएको छ कि:

कम्पाइलरले विश्वव्यापी वा साझा मेमोरीमा पढ्ने र लेख्ने अनुकूलन गर्न छनौट गर्न सक्छ... यी अप्टिमाइजेसनहरू कुञ्जी शब्द प्रयोग गरेर असक्षम गर्न सकिन्छ। volatile: ... यस चरको कुनै पनि सन्दर्भलाई वास्तविक मेमोरी पढ्न वा लेख्ने निर्देशनमा कम्पाइल गरिएको छ।

शुद्धता विचारहरू आवेदन आवश्यक छैन volatile। यदि कार्यान्वयन थ्रेडले अघिल्लो पढ्ने अपरेसनबाट क्यास गरिएको मान प्रयोग गर्छ भने, त्यसले थोरै पुरानो जानकारी प्रयोग गर्नेछ। तर अझै, यो कर्नेल कलको निश्चित क्षणमा ह्यास तालिकाको सही अवस्थाबाट जानकारी हो। यदि तपाइँ नवीनतम जानकारी प्रयोग गर्न आवश्यक छ भने, तपाइँ अनुक्रमणिका प्रयोग गर्न सक्नुहुन्छ volatile, तर त्यसपछि प्रदर्शन थोरै घट्नेछ: मेरो परीक्षण अनुसार, 32 मिलियन तत्वहरू मेटाउँदा, गति 500 ​​मिलियन मेटाउने/सेकेन्ड 450 मिलियन मेटाउने/सेकेन्डमा घट्यो।

उत्पादकता

परीक्षणमा 64 मिलियन तत्वहरू घुसाउने र तीमध्ये 32 मिलियन मेटाउने, बीच प्रतिस्पर्धा std::unordered_map र GPU को लागि लगभग कुनै ह्यास तालिका छैन:

std::unordered_map तत्वहरू सम्मिलित गर्न र हटाउन र त्यसपछि तिनीहरूलाई मुक्त गर्न 70 ms खर्च गरियो unordered_map (लाखौं तत्वहरूबाट छुटकारा पाउन धेरै समय लाग्छ, किनभने भित्र unordered_map धेरै मेमोरी आवंटन गरिएको छ)। इमानदारीपूर्वक भन्नु पर्दा, std:unordered_map पूर्णतया फरक प्रतिबन्धहरू। यो कार्यान्वयनको एकल CPU थ्रेड हो, कुनै पनि आकारको मुख्य-मानहरूलाई समर्थन गर्दछ, उच्च उपयोग दरहरूमा राम्रो प्रदर्शन गर्दछ, र धेरै मेटाइसकेपछि स्थिर प्रदर्शन देखाउँदछ।

GPU र अन्तर-कार्यक्रम संचारको लागि ह्यास तालिकाको अवधि 984 ms थियो। यसमा तालिकालाई मेमोरीमा राख्ने र मेटाउने समय (1 GB मेमोरी एक पटक आवंटित गर्ने, जुन CUDA मा केही समय लाग्छ), तत्वहरू घुसाउने र मेटाउने, र तिनीहरूलाई दोहोर्याउने समय समावेश गर्दछ। भिडियो कार्ड मेमोरीमा र बाट सबै प्रतिलिपिहरू पनि खातामा लिइन्छ।

ह्यास तालिका आफैले पूरा गर्न 271 एमएस लियो। यसमा भिडियो कार्डले तत्वहरू घुसाउन र मेटाउन बिताएको समय समावेश गर्दछ, र मेमोरीमा प्रतिलिपि गर्न र परिणाम तालिकामा दोहोर्याउन बिताएको समयलाई ध्यानमा राख्दैन। यदि GPU तालिका लामो समयसम्म जीवित छ भने, वा यदि ह्यास तालिका पूर्ण रूपमा भिडियो कार्डको मेमोरीमा समावेश छ (उदाहरणका लागि, ह्यास तालिका सिर्जना गर्न जुन अन्य GPU कोड द्वारा प्रयोग गरिनेछ र केन्द्रीय प्रोसेसरले होइन), त्यसपछि। परीक्षा परिणाम सान्दर्भिक छ।

भिडियो कार्डको लागि ह्यास तालिकाले उच्च थ्रुपुट र सक्रिय समानान्तरको कारणले उच्च प्रदर्शन देखाउँछ।

कमीकमजोरी

ह्यास टेबल आर्किटेक्चरसँग सचेत हुन केही मुद्दाहरू छन्:

• रैखिक जाँचलाई क्लस्टरिङले बाधा पुर्‍याएको छ, जसले तालिकामा कुञ्जीहरूलाई पूर्ण रूपमा कम राख्छ।
• प्रकार्य प्रयोग गरेर कुञ्जीहरू हटाइएको छैन delete र समय बित्दै जाँदा तिनीहरूले टेबललाई अव्यवस्थित पार्छन्।

नतिजाको रूपमा, ह्यास तालिकाको प्रदर्शन बिस्तारै घट्न सक्छ, विशेष गरी यदि यो लामो समयको लागि अवस्थित छ र धेरै घुसाउने र मेटाउने छ। यी बेफाइदाहरूलाई कम गर्ने एउटा तरिका एकदम कम उपयोग दरको साथ नयाँ तालिकामा पुन: ह्यास गर्नु र पुन: ह्यासिङको क्रममा हटाइएका कुञ्जीहरूलाई फिल्टर गर्नु हो।

वर्णन गरिएका मुद्दाहरू चित्रण गर्न, मैले 128 मिलियन स्लटहरू (लगभग 4 को उपयोग दर) भर्न नसकेसम्म 124 मिलियन तत्वहरू र 0,96 मिलियन तत्वहरू मार्फत लुप गर्नको लागि माथिको कोड प्रयोग गर्नेछु। यहाँ परिणाम तालिका छ, प्रत्येक पङ्क्ति एक ह्यास तालिकामा 4 मिलियन नयाँ तत्वहरू सम्मिलित गर्न CUDA कर्नेल कल हो:

उपयोग दर
सम्मिलन अवधि 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 मिलियन कुञ्जी/सेकेन्ड।)

जति उपयोगिता बढ्छ, कार्यसम्पादन घट्दै जान्छ। यो धेरै अवस्थामा वांछनीय छैन। यदि एप्लिकेसनले तालिकामा तत्वहरू सम्मिलित गर्छ र त्यसपछि तिनीहरूलाई खारेज गर्छ (उदाहरणका लागि, पुस्तकमा शब्दहरू गणना गर्दा), तब यो समस्या होइन। तर यदि एप्लिकेसनले लामो समयसम्म चल्ने ह्यास तालिका प्रयोग गर्दछ (उदाहरणका लागि, छविहरूको गैर-खाली भागहरू भण्डारण गर्न ग्राफिक्स सम्पादकमा जहाँ प्रयोगकर्ताले बारम्बार जानकारी घुसाउँदछ र मेटाउँछ), तब यो व्यवहार समस्याग्रस्त हुन सक्छ।

र 64 मिलियन सम्मिलितहरू (उपयोग कारक 0,5) पछि ह्यास टेबल प्रोबिङ गहिराइ मापन गरियो। औसत गहिराई 0,4774 थियो, त्यसैले धेरैजसो कुञ्जीहरू या त उत्तम सम्भावित स्लटमा थिए वा उत्तम स्थानबाट एक स्लट टाढा थिए। अधिकतम ध्वनि गहिराई 60 थियो।

त्यसपछि मैले १२४ मिलियन इन्सर्टहरू (उपयोग कारक ०.९७) भएको टेबुलमा प्रोबिङ गहिराइ मापन गरें। औसत गहिराई पहिले नै 124 थियो, र अधिकतम - 6474 (!!) रैखिक सेन्सिङ कार्यसम्पादन उच्च उपयोग दरहरूमा उल्लेखनीय रूपमा घट्छ।

यो ह्यास तालिकाको उपयोग दर कम राख्न सबै भन्दा राम्रो छ। तर त्यसपछि हामी मेमोरी खपतको खर्चमा प्रदर्शन बढाउँछौं। सौभाग्य देखि, 32-बिट कुञ्जी र मान को मामला मा, यो जायज हुन सक्छ। यदि माथिको उदाहरणमा, 128 मिलियन तत्वहरू भएको तालिकामा, हामीले 0,25 को उपयोगिता कारक राख्छौं, त्यसपछि हामी यसमा 32 मिलियन भन्दा बढी तत्वहरू राख्न सक्दैनौं, र बाँकी 96 मिलियन स्लटहरू हराउनेछन् - प्रत्येक जोडीको लागि 8 बाइटहरू। , 768 MB हराएको मेमोरी।

कृपया ध्यान दिनुहोस् कि हामी भिडियो कार्ड मेमोरीको हानिको बारेमा कुरा गर्दैछौं, जुन प्रणाली मेमोरी भन्दा बढी मूल्यवान स्रोत हो। यद्यपि धेरै आधुनिक डेस्कटप ग्राफिक्स कार्डहरू जसले CUDA लाई समर्थन गर्दछ कम्तिमा 4 GB मेमोरी छ (लेख्ने समयमा, NVIDIA 2080 Ti मा 11 GB छ), यो अझै पनि यस्तो रकम गुमाउनु बुद्धिमानी निर्णय हुनेछैन।

पछि म भिडियो कार्डहरूको लागि ह्यास तालिकाहरू सिर्जना गर्ने बारे थप लेख्नेछु जसमा गहिराइ जाँच गर्न समस्या छैन, साथै मेटाइएका स्लटहरू पुन: प्रयोग गर्ने तरिकाहरू।

ध्वनि गहिराई मापन

कुञ्जीको जाँच गहिराइ निर्धारण गर्न, हामी यसको वास्तविक तालिका अनुक्रमणिकाबाट कुञ्जीको ह्यास (यसको आदर्श तालिका अनुक्रमणिका) निकाल्न सक्छौं:

// 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 को बराबर छ।

निष्कर्षमा

यदि तपाइँसँग प्रश्न वा टिप्पणीहरू छन् भने, कृपया मलाई इमेल गर्नुहोस् twitter वा नयाँ विषय खोल्नुहोस् भण्डारहरू.

यो कोड उत्कृष्ट लेखबाट प्रेरणा अन्तर्गत लेखिएको थियो:

• संसारको सबैभन्दा सरल लक-मुक्त ह्यास तालिका
• एक लक-रहित प्रतीक्षा-मुक्त ह्यास तालिका

भविष्यमा, म भिडियो कार्डहरूको लागि ह्यास तालिका कार्यान्वयनको बारेमा लेख्न जारी राख्नेछु र तिनीहरूको कार्यसम्पादनको विश्लेषण गर्नेछु। मेरा योजनाहरूमा चेनिङ, रोबिन हुड ह्यासिङ, र GPU अनुकूल डेटा संरचनाहरूमा परमाणु अपरेसनहरू प्रयोग गरेर कोयल ह्यासिङ समावेश छ।

स्रोत: www.habr.com