Github ላይ ለጥፌዋለሁ አዲስ ፕሮጀክት ቀላል የጂፒዩ ሃሽ ሠንጠረዥ.

በሰከንድ በመቶ ሚሊዮኖች የሚቆጠሩ ማስገቢያዎችን ማካሄድ የሚችል ቀላል የጂፒዩ ሃሽ ሠንጠረዥ ነው። በእኔ የNVIDIA GTX 1060 ላፕቶፕ ላይ ኮዱ 64 ሚልዮን በዘፈቀደ የመነጩ የቁልፍ እሴት ጥንዶችን በ210 ms አካባቢ ያስገባ እና 32 ሚሊዮን ጥንዶችን በ64 ሚሴ አካባቢ ያስወግዳል።

ማለትም፣ የጭን ኮምፒውተር ፍጥነት በግምት 300 ሚሊዮን ገባዎች/ሴኮንድ እና 500 ሚሊዮን ማጥፋት/ሴኮንድ ነው።

ሠንጠረዡ በ CUDA ተጽፏል, ምንም እንኳን ተመሳሳይ ዘዴ ለ HLSL ወይም GLSL ሊተገበር ይችላል. በቪዲዮ ካርድ ላይ ከፍተኛ አፈጻጸም ለማረጋገጥ ትግበራው በርካታ ገደቦች አሉት፡-

• 32-ቢት ቁልፎች ብቻ እና ተመሳሳይ እሴቶች ይከናወናሉ.
• የሃሽ ጠረጴዛው ቋሚ መጠን አለው.
• እና ይህ መጠን ከኃይል ጋር ከሁለት ጋር እኩል መሆን አለበት.

ለቁልፍ እና እሴቶች፣ ቀላል ገዳቢ ምልክት ማድረጊያ ቦታ ማስያዝ ያስፈልግዎታል (ከላይ ባለው ኮድ ይህ 0xffffffff ነው)።

የሃሽ ጠረጴዛ ያለ መቆለፊያዎች

የሃሽ ጠረጴዛው ክፍት አድራሻዎችን ይጠቀማል መስመራዊ ምርመራማለትም፣ በቀላሉ በማህደረ ትውስታ ውስጥ የተከማቸ እና የላቀ የመሸጎጫ አፈጻጸም ያለው የቁልፍ እሴት ጥንዶች ድርድር ነው። በሰንሰለት መያያዝም ተመሳሳይ ነገር ሊባል አይችልም፣ ይህም በተገናኘ ዝርዝር ውስጥ ጠቋሚ መፈለግን ያካትታል። የሃሽ ጠረጴዛ ቀላል ድርድር አባሎችን ለማከማቸት ነው። KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

የሠንጠረዡ መጠን የሁለት ኃይል እንጂ ዋና ቁጥር አይደለም, ምክንያቱም አንድ ፈጣን መመሪያ የፓው2/AND ጭምብልን ለመተግበር በቂ ነው, ነገር ግን ሞጁል ኦፕሬተር በጣም ቀርፋፋ ነው. ይህ በመስመራዊ ፍለጋ ጉዳይ ላይ አስፈላጊ ነው ፣ ምክንያቱም በመስመር ሠንጠረዥ ፍለጋ ውስጥ የእቃ ማውጫው በእያንዳንዱ ማስገቢያ ውስጥ መጠቅለል አለበት። እና በውጤቱም, የክወና ወጪ በእያንዳንዱ ማስገቢያ ውስጥ ሞዱሎ ታክሏል.

ሠንጠረዡ ለእያንዳንዱ ኤለመንት ቁልፉን እና እሴቱን ብቻ ያከማቻል እንጂ የቁልፉን ሃሽ አይደለም። ሠንጠረዡ ባለ 32-ቢት ቁልፎችን ብቻ ስለሚያከማች ሃሽ በጣም በፍጥነት ይሰላል። ከላይ ያለው ኮድ Murmur3 hash ይጠቀማል፣ ይህም ጥቂት ፈረቃዎችን፣ 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 kernel ውስጥ ሳይሆን በአስተናጋጅ ኮድ ውስጥ ይከናወናል.

