මම ඒක Github එකේ දැම්මා නව ව්‍යාපෘතිය සරල GPU හැෂ් වගුවක්.

එය තත්පරයකට ඇතුළු කිරීම් මිලියන සිය ගණනක් සැකසීමට හැකි සරල GPU හැෂ් වගුවකි. මගේ NVIDIA GTX 1060 ලැප්ටොප් පරිගණකයේ, කේතය අහඹු ලෙස ජනනය කරන ලද යතුරු වටිනාකම් යුගල මිලියන 64ක් ms 210කින් පමණ ඇතුළු කරන අතර ms 32කින් පමණ යුගල මිලියන 64ක් ඉවත් කරයි.

එනම්, ලැප්ටොප් එකක වේගය ආසන්න වශයෙන් මිලියන 300 ඇතුළු කිරීම්/තත්පර සහ මිලියන 500 මකාදැමීම්/තත්පර වේ.

HLSL හෝ GLSL සඳහා එකම තාක්ෂණය යෙදිය හැකි වුවද වගුව CUDA හි ලියා ඇත. වීඩියෝ කාඩ්පතක ඉහළ කාර්යසාධනයක් සහතික කිරීම සඳහා ක්රියාත්මක කිරීම සඳහා සීමාවන් කිහිපයක් තිබේ:

සකසනු ලබන්නේ 32-bit යතුරු සහ එකම අගයන් පමණි.
හැෂ් වගුව ස්ථාවර ප්‍රමාණයක් ඇත.
තවද මෙම ප්රමාණය බලයට දෙකකට සමාන විය යුතුය.

යතුරු සහ අගයන් සඳහා, ඔබට සරල පරිසීමක සලකුණක් වෙන් කරවා ගත යුතුය (ඉහත කේතයේ මෙය 0xffffffff වේ).

අගුල් නොමැතිව හැෂ් මේසය

හැෂ් වගුව විවෘත ලිපින භාවිතා කරයි රේඛීය විමර්ශනය, එනම්, එය හුදෙක් මතකයේ ගබඩා කර ඇති සහ උසස් හැඹිලි කාර්ය සාධනයක් ඇති යතුරු-අගය යුගල අරාවකි. සම්බන්ධිත ලැයිස්තුවක පොයින්ටරයක් සෙවීම ඇතුළත් වන දම්වැල සම්බන්ධයෙන් ද එයම කිව නොහැක. හැෂ් වගුවක් යනු මූලද්‍රව්‍ය ගබඩා කරන සරල අරාවකි KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

මේසයේ ප්‍රමාණය දෙකේ බලයක් මිස ප්‍රථමක සංඛ්‍යාවක් නොවේ, මන්ද pow2/AND මාස්ක් එක යෙදීමට එක් වේගවත් උපදෙස් ප්‍රමාණවත් වන නමුත් මාපාංක ක්‍රියාකරු වඩා මන්දගාමී වේ. රේඛීය වගු සෙවීමේ දී, එක් එක් තව් තුළ තව් දර්ශකය ඔතා තිබිය යුතු බැවින්, රේඛීය පරීක්‍ෂණයේ දී මෙය වැදගත් වේ. එහි ප්‍රතිඵලයක් වශයෙන්, මෙහෙයුමේ පිරිවැය එක් එක් ස්ලට් එකෙහි මොඩියුල එකතු කරනු ලැබේ.

වගුව ගබඩා කරන්නේ එක් එක් මූලද්‍රව්‍ය සඳහා යතුර සහ අගය පමණක් මිස යතුරේ හැෂ් එකක් නොවේ. වගුව 32-bit යතුරු පමණක් ගබඩා කර ඇති බැවින්, හෑෂ් ඉතා ඉක්මනින් ගණනය කරනු ලැබේ. ඉහත කේතය Murmur3 හැෂ් භාවිතා කරයි, එය මාරු කිරීම්, XOR සහ ගුණ කිරීම් කිහිපයක් පමණක් සිදු කරයි.

