Ես այն տեղադրել եմ Github-ում նոր նախագիծ A Simple GPU Hash Table.

Դա 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 միջուկում, այլ հյուրընկալող կոդով:

Հոդվածում A Lock-Free Wait-Free Hash Table նկարագրում է, թե ինչպես փոփոխել կողպեքով պաշտպանված տվյալների կառուցվածքը:

Մրցունակություն

Վերոնշյալ գործառույթի կոդի հատվածներում 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 բոլորովին այլ սահմանափակումներ. Այն կատարման մեկ պրոցեսորի թեմա է, աջակցում է ցանկացած չափի առանցքային արժեքներին, լավ է աշխատում օգտագործման բարձր տեմպերով և ցույց է տալիս կայուն կատարում բազմակի ջնջումներից հետո:

GPU-ի և միջծրագրային հաղորդակցության համար հեշ աղյուսակի տևողությունը 984 ms էր: Սա ներառում է աղյուսակը հիշողության մեջ տեղադրելու և այն ջնջելու համար ծախսված ժամանակը (մեկ անգամ հատկացնելով 1 ԳԲ հիշողություն, որը որոշակի ժամանակ է պահանջում CUDA-ում), տարրերի տեղադրման և ջնջման և դրանց վրա կրկնվող ժամանակի վրա: Հաշվի են առնվում նաև վիդեո քարտի հիշողության բոլոր պատճենները:

Հեշ աղյուսակն ինքնին տևել է 271 ms՝ ավարտելու համար: Սա ներառում է վիդեո քարտի կողմից տարրեր տեղադրելու և ջնջելու ժամանակը, և հաշվի չի առնվում հիշողության մեջ պատճենելու և ստացված աղյուսակի վրայով կրկնվող ժամանակը: Եթե ​​GPU աղյուսակը երկար ժամանակ է ապրում, կամ եթե հեշ աղյուսակն ամբողջությամբ պարունակվում է վիդեո քարտի հիշողության մեջ (օրինակ՝ ստեղծելու հեշ աղյուսակ, որը կօգտագործվի այլ GPU կոդով և ոչ կենտրոնական պրոցեսորով), ապա թեստի արդյունքը տեղին է.

Վիդեոքարտի հեշ աղյուսակը ցույց է տալիս բարձր արդյունավետություն՝ շնորհիվ բարձր թողունակության և ակտիվ զուգահեռացման:

Սահմանափակումները

Հեշ սեղանի ճարտարապետությունը մի քանի խնդիրներ ունի, որոնց մասին պետք է տեղյակ լինել.

• Գծային զոնդավորումը խոչընդոտում է կլաստերավորումը, ինչը հանգեցնում է աղյուսակի ստեղների տեղադրմանը ավելի քիչ, քան կատարյալ:
• Ստեղները չեն հեռացվում ֆունկցիայի միջոցով delete ու ժամանակի ընթացքում սեղանը խառնում են:

Արդյունքում, հեշ աղյուսակի կատարումը կարող է աստիճանաբար վատթարանալ, հատկապես, եթե այն գոյություն ունի երկար ժամանակ և ունի բազմաթիվ ներդիրներ և ջնջումներ: Այս թերությունները մեղմելու եղանակներից մեկը նոր աղյուսակի մեջ կրկին հղկվելն է՝ օգտագործման բավականին ցածր արագությամբ և մաքրման ընթացքում հեռացված ստեղները զտել:

Նկարագրված խնդիրները լուսաբանելու համար ես կօգտագործեմ վերը նշված կոդը՝ 128 միլիոն տարրերով աղյուսակ ստեղծելու և 4 միլիոն տարրերի միջով անցնելու համար, մինչև որ լրացնեմ 124 միլիոն սլոտ (օգտագործման մակարդակը մոտ 0,96): Ահա արդյունքների աղյուսակը, յուրաքանչյուր տող CUDA միջուկի զանգ է՝ 4 միլիոն նոր տարրեր մեկ հեշ աղյուսակում տեղադրելու համար.

Օգտագործման տոկոսադրույքը
Տեղադրման տևողությունը 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 ՄԲ կորցրած հիշողություն։

Խնդրում ենք նկատի ունենալ, որ խոսքը վիդեո քարտի հիշողության կորստի մասին է, որն ավելի արժեքավոր ռեսուրս է, քան համակարգի հիշողությունը: Թեև ժամանակակից աշխատասեղանի գրաֆիկական քարտերի մեծ մասը, որոնք աջակցում են 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 կամ բացեք նոր թեմա պահոցներ.

Այս ծածկագիրը գրվել է հիանալի հոդվածների ոգեշնչմամբ.

• Աշխարհի ամենապարզ առանց կողպեքի հեշ աղյուսակը
• A Lock-Free Wait-Free Hash Table

Հետագայում ես կշարունակեմ գրել վիդեո քարտերի համար հեշ աղյուսակի ներդրման մասին և վերլուծել դրանց կատարումը: Իմ ծրագրերը ներառում են շղթայական կապ, Ռոբին Հուդի հաշինգ և կուկու հեշինգ՝ օգտագործելով ատոմային գործողությունները տվյալների կառուցվածքներում, որոնք հարմար են GPU-ին:

Source: www.habr.com