በጽሑፉ ከመቆለፊያ ነጻ የሆነ ነጻ የሃሽ ጠረጴዛ እንደዚህ ያለ በመቆለፊያ የተጠበቀ የውሂብ መዋቅር እንዴት እንደሚስተካከል ይገልጻል።

ተወዳዳሪነት

ከላይ ባለው የተግባር ኮድ ቅንጥቦች ውስጥ gpu_hashtable_insert(), _lookup() и _delete() በአንድ ጊዜ አንድ ቁልፍ-እሴት ማካሄድ። እና ዝቅ gpu_hashtable_insert(), _lookup() и _delete() የጥንዶች ድርድር በትይዩ ያካሂዳል፣ እያንዳንዱ ጥንድ በተለየ የጂፒዩ የማስፈጸሚያ ክር ውስጥ፡

// 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 እና ለጂፒዩ ምንም የሃሽ ጠረጴዛ የለም ማለት ይቻላል፡-

std::unordered_map ኤለመንቶችን በማስገባት እና በማስወገድ እና ከዚያም ነፃ በማውጣት 70 ሚሴ አውጥቷል። unordered_map (በሚሊዮን የሚቆጠሩ ንጥረ ነገሮችን ማስወገድ ብዙ ጊዜ ይወስዳል ፣ ምክንያቱም በውስጡ unordered_map ብዙ የማህደረ ትውስታ ምደባዎች ተደርገዋል). እውነት ለመናገር፣ std:unordered_map ሙሉ ለሙሉ የተለያዩ ገደቦች. እሱ ነጠላ የሲፒዩ አፈፃፀም ክር ነው ፣ ማንኛውንም መጠን ያላቸውን ቁልፍ እሴቶች ይደግፋል ፣ በከፍተኛ የአጠቃቀም ፍጥነት ጥሩ ይሰራል እና ከበርካታ ስረዛ በኋላ የተረጋጋ አፈፃፀም ያሳያል።

ለጂፒዩ እና ኢንተር-ፕሮግራም ግንኙነት የሃሽ ጠረጴዛው ቆይታ 984 ሚሴ ነበር። ይህም ሰንጠረዡን በማህደረ ትውስታ ውስጥ በማስቀመጥ እና በመሰረዝ (1 ጂቢ ማህደረ ትውስታን አንድ ጊዜ መመደብ, ይህም በ CUDA ውስጥ የተወሰነ ጊዜ ይወስዳል), ኤለመንቶችን ማስገባት እና መሰረዝ እና በእነሱ ላይ በመድገም ያሳለፈውን ጊዜ ይጨምራል. ሁሉም ቅጂዎች ወደ ቪዲዮ ካርድ ማህደረ ትውስታ በተጨማሪ ግምት ውስጥ ይገባሉ.

የሃሽ ጠረጴዛው ራሱ ለማጠናቀቅ 271 ሚሴ ፈጅቷል። ይህ በቪዲዮ ካርዱ ውስጥ ኤለመንቶችን በማስገባት እና በመሰረዝ ላይ ያለውን ጊዜ ይጨምራል, እና ወደ ማህደረ ትውስታ በመገልበጥ እና በተገኘው ሰንጠረዥ ላይ በመድገም ላይ ያለውን ጊዜ ግምት ውስጥ አያስገባም. የጂፒዩ ጠረጴዛው ለረጅም ጊዜ የሚኖር ከሆነ ወይም የሃሽ ጠረጴዛው ሙሉ በሙሉ በቪዲዮ ካርዱ ማህደረ ትውስታ ውስጥ የሚገኝ ከሆነ (ለምሳሌ በማዕከላዊ ፕሮሰሰር ሳይሆን በሌላ የጂፒዩ ኮድ ጥቅም ላይ የሚውል የሃሽ ሠንጠረዥ ለመፍጠር)። የፈተና ውጤቱ ተዛማጅ ነው.

