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

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

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

حالة جدول التجزئة

• المفتاح والمعنى فارغان. في هذه الحالة ، تتم تهيئة جدول التجزئة.
• تم تسجيل المفتاح ، ولكن القيمة ليست بعد. إذا كان هناك مؤشر ترابط آخر يقرأ البيانات في تلك اللحظة ، فإنه يُرجع قيمة فارغة. هذا أمر طبيعي ، سيحدث الشيء نفسه إذا نجحت سلسلة تنفيذ أخرى قبل ذلك بقليل ، ونحن نتحدث عن بنية بيانات تنافسية.
• يتم تسجيل كل من المفتاح والقيمة.
• القيمة متاحة لسلاسل التنفيذ الأخرى ، لكن المفتاح لم يتم بعد. يمكن أن يحدث هذا لأن نموذج البرمجة 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);
    }
}

تجزئة جدول البحث

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);
    }
}

تغيير حجم جدول التجزئة

القدرة التنافسية

// 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);
    }
}

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

أداء

سلبيات

• يتداخل التكتل مع الفحص الخطي ، لأن المفاتيح الموجودة في الجدول بعيدة عن المثالية.
• لا يتم حذف المفاتيح باستخدام الوظيفة deleteوبمرور الوقت يزدحم الجدول.

قياس عمق الاستشعار

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

استنتاج

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