ഞാൻ അത് Github-ൽ പോസ്റ്റ് ചെയ്തു പുതിയ പ്രോജക്റ്റ് ഒരു ലളിതമായ ജിപിയു ഹാഷ് ടേബിൾ.

സെക്കന്റിൽ ദശലക്ഷക്കണക്കിന് ഇൻസെർട്ടുകൾ പ്രോസസ്സ് ചെയ്യാൻ കഴിവുള്ള ഒരു ലളിതമായ ജിപിയു ഹാഷ് ടേബിളാണിത്. എന്റെ NVIDIA GTX 1060 ലാപ്‌ടോപ്പിൽ, കോഡ് ക്രമരഹിതമായി ജനറേറ്റുചെയ്‌ത 64 ദശലക്ഷം കീ-മൂല്യം ജോഡികളെ ഏകദേശം 210 ms-ൽ തിരുകുകയും ഏകദേശം 32 ms-ൽ 64 ദശലക്ഷം ജോഡികൾ നീക്കം ചെയ്യുകയും ചെയ്യുന്നു.

അതായത്, ഒരു ലാപ്‌ടോപ്പിലെ വേഗത ഏകദേശം 300 ദശലക്ഷം ഇൻസെർട്ടുകൾ/സെക്കന്റും 500 ദശലക്ഷം ഡിലീറ്റുകൾ/സെക്കന്റുമാണ്.

എച്ച്എൽഎസ്എൽ അല്ലെങ്കിൽ ജിഎൽഎസ്എൽ എന്നിവയിൽ ഇതേ സാങ്കേതികത പ്രയോഗിക്കാമെങ്കിലും പട്ടിക CUDA-യിലാണ് എഴുതിയിരിക്കുന്നത്. ഒരു വീഡിയോ കാർഡിൽ ഉയർന്ന പ്രകടനം ഉറപ്പാക്കുന്നതിന് നടപ്പിലാക്കുന്നതിന് നിരവധി പരിമിതികളുണ്ട്:

• 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() ജോഡികളുടെ ഒരു നിര സമാന്തരമായി പ്രോസസ്സ് ചെയ്യുക, ഓരോ ജോഡിയും പ്രത്യേക ജിപിയു എക്സിക്യൂഷൻ ത്രെഡിൽ:

// 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 ms ചെലവഴിച്ചു മൂലകങ്ങൾ ചേർക്കുകയും നീക്കം ചെയ്യുകയും തുടർന്ന് അവയെ സ്വതന്ത്രമാക്കുകയും ചെയ്തു unordered_map (ദശലക്ഷക്കണക്കിന് മൂലകങ്ങളിൽ നിന്ന് രക്ഷപ്പെടാൻ വളരെയധികം സമയമെടുക്കും, കാരണം ഉള്ളിൽ unordered_map ഒന്നിലധികം മെമ്മറി അലോക്കേഷനുകൾ നടത്തുന്നു). സത്യസന്ധമായി പറഞ്ഞാൽ, std:unordered_map തികച്ചും വ്യത്യസ്തമായ നിയന്ത്രണങ്ങൾ. ഇത് നിർവ്വഹണത്തിന്റെ ഒരു സിപിയു ത്രെഡ് ആണ്, ഏത് വലുപ്പത്തിലുമുള്ള കീ-മൂല്യങ്ങളെ പിന്തുണയ്ക്കുന്നു, ഉയർന്ന ഉപയോഗ നിരക്കിൽ മികച്ച പ്രകടനം നടത്തുന്നു, ഒന്നിലധികം ഇല്ലാതാക്കലുകൾക്ക് ശേഷം സ്ഥിരതയുള്ള പ്രകടനം കാണിക്കുന്നു.

GPU-നും ഇന്റർ-പ്രോഗ്രാം ആശയവിനിമയത്തിനുമുള്ള ഹാഷ് ടേബിളിന്റെ ദൈർഘ്യം 984 ms ആയിരുന്നു. ടേബിൾ മെമ്മറിയിൽ സ്ഥാപിക്കുന്നതിനും അത് ഇല്ലാതാക്കുന്നതിനും (ഒരു തവണ 1 GB മെമ്മറി അനുവദിക്കൽ, CUDA-യിൽ കുറച്ച് സമയമെടുക്കും), ഘടകങ്ങൾ ചേർക്കുകയും ഇല്ലാതാക്കുകയും, അവയ്ക്ക് മുകളിൽ ആവർത്തിക്കുകയും ചെയ്യുന്ന സമയം ഇതിൽ ഉൾപ്പെടുന്നു. വീഡിയോ കാർഡ് മെമ്മറിയിലേക്കുള്ള എല്ലാ പകർപ്പുകളും കണക്കിലെടുക്കുന്നു.

