நான் அதை Github இல் இடுகையிட்டேன் புதிய திட்டம் ஒரு எளிய GPU ஹாஷ் அட்டவணை.

இது ஒரு வினாடிக்கு நூற்றுக்கணக்கான மில்லியன் செருகல்களைச் செயலாக்கும் ஒரு எளிய GPU ஹாஷ் அட்டவணையாகும். எனது NVIDIA GTX 1060 மடிக்கணினியில், குறியீடு தோராயமாக உருவாக்கப்பட்ட 64 மில்லியன் முக்கிய மதிப்பு ஜோடிகளை சுமார் 210 ms இல் செருகுகிறது மற்றும் 32 மில்லியன் ஜோடிகளை சுமார் 64 ms இல் நீக்குகிறது.

அதாவது, மடிக்கணினியின் வேகம் தோராயமாக 300 மில்லியன் செருகல்கள்/வினாடி மற்றும் 500 மில்லியன் நீக்கல்கள்/வினாடி ஆகும்.

அட்டவணை CUDA இல் எழுதப்பட்டுள்ளது, இருப்பினும் அதே நுட்பத்தை HLSL அல்லது GLSL க்கும் பயன்படுத்தலாம். வீடியோ அட்டையில் உயர் செயல்திறனை உறுதிப்படுத்த செயல்படுத்த பல வரம்புகள் உள்ளன:

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() ஜோடிகளின் வரிசையை இணையாகச் செயல்படுத்தவும், ஒவ்வொரு ஜோடியும் தனித்தனி 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 ms ஆனது. வீடியோ கார்டு உறுப்புகளைச் செருகுவதற்கும் நீக்குவதற்கும் செலவழித்த நேரத்தை உள்ளடக்கியது, மேலும் நினைவகத்தில் நகலெடுக்கும் நேரத்தையும் அதன் விளைவாக வரும் அட்டவணையில் மீண்டும் மீண்டும் செய்வதையும் கணக்கில் எடுத்துக்கொள்ளாது. 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 மில்லியன் செருகல்களுடன் (பயன்பாட்டு காரணி 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 க்கு சமம்.

முடிவுக்கு

உங்களிடம் கேள்விகள் அல்லது கருத்துகள் இருந்தால், தயவுசெய்து எனக்கு மின்னஞ்சல் அனுப்பவும் ட்விட்டர் அல்லது ஒரு புதிய தலைப்பை திறக்கவும் களஞ்சியங்கள்.

இந்த குறியீடு சிறந்த கட்டுரைகளின் உத்வேகத்தின் கீழ் எழுதப்பட்டது:

எதிர்காலத்தில், வீடியோ கார்டுகளுக்கான ஹாஷ் அட்டவணை செயலாக்கங்களைப் பற்றி தொடர்ந்து எழுதுவேன் மற்றும் அவற்றின் செயல்திறனை பகுப்பாய்வு செய்வேன். எனது திட்டங்களில் செயினிங், ராபின் ஹூட் ஹேஷிங் மற்றும் குக்கூ ஹாஷிங் ஆகியவை ஜிபியு நட்பு தரவு கட்டமைப்புகளில் அணு செயல்பாடுகளைப் பயன்படுத்துகின்றன.

ஆதாரம்: www.habr.com

GPU க்கான எளிய ஹாஷ் அட்டவணை