Το δημοσίευσα στο Github νέο έργο Ένας απλός πίνακας κατακερματισμού GPU.

Είναι ένας απλός πίνακας κατακερματισμού GPU που μπορεί να επεξεργαστεί εκατοντάδες εκατομμύρια ένθετα ανά δευτερόλεπτο. Στον φορητό υπολογιστή NVIDIA GTX 1060, ο κωδικός εισάγει 64 εκατομμύρια ζεύγη κλειδιών-τιμών που δημιουργούνται τυχαία σε περίπου 210 ms και αφαιρεί 32 εκατομμύρια ζεύγη σε περίπου 64 ms.

Δηλαδή, η ταχύτητα σε ένα φορητό υπολογιστή είναι περίπου 300 εκατομμύρια ένθετα/δευτερόλεπτο και 500 εκατομμύρια διαγραφές/δευτ.

Ο πίνακας είναι γραμμένος σε CUDA, αν και η ίδια τεχνική μπορεί να εφαρμοστεί σε HLSL ή GLSL. Η υλοποίηση έχει αρκετούς περιορισμούς για να εξασφαλίσει υψηλή απόδοση σε μια κάρτα βίντεο:

Επεξεργάζονται μόνο κλειδιά 32-bit και οι ίδιες τιμές.
Ο πίνακας κατακερματισμού έχει σταθερό μέγεθος.
Και αυτό το μέγεθος πρέπει να είναι ίσο με δύο στην ισχύ.

Για κλειδιά και τιμές, πρέπει να κρατήσετε έναν απλό δείκτη οριοθέτησης (στον παραπάνω κωδικό είναι 0xffffffff).

Τραπέζι Hash χωρίς κλειδαριές

Ο πίνακας κατακερματισμού χρησιμοποιεί ανοιχτή διεύθυνση με γραμμική ανίχνευση, δηλαδή, είναι απλώς ένας πίνακας ζευγών κλειδιών-τιμών που αποθηκεύεται στη μνήμη και έχει ανώτερη απόδοση κρυφής μνήμης. Δεν μπορούμε να πούμε το ίδιο για το chaining, το οποίο περιλαμβάνει την αναζήτηση ενός δείκτη σε μια συνδεδεμένη λίστα. Ένας πίνακας κατακερματισμού είναι ένας απλός πίνακας που αποθηκεύει στοιχεία KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Το μέγεθος του πίνακα είναι δύναμη του δύο, όχι πρώτος αριθμός, γιατί αρκεί μια γρήγορη εντολή για την εφαρμογή της μάσκας pow2/AND, αλλά ο τελεστής συντελεστή είναι πολύ πιο αργός. Αυτό είναι σημαντικό στην περίπτωση της γραμμικής ανίχνευσης, καθώς σε μια γραμμική αναζήτηση πίνακα, ο δείκτης υποδοχής πρέπει να είναι τυλιγμένος σε κάθε υποδοχή. Και ως αποτέλεσμα, το κόστος της λειτουργίας προστίθεται modulo σε κάθε υποδοχή.

Ο πίνακας αποθηκεύει μόνο το κλειδί και την τιμή για κάθε στοιχείο, όχι έναν κατακερματισμό του κλειδιού. Δεδομένου ότι ο πίνακας αποθηκεύει μόνο κλειδιά 32-bit, ο κατακερματισμός υπολογίζεται πολύ γρήγορα. Ο παραπάνω κώδικας χρησιμοποιεί τον κατακερματισμό Murmur3, ο οποίος εκτελεί μόνο μερικές μετατοπίσεις, XOR και πολλαπλασιασμούς.

Ο πίνακας κατακερματισμού χρησιμοποιεί τεχνικές προστασίας κλειδώματος που είναι ανεξάρτητες από τη σειρά μνήμης. Ακόμα κι αν ορισμένες λειτουργίες εγγραφής διαταράσσουν τη σειρά άλλων τέτοιων λειτουργιών, ο πίνακας κατακερματισμού θα εξακολουθεί να διατηρεί τη σωστή κατάσταση. Θα μιλήσουμε για αυτό παρακάτω. Η τεχνική λειτουργεί εξαιρετικά με κάρτες βίντεο που εκτελούν χιλιάδες νήματα ταυτόχρονα.

Τα κλειδιά και οι τιμές στον πίνακα κατακερματισμού αρχικοποιούνται για άδεια.

