لقد نشرت على جيثب مشروع جديد جدول تجزئة GPU بسيط.

هذا جدول تجزئة GPU بسيط قادر على معالجة مئات الملايين من الإدخالات في الثانية. على جهاز الكمبيوتر المحمول NVIDIA GTX 1060 الخاص بي ، يقوم الكود بإدخال 64 مليون زوج قيم مفتاح تم إنشاؤه عشوائيًا في حوالي 210 مللي ثانية ويزيل 32 مليون زوج في حوالي 64 مللي ثانية.

أي أن السرعة على جهاز كمبيوتر محمول تبلغ حوالي 300 مليون إدخال / ثانية و 500 مليون عملية حذف / ثانية.

الجدول مكتوب بلغة CUDA ، على الرغم من أنه يمكن تطبيق نفس التقنية على HLSL أو GLSL. يحتوي التطبيق على العديد من القيود التي تضمن الأداء العالي على بطاقة الفيديو:

• تتم معالجة مفاتيح 32 بت ونفس القيم فقط.
• جدول التجزئة له حجم ثابت.
• وهذا الحجم يجب أن يساوي اثنين أس.

بالنسبة للمفاتيح والقيم ، تحتاج إلى حجز علامة تحديد بسيطة (في الكود أعلاه ، هذا هو 0xffffffff).

تجزئة الجدول بدون أقفال

يستخدم جدول التجزئة عنونة مفتوحة مع السبر الخطي، مما يعني أنها مجرد مجموعة من أزواج المفاتيح والقيمة المخزنة في الذاكرة وتتمتع بأداء فائق لذاكرة التخزين المؤقت. لا يمكن قول الشيء نفسه عن التسلسل ، والذي يتضمن البحث عن مؤشر في قائمة مرتبطة. جدول التجزئة هو مصفوفة بسيطة تخزن العناصر KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

حجم الجدول هو قوة اثنين ، وليس عددًا أوليًا ، لأن تعليمة واحدة سريعة تكفي لتطبيق قناع pow2 / AND ، وعامل modulo أبطأ بكثير. هذا مهم في حالة الفحص الخطي ، لأنه في بحث جدول خطي ، يجب لف فهرس الفتحة في كل فتحة. ونتيجة لذلك ، تتم إضافة تكلفة عملية النموذج في كل فتحة.

يخزن الجدول فقط المفتاح والقيمة لكل عنصر ، وليس تجزئة المفتاح. نظرًا لأن الجدول يخزن مفاتيح 32 بت فقط ، يتم حساب التجزئة بسرعة كبيرة. يستخدم الكود أعلاه تجزئة Murmur3 ، والتي تقوم فقط ببعض التحولات ، و XORs ، والمضاعفات.

يستخدم جدول التجزئة تقنية حماية قفل مستقلة عن ترتيب الذاكرة. حتى إذا كانت بعض عمليات الكتابة تنتهك ترتيب العمليات المماثلة الأخرى ، فسيظل جدول التجزئة يحتفظ بالحالة الصحيحة. سنتحدث عن هذا أدناه. تعمل هذه التقنية بشكل رائع مع بطاقات الفيديو التي تحتوي على آلاف سلاسل الرسائل التي تعمل بشكل متزامن.

تتم تهيئة المفاتيح والقيم في جدول التجزئة إلى قيمة خالية.

يمكن تعديل الكود للتعامل مع مفاتيح وقيم 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 ، تكتمل دائمًا أي عملية باستدعاء kernel قبل استدعاء kernel التالي (على الأقل داخل نفس مؤشر الترابط. في سلاسل العمليات المختلفة ، يتم تنفيذ النواة بالتوازي). إذا ، في المثال أعلاه ، تم استدعاء نواة واحدة مع A/0 B/1 A/2 C/3والآخر به A/4ثم المفتاح A سوف تحصل على القيمة 4.

الآن دعنا نتحدث عما إذا كان ينبغي للوظائف lookup() и delete() استخدم مؤشر عادي (عادي) أو متغير (متغير) لمجموعة من الأزواج في جدول تجزئة. وثائق كودا ينص علي:

