मैंने जीथब पर एक साधारण जीएपी हैश टेबल नामक एक नई परियोजना पोस्ट की ।

यह GPU के लिए एक सरल हैश तालिका है, जो प्रति सेकंड सैकड़ों लाखों आवेषणों को संसाधित करने में सक्षम है। NVIDIA GTX 1060 के साथ मेरे लैपटॉप पर, कोड ने लगभग 210 एमएस में 64 मिलियन यादृच्छिक रूप से उत्पन्न की-वैल्यू जोड़े सम्मिलित किए और लगभग 64 एमएस में 32 मिलियन जोड़े निकाले।

यही है, लैपटॉप पर गति लगभग 300 मिलियन आवेषण / सेकंड और 500 मिलियन निष्कासन / सेकंड है।

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

केवल 32-बिट कुंजियाँ और समान मान संसाधित किए जाते हैं।
हैश टेबल का एक निश्चित आकार होता है।
और यह आकार दो डिग्री के बराबर होना चाहिए।

कुंजियों और मूल्यों के लिए, आपको एक सरल परिसीमन मार्कर आरक्षित करना होगा (ऊपर कोड में यह 0xffffffff है)।

बिना ताले के हैश टेबल

हैश तालिका रेखीय संवेदन के साथ खुले पते का उपयोग करती है, अर्थात, यह कुंजी-मूल्य जोड़े की एक सरणी है जो स्मृति में संग्रहीत होती है और इसमें उत्कृष्ट कैश प्रदर्शन होता है। यह चेनिंग के मामले में नहीं है, जिसका अर्थ है कि एक लिंक्ड सूची में एक पॉइंटर की तलाश करना। एक हैश तालिका एक सरल सरणी भंडारण तत्व है KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

तालिका का आकार शक्ति में दो के बराबर है, और एक अभाज्य संख्या नहीं है, क्योंकि pow2 / AND-mask का उपयोग करने के लिए, एक त्वरित निर्देश पर्याप्त है, और मॉड्यूल ऑपरेटर बहुत धीमा है। यह रैखिक संवेदन के मामले में महत्वपूर्ण है, क्योंकि तालिका में एक रेखीय खोज के बाद, प्रत्येक स्लॉट में स्लॉट इंडेक्स को लपेटा जाना चाहिए। और परिणामस्वरूप, प्रत्येक स्लॉट में ऑपरेशन की लागत को मोडुलो जोड़ा जाता है।

तालिका प्रत्येक आइटम के लिए केवल कुंजी और मान संग्रहीत करती है, कुंजी हैश नहीं। चूंकि तालिका में केवल 32-बिट कुंजियाँ संग्रहीत हैं, इसलिए हैश की गणना बहुत तेज़ी से की जाती है। ऊपर दिया गया कोड मुरमुर 3 हैश का उपयोग करता है, जो केवल कुछ ऑफसेट, एक्सओआर और गुणन करता है।

हैश टेबल एक लॉक प्रोटेक्शन तकनीक का उपयोग करता है जो मेमोरी प्लेसमेंट ऑर्डर पर निर्भर नहीं करता है। यहां तक कि अगर कुछ लिखने के संचालन ऐसे अन्य संचालन के आदेश का उल्लंघन करते हैं, तो हैश टेबल अभी भी सही स्थिति बनाए रखेगा। हम इस बारे में नीचे बात करेंगे। तकनीक वीडियो कार्ड के साथ शानदार काम करती है जिसमें हजारों धागे प्रतिस्पर्धा करते हैं।

हैश तालिका में कुंजी और मान को खाली करने के लिए प्रारंभ किया जाता है।

कोड को संशोधित किया जा सकता है ताकि यह 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 691 एमएस खर्च किया गया unordered_map(लाखों तत्वों से रिहाई में बहुत समय लगता है, क्योंकि unordered_mapकई मेमोरी आवंटन अंदर किए जाते हैं )। ईमानदारी से, std:unordered_mapपूरी तरह से अलग सीमाएं हैं। यह निष्पादन का एक एकल सीपीयू-धागा है, यह किसी भी आकार के कुंजी-मूल्यों का समर्थन करता है, उच्च उपयोग दरों पर अच्छी तरह से काम करता है और कई विलोपन के बाद स्थिर प्रदर्शन दिखाता है।

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

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

वीडियो कार्ड के लिए हैश तालिका अपने उच्च बैंडविड्थ और सक्रिय समानांतरकरण के कारण उच्च प्रदर्शन को प्रदर्शित करता है।

नुकसान

हैश टेबल आर्किटेक्चर को ध्यान में रखने के लिए कई मुद्दे हैं:

क्लस्टरिंग रैखिक जांच के साथ हस्तक्षेप करता है, जिसके कारण तालिका में चाबियाँ आदर्श से बहुत दूर हैं।
फ़ंक्शन का उपयोग करके कुंजी को हटाया नहीं जाता है deleteऔर समय के साथ तालिका को अव्यवस्थित किया जाता है।

नतीजतन, हैश तालिका का प्रदर्शन धीरे-धीरे कम हो सकता है, खासकर अगर यह लंबे समय तक मौजूद है और इसमें कई सम्मिलन और विलोपन किए जाते हैं। इन कमियों को कम करने का एक तरीका एक नई तालिका को काफी कम उपयोग की दर के साथ फिर से जोड़ना और रीशश करने पर रिमोट कीज़ को फ़िल्टर करना है।

वर्णित समस्याओं का वर्णन करने के लिए, मैं उपरोक्त कोड का उपयोग 128 मिलियन तत्वों के लिए एक तालिका बनाने के लिए करता हूं, मैं 124 मिलियन स्लॉट भरने तक 4 मिलियन तत्वों को चक्रवाती रूप से सम्मिलित करूंगा (उपयोग लगभग 0.96 है)। यहाँ परिणाम तालिका है, प्रत्येक पंक्ति CUDA कोर के लिए एक हैश तालिका में 4 मिलियन नए तत्वों के सम्मिलन के साथ एक कॉल है:

प्रयोग दर	सम्मिलन 4 194 304 तत्वों की अवधि
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 एमबी की खोई गई मेमोरी।

कृपया ध्यान दें कि हम वीडियो कार्ड मेमोरी के नुकसान के बारे में बात कर रहे हैं, जो सिस्टम मेमोरी की तुलना में अधिक मूल्यवान संसाधन है। यद्यपि अधिकांश आधुनिक डेस्कटॉप ग्राफिक्स कार्ड जो CUDA का समर्थन करते हैं, उनमें कम से कम 4 जीबी मेमोरी होती है (लेखन के समय, NVIDIA 2080 तिवारी में 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 के लिए सरल हैश तालिका