გამოვაქვეყნე 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 ცხრილი დიდხანს ცხოვრობს, ან თუ ჰეშის ცხრილი მთლიანად შეიცავს ვიდეო ბარათის მეხსიერებას (მაგალითად, ჰეშის ცხრილის შესაქმნელად, რომელსაც გამოიყენებს სხვა 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 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-ს.

დასკვნა

თუ თქვენ გაქვთ შეკითხვები ან კომენტარები, გთხოვთ მომწეროთ მისამართზე Twitter ან გახსენით ახალი თემა საცავები.

ეს კოდი დაიწერა შესანიშნავი სტატიების შთაგონებით:

• მსოფლიოში ყველაზე მარტივი დაბლოკვის გარეშე ჰეშის ცხრილი
• დაბლოკვის გარეშე ლოდინის გარეშე ჰეშ მაგიდა

სამომავლოდ გავაგრძელებ ვიდეოკარტების ჰეშის ცხრილის დანერგვის შესახებ წერას და მათი შესრულების ანალიზს. ჩემი გეგმები მოიცავს ჯაჭვობას, რობინ ჰუდის ჰეშინგს და გუგულის ჰეშინგს ატომური ოპერაციების გამოყენებით მონაცემთა სტრუქტურებში, რომლებიც მოსახერხებელია GPU.

წყარო: www.habr.com