قد يقوم المترجم حسب تقديره بتحسين عمليات القراءة والكتابة في الذاكرة العالمية أو المشتركة ... يمكن تعطيل هذه التحسينات باستخدام الكلمة الأساسية volatile: ... يتم تجميع أي إشارة إلى هذا المتغير في ذاكرة حقيقية قراءة أو كتابة تعليمات.

اعتبارات الصواب لا تتطلب التطبيق volatile. إذا كان مؤشر ترابط التنفيذ يستخدم قيمة مخزنة مؤقتًا من عملية قراءة سابقة ، فسيستخدم معلومات قديمة قليلاً. ولكن لا تزال هذه معلومات من الحالة الصحيحة لجدول التجزئة عند نقطة معينة في استدعاء kernel. إذا كنت بحاجة إلى استخدام أحدث المعلومات ، يمكنك استخدام الفهرس volatile، ولكن بعد ذلك سينخفض ​​الأداء بشكل طفيف: وفقًا لاختباراتي ، عند حذف 32 مليون عنصر ، انخفضت السرعة من 500 مليون عملية حذف / ثانية إلى 450 مليون عملية حذف / ثانية.

أداء

في اختبار لإدخال 64 مليون عنصر وإزالة 32 مليون منهم ، المنافسة بين std::unordered_map وجدول تجزئة لوحدة معالجة الرسومات مفقود بالفعل:

std::unordered_map قضى 70 مللي ثانية في إدخال العناصر وحذفها ثم تحريرها unordered_map (التحرر من ملايين العناصر يستغرق الكثير من الوقت ، لأن الداخل unordered_map يتم إجراء عمليات تخصيص متعددة). لنكون صادقين ، في std:unordered_map قيود مختلفة جدًا. إنه خيط تنفيذي واحد لوحدة المعالجة المركزية ، ويدعم القيم الرئيسية من أي حجم ، ويعمل بشكل جيد بمعدلات استخدام عالية ، ويظهر أداءً ثابتًا بعد عمليات الحذف المتعددة.

كانت مدة جدول التجزئة لوحدة معالجة الرسومات (GPU) والتفاعل بين البرامج 984 مللي ثانية. يتضمن ذلك الوقت المستغرق في وضع الجدول في الذاكرة وحذفه (تخصيص لمرة واحدة لـ 1 جيجابايت من الذاكرة ، والذي يستغرق بعض الوقت في CUDA) ، وإدخال العناصر وحذفها ، والتكرار عليها. يتم أيضًا أخذ جميع النسخ من ذاكرة بطاقة الفيديو وإليها في الاعتبار.

استغرق جدول التجزئة نفسه 271 مللي ثانية. يتضمن ذلك الوقت الذي تقضيه بطاقة الرسومات في إدخال العناصر وحذفها ، ولا يأخذ في الاعتبار الوقت المستغرق في النسخ إلى الذاكرة والتكرار فوق الجدول الناتج. إذا كان جدول GPU يعمل لفترة طويلة ، أو إذا تم الاحتفاظ بجدول التجزئة بالكامل في ذاكرة بطاقة الفيديو (على سبيل المثال ، لإنشاء جدول تجزئة سيتم استخدامه بواسطة رمز GPU الآخر ، وليس وحدة المعالجة المركزية) ، ثم نتيجة الاختبار هي ذات الصلة.

يوضح جدول التجزئة الخاص ببطاقة الفيديو أداءً عاليًا بسبب النطاق الترددي الكبير والتوازي النشط.

القيود

تحتوي بنية جدول التجزئة على العديد من المشكلات التي يجب وضعها في الاعتبار:

• يتم إعاقة التحقيق الخطي عن طريق التجميع ، مما يؤدي إلى وضع المفاتيح في الجدول بعيدًا عن الوضع المثالي.
• لا تتم إزالة المفاتيح باستخدام الوظيفة delete وتكدس على الطاولة بمرور الوقت.

نتيجة لذلك ، يمكن أن يتدهور أداء جدول التجزئة تدريجيًا ، خاصةً إذا كان موجودًا لفترة طويلة وتحدث فيه العديد من عمليات الإدراج والحذف. تتمثل إحدى طرق التخفيف من أوجه القصور هذه في إعادة صياغتها في جدول جديد مع عامل استخدام منخفض إلى حد ما وتصفية المفاتيح المحذوفة عند إعادة الصياغة.

