مون ان کي Github تي پوسٽ ڪيو نئون پروجيڪٽ هڪ سادي GPU هش ٽيبل.

اهو هڪ سادي GPU هيش ٽيبل آهي جيڪو پروسيسنگ ڪرڻ جي قابل آهي سوين لکن داخلائن في سيڪنڊ. منهنجي NVIDIA GTX 1060 ليپ ٽاپ تي، ڪوڊ داخل ڪري ٿو 64 ملين بي ترتيب سان ٺاهيل ڪي-ويليو جوڙو اٽڪل 210 ms ۾ ۽ 32 ملين جوڙو ڪڍي ٿو اٽڪل 64 ms ۾.

اهو آهي، هڪ ليپ ٽاپ تي رفتار لڳ ڀڳ 300 ملين انسرٽس/سيڪنڊ ۽ 500 ملين ڊيليٽس/سيڪنڊ آهي.

ٽيبل CUDA ۾ لکيل آهي، جيتوڻيڪ ساڳي ٽيڪنڪ HLSL يا GLSL تي لاڳو ٿي سگهي ٿي. هڪ وڊيو ڪارڊ تي اعلي ڪارڪردگي کي يقيني بڻائڻ لاء عمل درآمد ۾ ڪيترائي حدون آهن:

صرف 32-bit چابيون ۽ ساڳيا قدر پروسيس ٿيل آهن.
هيش ٽيبل هڪ مقرر سائيز آهي.
۽ هي سائيز ٻن طاقتن جي برابر هجڻ گهرجي.

چابيون ۽ قدرن لاءِ، توھان کي ھڪ سادو ڊيليميٽر مارڪر رکڻو پوندو (مٿي ڏنل ڪوڊ ۾ اھو آھي 0xffffffff).

هيش ٽيبل بغير تالا

هيش ٽيبل استعمال ڪري ٿو اوپن ايڊريسنگ سان لڪير جي جاچ، اهو آهي، اهو صرف اهم-قدر جوڑوں جو هڪ صف آهي جيڪو ميموري ۾ ذخيرو ٿيل آهي ۽ اعلي ڪيش ڪارڪردگي آهي. ساڳيءَ ريت زنجير لاءِ به نه ٿو چئي سگهجي، جنهن ۾ ڳنڍيل لسٽ ۾ هڪ پوائنٽر ڳولڻ شامل آهي. هيش ٽيبل هڪ سادي صف آهي اسٽوريج عناصر KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

ٽيبل جي سائيز ٻن جي طاقت آھي، ھڪڙو پرائم نمبر نه آھي، ڇاڪاڻ ته ھڪڙو تيز ھدايت ڪافي آھي pow2 / AND ماسڪ لاڳو ڪرڻ لاء، پر ماڊيولس آپريٽر تمام گھڻو سست آھي. اهو لڪير جي جاچ جي صورت ۾ اهم آهي، ڇاڪاڻ ته هڪ لڪير ٽيبل جي نظر ۾ سلاٽ انڊيڪس هر سلاٽ ۾ لپي وڃي. ۽ نتيجي طور، آپريشن جي قيمت هر سلاٽ ۾ ماڊل شامل ڪئي وئي آهي.

ٽيبل صرف هر عنصر لاء اهم ۽ قيمت کي محفوظ ڪري ٿو، نه ڪي ڪي هيش. جيئن ته ٽيبل صرف 32-بٽ چيڪن کي محفوظ ڪري ٿو، هيش تمام جلدي حساب ڪيو ويندو آهي. مٿي ڏنل ڪوڊ استعمال ڪري ٿو Murmur3 هيش، جيڪو صرف چند شفٽ، XORs ۽ ضرب انجام ڏئي ٿو.

هيش ٽيبل استعمال ڪري ٿو تالا لڳائڻ واري حفاظتي ٽيڪنالاجي جيڪي ميموري آرڊر کان آزاد آهن. جيتوڻيڪ جيڪڏهن ڪجهه لکڻ جي عملن ٻين اهڙين عملن جي ترتيب کي خراب ڪن ٿا، هيش ٽيبل اڃا تائين صحيح حالت برقرار رکندو. اسان هيٺ ان بابت ڳالهائينداسين. ٽيڪنڪ ويڊيو ڪارڊن سان تمام سٺو ڪم ڪري ٿي جيڪي ھزارين موضوعن کي گڏ ڪن ٿا.

هيش ٽيبل ۾ ڪنجيون ۽ قدرن کي خالي ڪرڻ لاءِ شروعات ڪئي وئي آهي.

ڪوڊ تبديل ڪري سگھجي ٿو 64-bit ڪنجين ۽ قدرن کي پڻ سنڀالڻ لاءِ. ڪيز کي ايٽمي پڙهڻ، لکڻ، ۽ موازنہ-۽-سوپ آپريشن جي ضرورت آهي. ۽ قدرن کي ايٽمي پڙهڻ ۽ لکڻ جي عملن جي ضرورت آهي. خوشقسمتيءَ سان، CUDA ۾، 32- ۽ 64-bit قدرن لاءِ پڙھڻ-لکڻ جا عمل ايٽمي آھن جيستائين اھي قدرتي طور تي جڙيل آھن (ھيٺ ڏسو). هتي)، ۽ جديد وڊيو ڪارڊ 64-bit ايٽمي مقابلي ۽ مٽا سٽا واري عمل کي سپورٽ ڪن ٿا. يقينن، جڏهن 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);
        }
}