ለቪዲዮ ካርድ ያለው የሃሽ ሠንጠረዥ በከፍተኛ ፍጥነት እና ንቁ ትይዩ ምክንያት ከፍተኛ አፈፃፀም ያሳያል።

ችግሮች

የሃሽ ሠንጠረዥ አርክቴክቸር ልታስተውላቸው የሚገቡ ጥቂት ጉዳዮች አሉት፡

• መስመራዊ ፍተሻ በክላስተር ይስተጓጎላል፣ ይህ ደግሞ በሰንጠረዡ ውስጥ ያሉት ቁልፎች በትክክል ከተቀመጠው ያነሰ እንዲቀመጡ ያደርጋል።
• ተግባሩን በመጠቀም ቁልፎች አይወገዱም delete እና ከጊዜ በኋላ ጠረጴዛውን ያበላሻሉ.

በውጤቱም, የሃሽ ሠንጠረዥ አፈፃፀም ቀስ በቀስ ሊቀንስ ይችላል, በተለይም ለረጅም ጊዜ ካለ እና ብዙ ማስገቢያዎች እና መሰረዝ. እነዚህን ድክመቶች ለመቅረፍ አንዱ መንገድ ወደ አዲስ ሠንጠረዥ በመጠኑ ዝቅተኛ የመጠቀሚያ ፍጥነት እንደገና መታደስ እና በመልሶ ማቋቋም ጊዜ የተወገዱ ቁልፎችን ማጣራት ነው።

የተገለጹትን ጉዳዮች በምሳሌ ለማስረዳት፣ 128 ሚሊዮን ኤለመንቶችን የያዘ ሠንጠረዥ ለመፍጠር እና 4 ሚሊዮን ቦታዎችን እስከምሞላ ድረስ (የአጠቃቀም መጠን 124) እስኪሆን ድረስ በ0,96 ሚሊዮን ኤለመንቶች በኩል ሰንጠረዥ ለመፍጠር ከላይ ያለውን ኮድ እጠቀማለሁ። የውጤት ሠንጠረዥ እነሆ፣ እያንዳንዱ ረድፍ 4 ሚሊዮን አዳዲስ ንጥረ ነገሮችን ወደ አንድ የሃሽ ሠንጠረዥ ለማስገባት የCUDA kernel ጥሪ ነው።

የአጠቃቀም መጠን
የማስገባት ጊዜ 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 ሜባ የጠፋ ማህደረ ትውስታ።

እባክዎን ስለ ቪዲዮ ካርድ ማህደረ ትውስታ ማጣት እየተነጋገርን ነው, ይህም ከስርዓት ማህደረ ትውስታ የበለጠ ዋጋ ያለው ሃብት ነው. ምንም እንኳን 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 ጋር እኩል ነው.

መደምደሚያ

ጥያቄዎች ወይም አስተያየቶች ካሉዎት እባክዎን በኢሜል ይላኩልኝ። Twitter ወይም አዲስ ርዕስ ይክፈቱ ማከማቻዎች.

ይህ ኮድ ከጥሩ መጣጥፎች ተመስጦ ነው የተጻፈው፡-

• የአለማችን ቀላሉ ከመቆለፊያ-ነጻ የሃሽ ሠንጠረዥ
• ከመቆለፊያ ነጻ የሆነ ነጻ የሃሽ ጠረጴዛ

ለወደፊቱ, ስለ ሃሽ ሰንጠረዥ አተገባበር ለቪዲዮ ካርዶች መፃፍ እና አፈፃፀማቸውን መተንተን እቀጥላለሁ. የእኔ እቅዶች ለጂፒዩ ተስማሚ በሆኑ የመረጃ ቋቶች ውስጥ የአቶሚክ ኦፕሬሽኖችን በመጠቀም ሰንሰለት፣ ሮቢን ሁድ ሃሺንግ እና ኩክኩ ሃሺንግ ያካትታሉ።

ምንጭ: hab.com