لتوضيح المشكلات الموضحة ، باستخدام الكود أعلاه لإنشاء جدول به 128 مليون عنصر ، سوف أقوم بالتكرار خلال 4 ملايين عنصر حتى أقوم بملء 124 مليونًا من الفتحات (تبلغ نسبة الاستخدام حوالي 0,96). فيما يلي جدول النتائج ، كل سطر هو استدعاء لنواة CUDA لإدخال 4 ملايين عنصر جديد في جدول تجزئة واحد:

معدل الاستخدام
أدخل مدة 4،194،304 عناصر

0,00
11,608448 مللي ثانية (361,314798 مليون مفتاح / ثانية)

0,03
11,751424 مللي ثانية (356,918799 مليون مفتاح / ثانية)

0,06
11,942592 مللي ثانية (351,205515 مليون مفتاح / ثانية)

0,09
12,081120 مللي ثانية (347,178429 مليون مفتاح / ثانية)

0,12
12,242560 مللي ثانية (342,600233 مليون مفتاح / ثانية)

0,16
12,396448 مللي ثانية (338,347235 مليون مفتاح / ثانية)

0,19
12,533024 مللي ثانية (334,660176 مليون مفتاح / ثانية)

0,22
12,703328 مللي ثانية (330,173626 مليون مفتاح / ثانية)

0,25
12,884512 مللي ثانية (325,530693 مليون مفتاح / ثانية)

0,28
13,033472 مللي ثانية (321,810182 مليون مفتاح / ثانية)

0,31
13,239296 مللي ثانية (316,807174 مليون مفتاح / ثانية)

0,34
13,392448 مللي ثانية (313,184256 مليون مفتاح / ثانية)

0,37
13,624000 مللي ثانية (307,861434 مليون مفتاح / ثانية)

0,41
13,875520 مللي ثانية (302,280855 مليون مفتاح / ثانية)

0,44
14,126528 مللي ثانية (296,909756 مليون مفتاح / ثانية)

0,47
14,399328 مللي ثانية (291,284699 مليون مفتاح / ثانية)

0,50
14,690304 مللي ثانية (285,515123 مليون مفتاح / ثانية)

0,53
15,039136 مللي ثانية (278,892623 مليون مفتاح / ثانية)

0,56
15,478656 مللي ثانية (270,973402 مليون مفتاح / ثانية)

0,59
15,985664 مللي ثانية (262,379092 مليون مفتاح / ثانية)

0,62
16,668673 مللي ثانية (251,627968 مليون مفتاح / ثانية)

0,66
17,587200 مللي ثانية (238,486174 مليون مفتاح / ثانية)

0,69
18,690048 مللي ثانية (224,413765 مليون مفتاح / ثانية)

0,72
20,278816 مللي ثانية (206,831789 مليون مفتاح / ثانية)

0,75
22,545408 مللي ثانية (186,038058 مليون مفتاح / ثانية)

0,78
26,053312 مللي ثانية (160,989275 مليون مفتاح / ثانية)

0,81
31,895008 مللي ثانية (131,503463 مليون مفتاح / ثانية)

0,84
42,103294 مللي ثانية (99,619378 مليون مفتاح / ثانية)

0,87
61,849056 مللي ثانية (67,815164 مليون مفتاح / ثانية)

0,90
105,695999 مللي ثانية (39,682713 مليون مفتاح / ثانية)

0,94
240,204636 مللي ثانية (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.

اختتام

إذا كان لديك أي أسئلة أو تعليقات يرجى مراسلتي على البريد الإلكتروني تويتر أو افتح موضوعًا جديدًا في مستودعات.

هذا الرمز مستوحى من المقالات الرائعة:

• أبسط جدول تجزئة بدون قفل في العالم
• طاولة تجزئة خالية من القفل خالية من الانتظار

في المستقبل ، سأستمر في الكتابة عن تطبيقات جدول التجزئة لبطاقات الفيديو وتحليل أدائها. أخطط للتسلسل ، وتجزئة Robin Hood ، وتجزئة الوقواق باستخدام العمليات الذرية على هياكل البيانات الصديقة لبطاقات الرسومات.

المصدر: www.habr.com