හැෂ් වගුව මතක අනුපිළිවෙලින් ස්වාධීන වන අගුලු දැමීමේ ආරක්ෂණ ක්‍රම භාවිතා කරයි. සමහර ලිවීම් මෙහෙයුම් එවැනි අනෙකුත් මෙහෙයුම්වල අනුපිළිවෙලට බාධා කළත්, හැෂ් වගුව තවමත් නිවැරදි තත්ත්වය පවත්වා ගනී. අපි මේ ගැන පහතින් කතා කරමු. මෙම තාක්ෂණය සමගාමීව නූල් දහස් ගණනක් ධාවනය වන වීඩියෝ කාඩ්පත් සමඟ විශිෂ්ට ලෙස ක්රියා කරයි.

හැෂ් වගුවේ ඇති යතුරු සහ අගයන් හිස් කිරීමට ආරම්භ කර ඇත.

64-bit යතුරු සහ අගයන් හැසිරවීමට කේතය වෙනස් කළ හැක. යතුරු සඳහා පරමාණුක කියවීම, ලිවීම සහ සංසන්දනය කිරීම සහ හුවමාරු කිරීමේ මෙහෙයුම් අවශ්‍ය වේ. සහ අගයන් සඳහා පරමාණුක කියවීම සහ ලිවීමේ මෙහෙයුම් අවශ්‍ය වේ. වාසනාවකට මෙන්, CUDA හි, 32- සහ 64-bit අගයන් සඳහා කියවීමේ-ලිවීමේ මෙහෙයුම් ඒවා ස්වභාවිකව පෙලගැසී ඇති තාක් පරමාණුක වේ (පහත බලන්න). මෙහි), සහ නවීන වීඩියෝ කාඩ්පත් 64-bit පරමාණුක සංසන්දනය සහ හුවමාරු මෙහෙයුම් සඳහා සහය දක්වයි. ඇත්ත වශයෙන්ම, බිටු 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(), යතුරක් මත ගැළපීමක් සොයාගත් විට, එය එහි අගය හිස් කරයි.

ඉහත සඳහන් කළ පරිදි, යතුරක් slot එකකට ලියා ඇති විට, එය තවදුරටත් චලනය නොවේ. වගුවෙන් මූලද්‍රව්‍යයක් මකා දැමූ විට පවා, යතුර එහි ස්ථානයේ පවතී, එහි අගය හිස් වේ. මෙයින් අදහස් කරන්නේ අපට slot අගය සඳහා පරමාණුක ලිවීමේ මෙහෙයුමක් භාවිතා කිරීමට අවශ්‍ය නොවන බවයි, මන්ද වත්මන් අගය හිස්ද නැද්ද යන්න ගැටළුවක් නොවේ - එය තවමත් හිස් වනු ඇත.

හැෂ් වගුවක ප්‍රමාණය වෙනස් කිරීම

ඔබට විශාල වගුවක් සාදා පැරණි වගුවේ ඇති හිස් නොවන මූලද්‍රව්‍ය එයට ඇතුළත් කිරීමෙන් ඔබට හැෂ් වගුවක ප්‍රමාණය වෙනස් කළ හැකිය. නියැදි කේතය සරලව තබා ගැනීමට මට අවශ්‍ය නිසා මම මෙම ක්‍රියාකාරීත්වය ක්‍රියාත්මක නොකළෙමි. එපමනක් නොව, 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.

දැන් අපි Functions කළ යුතුද යන්න ගැන කතා කරමු lookup() и delete() හැෂ් වගුවේ යුගල අරාවකට සරල හෝ වාෂ්පශීලී දර්ශකයක් භාවිතා කරන්න. CUDA ලේඛනගත කිරීම ප්‍රකාශ කරන්නේ:

සම්පාදකයා ගෝලීය හෝ හවුල් මතකයට කියවීම් සහ ලිවීම් ප්‍රශස්ත කිරීමට තෝරාගත හැක... මෙම ප්‍රශස්තකරණයන් මූල පදය භාවිතයෙන් අක්‍රිය කළ හැක. volatile: ... මෙම විචල්‍යය පිළිබඳ ඕනෑම සඳහනක් සැබෑ මතක කියවීමට හෝ ලිවීමට උපදෙස් බවට සම්පාදනය කෙරේ.