Ο κώδικας μπορεί να τροποποιηθεί για να χειρίζεται επίσης κλειδιά και τιμές 64-bit. Τα κλειδιά απαιτούν ατομικές λειτουργίες ανάγνωσης, εγγραφής και σύγκρισης και ανταλλαγής. Και οι τιμές απαιτούν ατομικές λειτουργίες ανάγνωσης και εγγραφής. Ευτυχώς, στο CUDA, οι λειτουργίες ανάγνωσης-εγγραφής για τιμές 32- και 64-bit είναι ατομικές, εφόσον είναι φυσικά ευθυγραμμισμένες (δείτε παρακάτω). εδώ), και οι σύγχρονες κάρτες γραφικών υποστηρίζουν λειτουργίες ατομικής σύγκρισης και ανταλλαγής 64-bit. Φυσικά, όταν μετακινηθείτε στα 64 bit, η απόδοση θα μειωθεί ελαφρώς.

Κατάσταση πίνακα κατακερματισμού

Κάθε ζεύγος κλειδιού-τιμής σε έναν πίνακα κατακερματισμού μπορεί να έχει μία από τις τέσσερις καταστάσεις:

Το κλειδί και η τιμή είναι κενά. Σε αυτήν την κατάσταση, ο πίνακας κατακερματισμού αρχικοποιείται.
Το κλειδί έχει καταγραφεί, αλλά η τιμή δεν έχει γραφτεί ακόμα. Εάν ένα άλλο νήμα διαβάζει δεδομένα αυτήν τη στιγμή, επιστρέφει κενό. Αυτό είναι φυσιολογικό, το ίδιο θα είχε συμβεί εάν ένα άλλο νήμα εκτέλεσης είχε λειτουργήσει λίγο νωρίτερα, και μιλάμε για μια παράλληλη δομή δεδομένων.
Καταγράφονται τόσο το κλειδί όσο και η τιμή.
Η τιμή είναι διαθέσιμη σε άλλα νήματα εκτέλεσης, αλλά το κλειδί δεν είναι ακόμα. Αυτό μπορεί να συμβεί επειδή το μοντέλο προγραμματισμού 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 και με τον καιρό μπουκώνουν το τραπέζι.

Ως αποτέλεσμα, η απόδοση ενός πίνακα κατακερματισμού μπορεί σταδιακά να υποβαθμιστεί, ειδικά εάν υπάρχει για μεγάλο χρονικό διάστημα και έχει πολλά ένθετα και διαγραφές. Ένας τρόπος για να μετριαστούν αυτά τα μειονεκτήματα είναι να κάνετε rehash σε έναν νέο πίνακα με αρκετά χαμηλό ποσοστό χρήσης και να φιλτράρετε τα κλειδιά που αφαιρέθηκαν κατά τη διάρκεια της ανανέωσης.

Για να επεξηγήσω τα ζητήματα που περιγράφονται, θα χρησιμοποιήσω τον παραπάνω κώδικα για να δημιουργήσω έναν πίνακα με 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-bit, αυτό μπορεί να δικαιολογηθεί. Εάν στο παραπάνω παράδειγμα, σε έναν πίνακα με 128 εκατομμύρια στοιχεία, διατηρήσουμε τον συντελεστή χρήσης 0,25, τότε δεν μπορούμε να τοποθετήσουμε περισσότερα από 32 εκατομμύρια στοιχεία σε αυτόν και οι υπόλοιπες 96 εκατομμύρια υποδοχές θα χαθούν - 8 byte για κάθε ζεύγος , 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.

Συμπέρασμα

Εάν έχετε ερωτήσεις ή σχόλια, στείλτε μου email στο Twitter ή ανοίξτε ένα νέο θέμα αποθετήρια.

Αυτός ο κώδικας γράφτηκε με έμπνευση από εξαιρετικά άρθρα:

Στο μέλλον, θα συνεχίσω να γράφω για εφαρμογές κατακερματισμού πίνακα για κάρτες βίντεο και να αναλύω την απόδοσή τους. Τα σχέδιά μου περιλαμβάνουν το chaining, το Robin Hood hash και το cuckoo hashing με χρήση ατομικών λειτουργιών σε δομές δεδομένων που είναι φιλικές προς τις GPU.

Πηγή: www.habr.com

Απλός πίνακας κατακερματισμού για GPU