ഹാഷ് ടേബിൾ തന്നെ പൂർത്തിയാക്കാൻ 271 ms എടുത്തു. വീഡിയോ കാർഡ് ഘടകങ്ങൾ ചേർക്കുകയും ഇല്ലാതാക്കുകയും ചെയ്യുന്ന സമയം ഇതിൽ ഉൾപ്പെടുന്നു, കൂടാതെ മെമ്മറിയിലേക്ക് പകർത്തുന്നതിനും ഫലമായുണ്ടാകുന്ന പട്ടികയിൽ ആവർത്തിക്കുന്നതിനും ചെലവഴിച്ച സമയം കണക്കിലെടുക്കുന്നില്ല. ജിപിയു ടേബിൾ ദീർഘകാലം നിലനിൽക്കുകയാണെങ്കിൽ, അല്ലെങ്കിൽ വീഡിയോ കാർഡിന്റെ മെമ്മറിയിൽ ഹാഷ് ടേബിൾ മുഴുവനായും അടങ്ങിയിട്ടുണ്ടെങ്കിൽ (ഉദാഹരണത്തിന്, മറ്റ് ജിപിയു കോഡ് ഉപയോഗിക്കുന്ന ഒരു ഹാഷ് ടേബിൾ സൃഷ്ടിക്കാൻ, സെൻട്രൽ പ്രോസസർ അല്ല), തുടർന്ന് പരിശോധന ഫലം പ്രസക്തമാണ്.

ഉയർന്ന ത്രൂപുട്ടും സജീവമായ സമാന്തരവൽക്കരണവും കാരണം ഒരു വീഡിയോ കാർഡിനുള്ള ഹാഷ് ടേബിൾ ഉയർന്ന പ്രകടനം പ്രകടമാക്കുന്നു.

അസൗകര്യങ്ങൾ

ഹാഷ് ടേബിൾ ആർക്കിടെക്ചറിന് അറിഞ്ഞിരിക്കേണ്ട ചില പ്രശ്നങ്ങളുണ്ട്:

• ലീനിയർ പ്രോബിംഗ് ക്ലസ്റ്ററിംഗ് തടസ്സപ്പെടുത്തുന്നു, ഇത് ടേബിളിലെ കീകൾ പൂർണ്ണതയിൽ കുറവായി സ്ഥാപിക്കാൻ കാരണമാകുന്നു.
• ഫംഗ്ഷൻ ഉപയോഗിച്ച് കീകൾ നീക്കം ചെയ്യപ്പെടുന്നില്ല 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 ദശലക്ഷം ഇൻസേർട്ടുകളുള്ള ഒരു ടേബിളിൽ ഞാൻ പ്രോബിംഗ് ഡെപ്ത് അളന്നു (ഉപയോഗ ഘടകം 0,97). ശരാശരി ആഴം ഇതിനകം 10,1757 ആയിരുന്നു, പരമാവധി - 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 ന് തുല്യമാണ്.

തീരുമാനം

നിങ്ങൾക്ക് ചോദ്യങ്ങളോ അഭിപ്രായങ്ങളോ ഉണ്ടെങ്കിൽ, ദയവായി എനിക്ക് ഇമെയിൽ ചെയ്യുക ട്വിറ്റർ അല്ലെങ്കിൽ ഒരു പുതിയ വിഷയം തുറക്കുക ശേഖരങ്ങൾ.

മികച്ച ലേഖനങ്ങളിൽ നിന്ന് പ്രചോദനം ഉൾക്കൊണ്ടാണ് ഈ കോഡ് എഴുതിയത്:

• ലോകത്തിലെ ഏറ്റവും ലളിതമായ ലോക്ക്-ഫ്രീ ഹാഷ് ടേബിൾ
• ഒരു ലോക്ക്-ഫ്രീ വെയ്റ്റ്-ഫ്രീ ഹാഷ് ടേബിൾ

ഭാവിയിൽ, വീഡിയോ കാർഡുകൾക്കായുള്ള ഹാഷ് ടേബിൾ നിർവ്വഹണങ്ങളെക്കുറിച്ച് ഞാൻ എഴുതുന്നത് തുടരുകയും അവയുടെ പ്രകടനം വിശകലനം ചെയ്യുകയും ചെയ്യും. എന്റെ പ്ലാനുകളിൽ GPU ഫ്രണ്ട്‌ലി ആയ ഡാറ്റ സ്ട്രക്ച്ചറുകളിലെ ആറ്റോമിക് ഓപ്പറേഷനുകൾ ഉപയോഗിച്ച് ചെയിനിംഗ്, റോബിൻ ഹുഡ് ഹാഷിംഗ്, കുക്കൂ ഹാഷിംഗ് എന്നിവ ഉൾപ്പെടുന്നു.

അവലംബം: www.habr.com