නිවැරදි සලකා බැලීම් අයදුම් කිරීම අවශ්ය නොවේ volatile. ක්‍රියාත්මක කිරීමේ නූල් පෙර කියවූ මෙහෙයුමකින් හැඹිලිගත අගයක් භාවිතා කරන්නේ නම්, එය තරමක් යල් පැන ගිය තොරතුරු භාවිතා කරයි. නමුත් තවමත්, මෙය කර්නල් ඇමතුමේ නිශ්චිත මොහොතක හැෂ් වගුවේ නිවැරදි තත්වයෙන් තොරතුරු වේ. ඔබට නවතම තොරතුරු භාවිතා කිරීමට අවශ්‍ය නම්, ඔබට දර්ශකය භාවිතා කළ හැකිය volatile, නමුත් එවිට කාර්ය සාධනය තරමක් අඩු වනු ඇත: මගේ පරීක්ෂණ වලට අනුව, මූලද්‍රව්‍ය මිලියන 32 ක් මකා දැමීමේදී, වේගය මකාදැමීම් මිලියන 500 / තත්පර සිට මිලියන 450 මකාදැමීම් / තත්පර දක්වා අඩු විය.

ඵලදායිතාව

මුලද්‍රව්‍ය මිලියන 64 ක් ඇතුළු කිරීම සහ ඒවායින් මිලියන 32 ක් මකා දැමීම සඳහා වන පරීක්ෂණයේදී තරඟය අතර std::unordered_map GPU සඳහා ප්‍රායෝගිකව හැෂ් වගුවක් නොමැත:

std::unordered_map මූලද්‍රව්‍ය ඇතුළු කිරීම සහ ඉවත් කිරීම සහ ඒවා නිදහස් කිරීම සඳහා ms 70 ක් වැය කරන ලදී unordered_map (මූලද්‍රව්‍ය මිලියන ගණනක් ඉවත් කිරීමට බොහෝ කාලයක් ගතවේ, මන්ද ඇතුළත unordered_map බහු මතක වෙන් කිරීම් සිදු කරනු ලැබේ). අවංකවම කියනවා නම්, std:unordered_map සම්පූර්ණයෙන්ම වෙනස් සීමා කිරීම්. එය ක්‍රියාත්මක කිරීමේ තනි CPU නූලකි, ඕනෑම ප්‍රමාණයක යතුරු අගයන් සඳහා සහය දක්වයි, ඉහළ උපයෝගිතා අනුපාතවලදී හොඳින් ක්‍රියා කරයි, සහ බහු මකාදැමීමෙන් පසු ස්ථාවර කාර්ය සාධනයක් පෙන්වයි.

GPU සහ අන්තර් වැඩසටහන් සන්නිවේදනය සඳහා හැෂ් වගුවේ කාලසීමාව 984 ms විය. මේසය මතකයේ තබා එය මකා දැමීම (එක් වරක් මතකය 1 GB වෙන් කිරීම, CUDA හි යම් කාලයක් ගත වේ), මූලද්‍රව්‍ය ඇතුළු කිරීම සහ මකා දැමීම සහ ඒවා මත නැවත නැවත කිරීම මෙයට ඇතුළත් වේ. වීඩියෝ කාඩ්පත් මතකයට සහ ඉන් ලැබෙන සියලුම පිටපත් ද සැලකිල්ලට ගනී.

හැෂ් වගුව සම්පූර්ණ කිරීමට ms 271ක් ගත විය. වීඩියෝ කාඩ්පත මඟින් මූලද්‍රව්‍ය ඇතුළු කිරීම සහ මකා දැමීම සඳහා ගත කරන කාලය මෙයට ඇතුළත් වන අතර, මතකයට පිටපත් කිරීමට සහ ලැබෙන වගුව හරහා පුනරාවර්තනය කිරීමට ගත කරන කාලය සැලකිල්ලට නොගනී. GPU වගුව දිගු කාලයක් ජීවත් වන්නේ නම්, හෝ හෑෂ් වගුව සම්පූර්ණයෙන්ම වීඩියෝ කාඩ්පතේ මතකයේ තිබේ නම් (උදාහරණයක් ලෙස, මධ්‍යම ප්‍රොසෙසරය නොව වෙනත් GPU කේතය භාවිතා කරන හැෂ් වගුවක් සෑදීමට), එවිට පරීක්ෂණ ප්රතිඵලය අදාළ වේ.