ٽيبل ۾ محفوظ ڪيل ڪيئي جي قيمت معلوم ڪرڻ لاءِ، اسان ان ڪيئي جي هيش سان شروع ٿيندڙ صف جي ذريعي ٻيهر ورجائيندا آهيون جنهن کي اسان ڳولي رهيا آهيون. هر سلاٽ ۾، اسان چيڪ ڪريون ٿا ته چيڪ اهو آهي جيڪو اسان ڳولي رهيا آهيون، ۽ جيڪڏهن ائين آهي، اسان ان جي قيمت واپس ڪنداسين. اسان اهو به چيڪ ڪريون ٿا ته ڪيچي خالي آهي، ۽ جيڪڏهن ائين آهي، ته اسان ڳولا کي ختم ڪريون ٿا.

جيڪڏهن اسان ڪيڏي نه ڳولي سگهون ٿا، ڪوڊ هڪ خالي قيمت واپس ڪري ٿو.

اهي سڀئي سرچ آپريشن هڪ ئي وقت اندر داخل ڪرڻ ۽ حذف ڪرڻ جي ذريعي ڪري سگهجن ٿا. جدول ۾ هر هڪ جوڙو هوندو چار رياستن مان هڪ هوندي جيڪا وهڪري لاءِ مٿي بيان ڪئي وئي آهي.

هڪ hash ٽيبل ۾ ختم ڪرڻ

چيڪ کي حذف ڪرڻ لاء ڪوڊ:

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. جيڪڏھن execution thread اڳئين پڙھندڙ آپريشن مان ڪيش ڪيل قدر استعمال ڪري ٿي، پوءِ اھو ٿوري پراڻي معلومات استعمال ڪندي. پر اڃا تائين، اها معلومات آهي هيش ٽيبل جي صحيح حالت مان ڪنيل ڪال جي هڪ خاص لمحي تي. جيڪڏهن توهان کي جديد معلومات استعمال ڪرڻ جي ضرورت آهي، توهان انڊيڪس استعمال ڪري سگهو ٿا 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 ايم ايس ورتو. ھن ۾ شامل آھي وقت گذريو وڊيو ڪارڊ پاران عناصر داخل ڪرڻ ۽ حذف ڪرڻ، ۽ ميموري ۾ ڪاپي ڪرڻ ۽ نتيجي واري ٽيبل تي ورجائڻ ۾ خرچ ٿيل وقت کي حساب ۾ نٿو رکي. جيڪڏهن GPU ٽيبل گهڻي وقت تائين رهي ٿو، يا جيڪڏهن هيش ٽيبل مڪمل طور تي وڊيو ڪارڊ جي ميموري ۾ موجود آهي (مثال طور، هڪ هيش ٽيبل ٺاهڻ لاءِ جيڪو استعمال ڪيو ويندو ٻيو GPU ڪوڊ ۽ نه مرڪزي پروسيسر)، پوءِ. امتحان جو نتيجو لاڳاپيل آهي.

ويڊيو ڪارڊ لاءِ هيش ٽيبل اعلي ڪارڪردگي ۽ فعال متوازي ڪرڻ جي ڪري اعليٰ ڪارڪردگي ڏيکاري ٿو.

shortcomings

هيش ٽيبل آرڪيٽيڪچر ۾ ڪجھ مسئلا آهن جن کان آگاهي ٿيڻ گهرجي:

لڪير جي جاچ کي ڪلسٽرنگ جي ذريعي روڪيو ويو آهي، جنهن جي ڪري ٽيبل ۾ چابيون مڪمل طور تي گهٽ رکيا وڃن ٿيون.
فنڪشن استعمال ڪندي چابيون ختم نه ڪيون ويون آهن delete ۽ وقت گذرڻ سان گڏ اهي ميز کي ڇڪيندا آهن.

نتيجي طور، هيش ٽيبل جي ڪارڪردگي تيزيء سان خراب ٿي سگهي ٿي، خاص طور تي جيڪڏهن اهو هڪ ڊگهي وقت تائين موجود آهي ۽ ڪيترن ئي داخل ۽ حذف ٿيل آهن. انهن نقصانن کي گهٽائڻ جو هڪ طريقو اهو آهي ته هڪ نئين ٽيبل ۾ ٻيهر هٽايو وڃي هڪ تمام گهٽ استعمال جي شرح سان ۽ ٻيهر هٽائڻ دوران هٽايل ڪنجين کي فلٽر ڪيو وڃي.

بيان ڪيل مسئلن کي بيان ڪرڻ لاءِ، مان 128 ملين عنصرن سان ٽيبل ٺاهڻ لاءِ مٿي ڏنل ڪوڊ استعمال ڪندس ۽ 4 ملين عناصر جي ذريعي لوپ ڪندس جيستائين مان 124 ملين سلاٽ (استعمال جي شرح اٽڪل 0,96). هتي نتيجو جدول آهي، هر قطار هڪ CUDA ڪرنل ڪال آهي 4 ملين نوان عنصر هڪ هيش ٽيبل ۾ داخل ڪرڻ لاءِ:

استعمال جي شرح
داخل ٿيڻ جي مدت 4 عناصر

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-bit چاٻين ۽ قدرن جي صورت ۾، اھو صحيح ٿي سگھي ٿو. جيڪڏهن مٿين مثال ۾، 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

GPU لاءِ سادي هيش ٽيبل