වීඩියෝ කාඩ්පතක් සඳහා හැෂ් වගුව ඉහළ ප්‍රතිදානය සහ ක්‍රියාකාරී සමාන්තරකරණය හේතුවෙන් ඉහළ කාර්ය සාධනයක් පෙන්නුම් කරයි.

අඩුපාඩු

හැෂ් වගු ගෘහ නිර්මාණ ශිල්පයට දැනගත යුතු ගැටළු කිහිපයක් තිබේ:

රේඛීය ගවේෂණය පොකුරු කිරීම මගින් බාධා ඇති කරයි, එමඟින් මේසයේ යතුරු පරිපූර්ණව වඩා අඩුවෙන් තැබීමට හේතු වේ.
කාර්යය භාවිතයෙන් යතුරු ඉවත් නොකෙරේ delete කාලයත් සමඟ ඔවුන් මේසය අවුල් කරයි.

එහි ප්‍රතිඵලයක් වශයෙන්, හෑෂ් වගුවක ක්‍රියාකාරීත්වය ක්‍රමයෙන් පිරිහීමට ලක් විය හැක, විශේෂයෙන්ම එය දිගු කාලයක් පවතින විට සහ බොහෝ ඇතුළු කිරීම් සහ මකාදැමීම් තිබේ නම්. මෙම අවාසි අවම කර ගැනීමට එක් ක්‍රමයක් නම් තරමක් අඩු උපයෝගිතා අනුපාතයක් සහිත නව වගුවකට නැවත සකස් කිරීම සහ නැවත සකස් කිරීමේදී ඉවත් කරන ලද යතුරු පෙරීමයි.

විස්තර කර ඇති ගැටළු නිදර්ශනය කිරීම සඳහා, මම ඉහත කේතය භාවිතා කර මූලද්‍රව්‍ය මිලියන 128ක් සහිත වගුවක් සාදා මිලියන 4ක මූලද්‍රව්‍ය හරහා ලූප් කරන්නෙමි. මෙන්න ප්‍රතිඵල වගුව, සෑම පේළියක්ම එක් හැෂ් වගුවකට නව මූලද්‍රව්‍ය මිලියන 124ක් ඇතුළු කිරීමට 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-bit යතුරු සහ අගයන් සම්බන්ධයෙන්, මෙය සාධාරණීකරණය කළ හැකිය. ඉහත උදාහරණයේ, මූලද්‍රව්‍ය මිලියන 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

GPU සඳහා සරල හැෂ් වගුව

අගුල් නොමැතිව හැෂ් මේසය

හැෂ් වගු තත්ත්වය

හැෂ් වගුවකට ඇතුල් කිරීම

හැෂ් මේස බැලීම

හැෂ් වගුවක මකා දැමීම

හැෂ් වගුවක ප්‍රමාණය වෙනස් කිරීම

තරඟකාරිත්වය

ඵලදායිතාව

අඩුපාඩු

ශබ්ද ගැඹුර මැනීම

නිගමනය

අදහස් එක් කරන්න

GPU සඳහා සරල හැෂ් වගුව

අගුල් නොමැතිව හැෂ් මේසය

හැෂ් වගු තත්ත්වය

හැෂ් වගුවකට ඇතුල් කිරීම

හැෂ් මේස බැලීම

හැෂ් වගුවක මකා දැමීම

හැෂ් වගුවක ප්‍රමාණය වෙනස් කිරීම

තරඟකාරිත්වය

ඵලදායිතාව

අඩුපාඩු

ශබ්ද ගැඹුර මැනීම

නිගමනය

අදහස් එක් කරන්න ඊ-මේල් මගින් පිලිතුරු දෙන්න අවලංගු

අදහස් එක් කරන්න