ຂ້ອຍໄດ້ໂພດມັນຢູ່ໃນ Github ໂຄງການໃຫມ່ A Simple GPU Hash Table.

ມັນເປັນຕາຕະລາງ GPU hash ງ່າຍດາຍທີ່ສາມາດປະມວນຜົນຫຼາຍຮ້ອຍລ້ານຂອງ inserts ຕໍ່ວິນາທີ. ໃນຄອມພິວເຕີໂນດບຸກ NVIDIA GTX 1060 ຂອງຂ້ອຍ, ລະຫັດໃສ່ 64 ລ້ານຄູ່ຄີ-ມູນຄ່າທີ່ສ້າງຂຶ້ນແບບສຸ່ມໃນປະມານ 210 ms ແລະເອົາ 32 ລ້ານຄູ່ໃນປະມານ 64 ms.

ນັ້ນແມ່ນ, ຄວາມໄວໃນຄອມພິວເຕີໂນດບຸກແມ່ນປະມານ 300 ລ້ານ inserts/sec ແລະ 500 ລ້ານ deletes/sec.

ຕາຕະລາງແມ່ນຂຽນໃນ CUDA, ເຖິງແມ່ນວ່າເຕັກນິກດຽວກັນສາມາດນໍາໃຊ້ກັບ HLSL ຫຼື GLSL. ການຈັດຕັ້ງປະຕິບັດມີຂໍ້ຈຳກັດຫຼາຍຢ່າງເພື່ອຮັບປະກັນປະສິດທິພາບສູງໃນບັດວີດີໂອ:

ພຽງແຕ່ 32-bit keys ແລະຄ່າດຽວກັນໄດ້ຖືກປະມວນຜົນ.
ຕາຕະລາງ hash ມີຂະຫນາດຄົງທີ່.
ແລະຂະຫນາດນີ້ຕ້ອງເທົ່າກັບສອງກັບພະລັງງານ.

ສໍາລັບລະຫັດແລະຄ່າ, ທ່ານຈໍາເປັນຕ້ອງສະຫງວນຕົວກໍານົດຕົວຂັ້ນງ່າຍດາຍ (ໃນລະຫັດຂ້າງເທິງນີ້ແມ່ນ 0xffffffff).

ຕາຕະລາງ hash ໂດຍບໍ່ມີການ locks

ຕາຕະລາງ hash ໃຊ້ການເປີດທີ່ຢູ່ກັບ ການສອບເສັງເສັ້ນຊື່, ນັ້ນແມ່ນ, ມັນເປັນພຽງແຕ່ array ຂອງຄູ່ key-value ທີ່ຖືກເກັບໄວ້ໃນຫນ່ວຍຄວາມຈໍາແລະມີປະສິດທິພາບ cache ດີກວ່າ. ສິ່ງດຽວກັນບໍ່ສາມາດເວົ້າໄດ້ສໍາລັບລະບົບຕ່ອງໂສ້, ເຊິ່ງກ່ຽວຂ້ອງກັບການຊອກຫາຕົວຊີ້ໃນບັນຊີລາຍຊື່ທີ່ເຊື່ອມໂຍງ. ຕາຕະລາງ hash ເປັນ array ງ່າຍດາຍເກັບຮັກສາອົງປະກອບ KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

ຂະຫນາດຂອງຕາຕະລາງແມ່ນພະລັງງານຂອງສອງ, ບໍ່ແມ່ນຕົວເລກຕົ້ນຕໍ, ເພາະວ່າຄໍາແນະນໍາທີ່ໄວອັນຫນຶ່ງແມ່ນພຽງພໍທີ່ຈະນໍາໃຊ້ຫນ້າກາກ pow2 / AND, ແຕ່ຕົວປະຕິບັດການ modulus ແມ່ນຊ້າກວ່າຫຼາຍ. ນີ້ແມ່ນສິ່ງສໍາຄັນໃນກໍລະນີຂອງການສືບສວນເສັ້ນ, ເນື່ອງຈາກວ່າໃນຕາຕະລາງເສັ້ນຊື່ຂອງຕາຕະລາງການເບິ່ງດັດຊະນີສະລັອດຕິງຕ້ອງຖືກຫໍ່ຢູ່ໃນແຕ່ລະຊ່ອງ. ແລະດັ່ງນັ້ນ, ຄ່າໃຊ້ຈ່າຍຂອງການດໍາເນີນງານແມ່ນເພີ່ມ modulo ໃນແຕ່ລະຊ່ອງສຽບ.

ຕາຕະລາງພຽງແຕ່ເກັບຮັກສາລະຫັດແລະຄ່າສໍາລັບແຕ່ລະອົງປະກອບ, ບໍ່ແມ່ນ hash ຂອງກະແຈ. ເນື່ອງຈາກຕາຕະລາງພຽງແຕ່ເກັບຮັກສາກະແຈ 32-bit, hash ໄດ້ຖືກຄິດໄລ່ຢ່າງໄວວາ. ລະຫັດຂ້າງເທິງໃຊ້ Murmur3 hash, ເຊິ່ງປະຕິບັດພຽງແຕ່ສອງສາມ shifts, XORs ແລະການຄູນ.

ຕາຕະລາງ hash ໃຊ້ເຕັກນິກການປ້ອງກັນການລັອກທີ່ບໍ່ຂຶ້ນກັບຄໍາສັ່ງຂອງຫນ່ວຍຄວາມຈໍາ. ເຖິງແມ່ນວ່າບາງການດໍາເນີນງານການຂຽນຂັດຂວາງຄໍາສັ່ງຂອງການດໍາເນີນງານດັ່ງກ່າວ, ຕາຕະລາງ hash ຈະຍັງຄົງຮັກສາສະຖານະທີ່ຖືກຕ້ອງ. ພວກເຮົາຈະເວົ້າກ່ຽວກັບເລື່ອງນີ້ຂ້າງລຸ່ມນີ້. ເຕັກນິກດັ່ງກ່າວເຮັດວຽກໄດ້ດີກັບບັດວີດີໂອທີ່ແລ່ນຫຼາຍພັນກະທູ້ພ້ອມກັນ.

ປຸ່ມແລະຄ່າໃນຕາຕະລາງ hash ໄດ້ຖືກເລີ່ມຕົ້ນໃຫ້ຫວ່າງເປົ່າ.

ລະຫັດສາມາດຖືກດັດແປງເພື່ອຈັດການກັບຄີ 64-bit ແລະຄ່າເຊັ່ນດຽວກັນ. ກະແຈຕ້ອງການການອ່ານ, ຂຽນ, ແລະການປຽບທຽບປະລໍາມະນູ. ແລະຄ່າຕ່າງໆຮຽກຮ້ອງໃຫ້ມີການປະຕິບັດການອ່ານແລະຂຽນປະລໍາມະນູ. ໂຊກດີ, ໃນ CUDA, ການປະຕິບັດການອ່ານ - ຂຽນສໍາລັບຄ່າ 32- ແລະ 64-bit ແມ່ນປະລໍາມະນູຕາບໃດທີ່ພວກມັນສອດຄ່ອງຕາມທໍາມະຊາດ (ເບິ່ງຂ້າງລຸ່ມນີ້). ທີ່ນີ້), ແລະບັດວີດີໂອທີ່ທັນສະໄຫມສະຫນັບສະຫນູນການປະຕິບັດການປຽບທຽບແລະແລກປ່ຽນປະລໍາມະນູ 64-bit. ແນ່ນອນ, ເມື່ອຍ້າຍໄປ 64 bits, ການປະຕິບັດຈະຫຼຸດລົງເລັກນ້ອຍ.

ສະຖານະຕາຕະລາງ Hash

ແຕ່ລະຄູ່ຄີ-ຄ່າໃນຕາຕະລາງ hash ສາມາດມີຫນຶ່ງໃນສີ່ລັດ:

ລະຫັດ ແລະຄ່າຫວ່າງເປົ່າ. ໃນສະຖານະນີ້, ຕາຕະລາງ hash ແມ່ນເລີ່ມຕົ້ນ.
ຄີໄດ້ຖືກຂຽນລົງ, ແຕ່ມູນຄ່າຍັງບໍ່ທັນໄດ້ຂຽນ. ຖ້າກະທູ້ອື່ນກຳລັງອ່ານຂໍ້ມູນຢູ່, ມັນຈະກັບຄືນມາຫວ່າງເປົ່າ. ນີ້ແມ່ນເລື່ອງປົກກະຕິ, ສິ່ງດຽວກັນຈະເກີດຂຶ້ນຖ້າຫາກວ່າກະທູ້ອື່ນຂອງການປະຕິບັດໄດ້ເຮັດວຽກເລັກນ້ອຍກ່ອນຫນ້ານີ້, ແລະພວກເຮົາເວົ້າກ່ຽວກັບໂຄງສ້າງຂໍ້ມູນພ້ອມກັນ.
ທັງຄີແລະມູນຄ່າຖືກບັນທຶກ.
ມູນຄ່າແມ່ນມີຢູ່ໃນກະທູ້ອື່ນຂອງການປະຕິບັດ, ແຕ່ຄີແມ່ນຍັງບໍ່ທັນ. ນີ້ສາມາດເກີດຂຶ້ນໄດ້ເພາະວ່າຮູບແບບການຂຽນໂປລແກລມ CUDA ມີຮູບແບບຫນ່ວຍຄວາມຈໍາທີ່ມີຄໍາສັ່ງວ່າງ. ນີ້ແມ່ນເລື່ອງປົກກະຕິ; ໃນກໍລະນີໃດກໍ່ຕາມ, ກະແຈຍັງຫວ່າງເປົ່າ, ເຖິງແມ່ນວ່າຄ່າຈະບໍ່ເປັນດັ່ງນັ້ນ.

ຄວາມແຕກຕ່າງທີ່ ສຳ ຄັນແມ່ນວ່າເມື່ອລະຫັດຖືກຂຽນໃສ່ສະລັອດຕິງ, ມັນບໍ່ໄດ້ຍ້າຍອີກຕໍ່ໄປ - ເຖິງແມ່ນວ່າລະຫັດຈະຖືກລຶບ, ພວກເຮົາຈະເວົ້າກ່ຽວກັບເລື່ອງນີ້ຂ້າງລຸ່ມນີ້.

ລະຫັດຕາຕະລາງ hash ເຖິງແມ່ນເຮັດວຽກກັບຕົວແບບຫນ່ວຍຄວາມຈໍາທີ່ມີຄໍາສັ່ງວ່າງທີ່ລໍາດັບທີ່ຫນ່ວຍຄວາມຈໍາຖືກອ່ານແລະຂຽນແມ່ນບໍ່ຮູ້. ເມື່ອພວກເຮົາເບິ່ງການແຊກ, ຊອກຫາ, ແລະການລຶບໃນຕາຕະລາງ hash, ຈື່ໄວ້ວ່າແຕ່ລະຄູ່ຄີ-ຄ່າແມ່ນຢູ່ໃນຫນຶ່ງໃນສີ່ລັດທີ່ອະທິບາຍຂ້າງເທິງ.

ການໃສ່ເຂົ້າໄປໃນຕາຕະລາງ hash

ຟັງຊັນ CUDA ທີ່ແຊກຄູ່ຄີ-ຄ່າເຂົ້າໄປໃນຕາຕະລາງ hash ເບິ່ງຄືວ່າ:

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

ເພື່ອໃສ່ກະແຈ, ລະຫັດຈະເຮັດຊ້ຳຜ່ານອາເຣຕາຕະລາງ hash ໂດຍເລີ່ມຕົ້ນດ້ວຍ hash ຂອງກະແຈທີ່ໃສ່. ແຕ່ລະຊ່ອງໃນ array ປະຕິບັດການປຽບທຽບແລະສະຫຼັບປະລໍາມະນູທີ່ປຽບທຽບກະແຈໃນຊ່ອງນັ້ນຫວ່າງເປົ່າ. ຖ້າກວດພົບບໍ່ກົງກັນ, ກະແຈໃນສະລັອດຕິງຖືກປັບປຸງດ້ວຍກະແຈທີ່ໃສ່, ແລະຫຼັງຈາກນັ້ນກະແຈສະລັອດຕິງເດີມຈະຖືກສົ່ງຄືນ. ຖ້າກະແຈຕົ້ນສະບັບນີ້ຫວ່າງເປົ່າ ຫຼືກົງກັບກະແຈທີ່ໃສ່ແລ້ວ, ລະຫັດໄດ້ພົບຊ່ອງສຽບທີ່ເໝາະສົມສຳລັບການໃສ່ ແລະໃສ່ຄ່າທີ່ໃສ່ໃສ່ໃນຊ່ອງໃສ່.

ຖ້າຢູ່ໃນຫນຶ່ງ kernel ໂທ gpu_hashtable_insert() ມີຫຼາຍອົງປະກອບທີ່ມີລະຫັດດຽວກັນ, ຫຼັງຈາກນັ້ນຄ່າໃດໆຂອງພວກມັນສາມາດຖືກຂຽນໃສ່ຊ່ອງສຽບ. ອັນນີ້ຖືວ່າເປັນເລື່ອງປົກກະຕິ: ຫນຶ່ງໃນຄໍາສໍາຄັນທີ່ຂຽນໃນລະຫວ່າງການໂທຈະປະສົບຜົນສໍາເລັດ, ແຕ່ວ່າທັງຫມົດນີ້ເກີດຂື້ນໃນຂະຫນານພາຍໃນຫຼາຍໆຫົວຂໍ້ຂອງການປະຕິບັດ, ພວກເຮົາບໍ່ສາມາດຄາດເດົາໄດ້ວ່າການຂຽນຫນ່ວຍຄວາມຈໍາໃດຈະເປັນອັນສຸດທ້າຍ.

ການຊອກຫາຕາຕະລາງ Hash

ລະຫັດສໍາລັບການຊອກຫາກະແຈ:

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

ເພື່ອຊອກຫາຄ່າຂອງກະແຈທີ່ເກັບໄວ້ໃນຕາຕະລາງ, ພວກເຮົາ iterate ຜ່ານ array ເລີ່ມຕົ້ນດ້ວຍ hash ຂອງ key ທີ່ພວກເຮົາກໍາລັງຊອກຫາ. ໃນແຕ່ລະຊ່ອງ, ພວກເຮົາກວດເບິ່ງວ່າກະແຈແມ່ນອັນທີ່ພວກເຮົາກໍາລັງຊອກຫາ, ແລະຖ້າເປັນດັ່ງນັ້ນ, ພວກເຮົາສົ່ງຄືນມູນຄ່າຂອງມັນ. ພວກເຮົາຍັງກວດເບິ່ງວ່າກະແຈຫວ່າງເປົ່າ, ແລະຖ້າເປັນດັ່ງນັ້ນ, ພວກເຮົາຍົກເລີກການຄົ້ນຫາ.

ຖ້າພວກເຮົາບໍ່ສາມາດຊອກຫາລະຫັດໄດ້, ລະຫັດຈະສົ່ງຄືນຄ່າຫວ່າງເປົ່າ.

ທັງຫມົດຂອງການຄົ້ນຫາເຫຼົ່ານີ້ສາມາດດໍາເນີນການພ້ອມກັນໂດຍຜ່ານການແຊກແລະການລຶບ. ແຕ່ລະຄູ່ໃນຕາຕະລາງຈະມີຫນຶ່ງໃນສີ່ລັດທີ່ອະທິບາຍຂ້າງເທິງສໍາລັບການໄຫຼ.

ການລຶບໃນຕາຕະລາງ 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(), ເວັ້ນເສຍແຕ່ວ່າເມື່ອພົບຄໍາທີ່ກົງກັນ, ມັນເຮັດໃຫ້ມູນຄ່າຂອງມັນຫວ່າງເປົ່າ.

ດັ່ງທີ່ໄດ້ກ່າວມາຂ້າງເທິງ, ເມື່ອລະຫັດຖືກຂຽນໃສ່ຊ່ອງ, ມັນຈະບໍ່ຖືກຍ້າຍອີກຕໍ່ໄປ. ເຖິງແມ່ນວ່າໃນເວລາທີ່ອົງປະກອບຖືກລົບອອກຈາກຕາຕະລາງ, ກຸນແຈຍັງຄົງຢູ່ໃນສະຖານທີ່, ມູນຄ່າຂອງມັນພຽງແຕ່ກາຍເປັນຫວ່າງເປົ່າ. ນີ້ຫມາຍຄວາມວ່າພວກເຮົາບໍ່ຈໍາເປັນຕ້ອງໃຊ້ການປະຕິບັດການຂຽນປະລໍາມະນູສໍາລັບຄ່າຊ່ອງ, ເພາະວ່າມັນບໍ່ສໍາຄັນວ່າຄ່າປະຈຸບັນແມ່ນຫວ່າງຫຼືບໍ່ - ມັນຈະກາຍເປັນຫວ່າງເປົ່າ.

ປັບຂະໜາດຕາຕະລາງ hash

ທ່ານສາມາດປ່ຽນຂະຫນາດຂອງຕາຕະລາງ hash ໂດຍການສ້າງຕາຕະລາງຂະຫນາດໃຫຍ່ແລະໃສ່ອົງປະກອບທີ່ບໍ່ຫວ່າງເປົ່າຈາກຕາຕະລາງເກົ່າເຂົ້າໄປໃນມັນ. ຂ້າພະເຈົ້າບໍ່ໄດ້ປະຕິບັດຫນ້າທີ່ນີ້ເພາະວ່າຂ້ອຍຕ້ອງການຮັກສາລະຫັດຕົວຢ່າງງ່າຍດາຍ. ຍິ່ງໄປກວ່ານັ້ນ, ໃນໂຄງການ CUDA, ການຈັດສັນຫນ່ວຍຄວາມຈໍາມັກຈະເຮັດຢູ່ໃນລະຫັດເຈົ້າພາບແທນທີ່ຈະຢູ່ໃນ CUDA kernel.

ບົດຂຽນ ຕາຕະລາງ Hash ທີ່ບໍ່ມີການລັອກ ອະທິບາຍວິທີການດັດແປງໂຄງສ້າງຂໍ້ມູນທີ່ມີການປ້ອງກັນການລັອກດັ່ງກ່າວ.

ຄວາມສາມາດໃນການແຂ່ງຂັນ

ໃນ snippet ລະຫັດຟັງຊັນຂ້າງເທິງ gpu_hashtable_insert(), _lookup() и _delete() ປະມວນຜົນຄູ່ຄີ-ຄ່າເທື່ອລະອັນ. ແລະຕ່ໍາກວ່າ gpu_hashtable_insert(), _lookup() и _delete() ປະມວນຜົນ array ຂອງຄູ່ຂະຫນານ, ແຕ່ລະຄູ່ໃນກະທູ້ປະຕິບັດ 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);
    }
}

ຕາຕະລາງ hash ທົນທານຕໍ່ lock ສະຫນັບສະຫນູນການແຊກ, ຊອກຫາ, ແລະການລຶບພ້ອມກັນ. ເນື່ອງຈາກວ່າຄູ່ຄີ-ຄ່າສະເຫມີຢູ່ໃນຫນຶ່ງໃນສີ່ລັດແລະກະແຈບໍ່ຍ້າຍອອກ, ຕາຕະລາງຮັບປະກັນຄວາມຖືກຕ້ອງເຖິງແມ່ນວ່າການດໍາເນີນການປະເພດຕ່າງໆຖືກນໍາໃຊ້ພ້ອມກັນ.

ແນວໃດກໍ່ຕາມ, ຖ້າພວກເຮົາປະມວນຜົນຊຸດການແຊກແລະການລຶບແບບຂະຫນານ, ແລະຖ້າ input array ຂອງຄູ່ມີລະຫັດທີ່ຊ້ໍາກັນ, ຫຼັງຈາກນັ້ນ, ພວກເຮົາຈະບໍ່ສາມາດຄາດເດົາໄດ້ວ່າຄູ່ໃດຈະ "ຊະນະ" - ຈະຖືກຂຽນໄວ້ໃນຕາຕະລາງ hash. ສົມມຸດວ່າພວກເຮົາເອີ້ນວ່າລະຫັດແຊກທີ່ມີ array input ຂອງຄູ່ A/0 B/1 A/2 C/3 A/4. ເມື່ອລະຫັດສໍາເລັດ, ຄູ່ B/1 и C/3 ມີການຮັບປະກັນວ່າມີຢູ່ໃນຕາຕະລາງ, ແຕ່ໃນເວລາດຽວກັນຄູ່ໃດຈະປາກົດຢູ່ໃນມັນ A/0, A/2 ຫຼື A/4. ນີ້ອາດຈະເປັນຫຼືອາດຈະບໍ່ເປັນບັນຫາ - ມັນທັງຫມົດແມ່ນຂຶ້ນກັບຄໍາຮ້ອງສະຫມັກ. ເຈົ້າອາດຈະຮູ້ລ່ວງໜ້າວ່າບໍ່ມີກະແຈທີ່ຊໍ້າກັນຢູ່ໃນ array ການປ້ອນຂໍ້ມູນ, ຫຼືເຈົ້າອາດຈະບໍ່ສົນໃຈວ່າຄ່າໃດຖືກຂຽນຫຼ້າສຸດ.

ຖ້ານີ້ແມ່ນບັນຫາສໍາລັບທ່ານ, ຫຼັງຈາກນັ້ນທ່ານຈໍາເປັນຕ້ອງແຍກຄູ່ທີ່ຊ້ໍາກັນເຂົ້າໄປໃນການໂທລະບົບ CUDA ທີ່ແຕກຕ່າງກັນ. ໃນ CUDA, ການດໍາເນີນການໃດໆທີ່ເອີ້ນ kernel ສະເຫມີຈະສໍາເລັດກ່ອນທີ່ຈະເອີ້ນ kernel ຕໍ່ໄປ (ຢ່າງຫນ້ອຍພາຍໃນຫນຶ່ງ thread. ໃນ thread ທີ່ແຕກຕ່າງກັນ, kernels ຖືກປະຕິບັດໃນຂະຫນານ). ໃນຕົວຢ່າງຂ້າງເທິງ, ຖ້າທ່ານໂທຫາຫນຶ່ງ kernel ກັບ A/0 B/1 A/2 C/3, ແລະອື່ນໆທີ່ມີ A/4, ຫຼັງຈາກນັ້ນທີ່ສໍາຄັນ A ຈະໄດ້ຮັບມູນຄ່າ 4.

ຕອນນີ້ໃຫ້ເວົ້າກ່ຽວກັບຫນ້າທີ່ຄວນ lookup() и delete() ໃຊ້ຕົວຊີ້ແບບທຳມະດາ ຫຼືຕົວຊີ້ການຜັນຜວນໄປຫາ array ຂອງຄູ່ໃນຕາຕະລາງ hash. ເອກະສານ CUDA ລະບຸວ່າ:

ຜູ້ສັງລວມອາດຈະເລືອກທີ່ຈະເພີ່ມປະສິດທິພາບການອ່ານແລະການຂຽນໃສ່ຫນ່ວຍຄວາມຈໍາທົ່ວໂລກຫຼືແບ່ງປັນ ... ການເພີ່ມປະສິດທິພາບເຫຼົ່ານີ້ສາມາດຖືກປິດການໃຊ້ງານໂດຍໃຊ້ຄໍາສໍາຄັນ. volatile: ... ການອ້າງອິງໃດໆກັບຕົວແປນີ້ຖືກລວບລວມເຂົ້າໄປໃນຫນ່ວຍຄວາມຈໍາທີ່ແທ້ຈິງທີ່ອ່ານຫຼືຂຽນຄໍາແນະນໍາ.

ການພິຈາລະນາຄວາມຖືກຕ້ອງບໍ່ຮຽກຮ້ອງໃຫ້ມີຄໍາຮ້ອງສະຫມັກ volatile. ຖ້າກະທູ້ປະຕິບັດໃຊ້ຄ່າຈາກຖານຄວາມຈໍາຈາກການດໍາເນີນການທີ່ອ່ານກ່ອນຫນ້ານັ້ນ, ມັນຈະໃຊ້ຂໍ້ມູນລ້າສະໄຫມເລັກນ້ອຍ. ແຕ່ຢ່າງໃດກໍ່ຕາມ, ນີ້ແມ່ນຂໍ້ມູນຈາກສະຖານະທີ່ຖືກຕ້ອງຂອງຕາຕະລາງ hash ໃນຊ່ວງເວລາທີ່ແນ່ນອນຂອງການໂທຫາເມັດ. ຖ້າທ່ານຕ້ອງການໃຊ້ຂໍ້ມູນຫຼ້າສຸດ, ທ່ານສາມາດນໍາໃຊ້ດັດສະນີ volatile, ແຕ່ຫຼັງຈາກນັ້ນການປະຕິບັດຈະຫຼຸດລົງເລັກນ້ອຍ: ອີງຕາມການທົດສອບຂອງຂ້ອຍ, ເມື່ອລຶບ 32 ລ້ານອົງປະກອບ, ຄວາມໄວຫຼຸດລົງຈາກ 500 ລ້ານການລຶບ / ວິນາທີເປັນ 450 ລ້ານລຶບ / ວິນາທີ.

ຜະລິດຕະພັນ

ໃນການທົດສອບການໃສ່ 64 ລ້ານອົງປະກອບແລະການລົບ 32 ລ້ານຂອງພວກເຂົາ, ການແຂ່ງຂັນລະຫວ່າງ std::unordered_map ແລະເກືອບບໍ່ມີຕາຕະລາງ hash ສໍາລັບ GPU:

std::unordered_map ໃຊ້ເວລາ 70 ms ໃສ່ແລະເອົາອົງປະກອບອອກແລະຫຼັງຈາກນັ້ນປ່ອຍໃຫ້ເຂົາເຈົ້າ unordered_map (ການກໍາຈັດຫຼາຍລ້ານອົງປະກອບໃຊ້ເວລາຫຼາຍ, ເພາະວ່າພາຍໃນ unordered_map ການຈັດສັນຫນ່ວຍຄວາມຈໍາຫຼາຍແມ່ນເຮັດ). ເວົ້າດ້ວຍຄວາມຊື່ສັດ, std:unordered_map ຂໍ້ຈໍາກັດທີ່ແຕກຕ່າງກັນຫມົດ. ມັນເປັນກະທູ້ CPU ດຽວຂອງການປະຕິບັດ, ສະຫນັບສະຫນູນ key-value ຂອງຂະຫນາດໃດກໍ່ຕາມ, ປະຕິບັດໄດ້ດີໃນອັດຕາການນໍາໃຊ້ສູງ, ແລະສະແດງໃຫ້ເຫັນປະສິດທິພາບທີ່ຫມັ້ນຄົງຫຼັງຈາກການລຶບຫຼາຍ.

ໄລຍະເວລາຂອງຕາຕະລາງ hash ສໍາລັບ GPU ແລະການສື່ສານລະຫວ່າງໂຄງການແມ່ນ 984 ms. ນີ້ປະກອບມີເວລາທີ່ໃຊ້ໃນການຈັດວາງຕາຕະລາງໃນຫນ່ວຍຄວາມຈໍາແລະການລຶບມັນ (ການຈັດສັນ 1 GB ຂອງຫນ່ວຍຄວາມຈໍາຫນຶ່ງຄັ້ງ, ເຊິ່ງໃຊ້ເວລາບາງເວລາໃນ CUDA), ການໃສ່ແລະລຶບອົງປະກອບ, ແລະເຮັດຊ້ໍາອີກ. ສໍາເນົາທັງຫມົດໄປແລະຈາກຫນ່ວຍຄວາມຈໍາບັດວິດີໂອໄດ້ຖືກເອົາເຂົ້າໄປໃນບັນຊີ.

ຕາຕະລາງ hash ຕົວຂອງມັນເອງໃຊ້ເວລາ 271 ms ເພື່ອເຮັດໃຫ້ສໍາເລັດ. ນີ້ປະກອບມີເວລາທີ່ໃຊ້ໂດຍການໃສ່ແລະລຶບອົງປະກອບຂອງບັດວີດີໂອ, ແລະບໍ່ໄດ້ຄໍານຶງເຖິງເວລາທີ່ໃຊ້ໃນການຄັດລອກເຂົ້າໄປໃນຫນ່ວຍຄວາມຈໍາແລະ iterating ໃນໄລຍະຕາຕະລາງຜົນໄດ້ຮັບ. ຖ້າຕາຕະລາງ GPU ມີຊີວິດຢູ່ເປັນເວລາດົນນານ, ຫຼືຖ້າຕາຕະລາງ hash ແມ່ນບັນຈຸຢູ່ໃນຫນ່ວຍຄວາມຈໍາຂອງບັດວີດີໂອທັງຫມົດ (ຕົວຢ່າງ, ເພື່ອສ້າງຕາຕະລາງ hash ທີ່ຈະໃຊ້ໂດຍລະຫັດ GPU ອື່ນແລະບໍ່ແມ່ນໂປເຊດເຊີກາງ), ຫຼັງຈາກນັ້ນ. ຜົນການທົດສອບແມ່ນກ່ຽວຂ້ອງ.

ຕາຕະລາງ hash ສໍາລັບບັດວີດີໂອສະແດງໃຫ້ເຫັນປະສິດທິພາບສູງເນື່ອງຈາກການສົ່ງຜ່ານສູງແລະການຂະຫນານການເຄື່ອນໄຫວ.

ຂໍ້ບົກພ່ອງ

ສະຖາປັດຕະຍະກໍາຕາຕະລາງ hash ມີບັນຫາຈໍານວນຫນຶ່ງທີ່ຕ້ອງລະວັງ:

Linear probing ຖືກຂັດຂວາງໂດຍການເປັນກຸ່ມ, ເຊິ່ງເຮັດໃຫ້ກະແຈໃນຕາຕະລາງຖືກວາງໄວ້ຫນ້ອຍກວ່າຢ່າງສົມບູນ.
ກະແຈບໍ່ໄດ້ຖືກຖອດອອກໂດຍໃຊ້ຟັງຊັນ delete ແລະເມື່ອເວລາຜ່ານໄປ, ພວກເຂົາກໍ່ເຮັດໃຫ້ໂຕະ.

ດັ່ງນັ້ນ, ການປະຕິບັດຂອງຕາຕະລາງ hash ສາມາດຄ່ອຍໆຫຼຸດລົງ, ໂດຍສະເພາະຖ້າມັນມີຢູ່ໃນເວລາດົນນານແລະມີການໃສ່ແລະລົບຈໍານວນຫລາຍ. ວິທີຫນຶ່ງເພື່ອຫຼຸດຜ່ອນຂໍ້ເສຍເຫຼົ່ານີ້ແມ່ນການ rehash ເຂົ້າໄປໃນຕາຕະລາງໃຫມ່ທີ່ມີອັດຕາການນໍາໃຊ້ຕ່ໍາພໍສົມຄວນແລະການກັ່ນຕອງອອກຄີອອກໃນລະຫວ່າງການ rehashing.

ເພື່ອສະແດງໃຫ້ເຫັນບັນຫາທີ່ອະທິບາຍ, ຂ້ອຍຈະໃຊ້ລະຫັດຂ້າງເທິງເພື່ອສ້າງຕາຕະລາງທີ່ມີ 128 ລ້ານອົງປະກອບແລະ loop ຜ່ານ 4 ລ້ານອົງປະກອບຈົນກ່ວາຂ້ອຍໄດ້ຕື່ມຂໍ້ມູນໃສ່ 124 ລ້ານຊ່ອງ (ອັດຕາການນໍາໃຊ້ປະມານ 0,96). ນີ້ແມ່ນຕາຕະລາງຜົນໄດ້ຮັບ, ແຕ່ລະແຖວແມ່ນ CUDA kernel ໂທຫາເພື່ອແຊກ 4 ລ້ານອົງປະກອບໃຫມ່ເຂົ້າໄປໃນຕາຕະລາງ hash:

ອັດຕາການນໍາໃຊ້
ໄລຍະເວລາການແຊກ 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 ລ້ານກະແຈ/ວິນາທີ).

ເມື່ອການນໍາໃຊ້ເພີ່ມຂຶ້ນ, ການປະຕິບັດຫຼຸດລົງ. ນີ້ບໍ່ເປັນທີ່ຕ້ອງການໃນກໍລະນີຫຼາຍທີ່ສຸດ. ຖ້າແອັບພລິເຄຊັນໃສ່ອົງປະກອບເຂົ້າໄປໃນຕາຕະລາງແລະຫຼັງຈາກນັ້ນຖິ້ມພວກມັນ (ຕົວຢ່າງ, ເມື່ອນັບຄໍາໃນປື້ມ), ນີ້ບໍ່ແມ່ນບັນຫາ. ແຕ່ຖ້າແອັບພລິເຄຊັນໃຊ້ຕາຕະລາງ hash ທີ່ໃຊ້ເວລາດົນນານ (ຕົວຢ່າງເຊັ່ນໃນບັນນາທິການຮູບພາບເພື່ອເກັບສ່ວນທີ່ບໍ່ຫວ່າງເປົ່າຂອງຮູບພາບທີ່ຜູ້ໃຊ້ມັກຈະໃສ່ແລະລຶບຂໍ້ມູນ), ຫຼັງຈາກນັ້ນພຶດຕິກໍານີ້ສາມາດມີບັນຫາ.

ແລະວັດແທກຄວາມເລິກຂອງຕາຕະລາງ hash ຫຼັງຈາກ 64 ລ້ານ inserts (ປັດໄຈການນໍາໃຊ້ 0,5). ຄວາມເລິກສະເລ່ຍແມ່ນ 0,4774, ດັ່ງນັ້ນກະແຈສ່ວນໃຫຍ່ຢູ່ໃນຊ່ອງທີ່ດີທີ່ສຸດທີ່ເປັນໄປໄດ້ຫຼືຫນຶ່ງຊ່ອງຫ່າງຈາກຕໍາແຫນ່ງທີ່ດີທີ່ສຸດ. ຄວາມເລິກຂອງສຽງສູງສຸດແມ່ນ 60.

ຫຼັງຈາກນັ້ນຂ້າພະເຈົ້າໄດ້ວັດແທກຄວາມເລິກຂອງ probing ໃນຕາຕະລາງທີ່ມີ 124 ລ້ານ inserts (ປັດໄຈການນໍາໃຊ້ 0,97). ຄວາມເລິກສະເລ່ຍແມ່ນແລ້ວ 10,1757, ແລະສູງສຸດ - 6474 (!!). ປະສິດທິພາບການຮັບຮູ້ເສັ້ນຊື່ຫຼຸດລົງຢ່າງຫຼວງຫຼາຍໃນອັດຕາການນໍາໃຊ້ສູງ.

ມັນເປັນສິ່ງທີ່ດີທີ່ສຸດທີ່ຈະຮັກສາອັດຕາການນໍາໃຊ້ຂອງຕາຕະລາງ hash ນີ້ຕໍ່າ. ແຕ່ຫຼັງຈາກນັ້ນພວກເຮົາເພີ່ມທະວີການປະຕິບັດໃນຄ່າໃຊ້ຈ່າຍຂອງການບໍລິໂພກຄວາມຊົງຈໍາ. ໂຊກດີ, ໃນກໍລະນີຂອງກະແຈ 32-bit ແລະຄ່າ, ນີ້ສາມາດສົມເຫດສົມຜົນ. ຖ້າໃນຕົວຢ່າງຂ້າງເທິງ, ໃນຕາຕະລາງທີ່ມີ 128 ລ້ານອົງປະກອບ, ພວກເຮົາຮັກສາປັດໃຈການນໍາໃຊ້ 0,25, ຫຼັງຈາກນັ້ນພວກເຮົາສາມາດວາງບໍ່ມີຫຼາຍກ່ວາ 32 ລ້ານອົງປະກອບໃນນັ້ນ, ແລະສ່ວນທີ່ເຫຼືອ 96 ລ້ານຊ່ອງຈະສູນເສຍ - 8 bytes ສໍາລັບແຕ່ລະຄູ່. , 768 MB ຂອງຄວາມຊົງຈໍາທີ່ສູນເສຍໄປ.

ກະລຸນາຮັບຊາບວ່າພວກເຮົາກໍາລັງເວົ້າກ່ຽວກັບການສູນເສຍຄວາມຊົງຈໍາຂອງບັດວີດີໂອ, ເຊິ່ງເປັນຊັບພະຍາກອນທີ່ມີຄຸນຄ່າຫຼາຍກ່ວາຄວາມຊົງຈໍາຂອງລະບົບ. ເຖິງແມ່ນວ່າບັດກາຟິກ desktop ທີ່ທັນສະໄຫມທີ່ສຸດທີ່ສະຫນັບສະຫນູນ CUDA ມີຫນ່ວຍຄວາມຈໍາຢ່າງຫນ້ອຍ 4 GB (ໃນເວລາຂຽນ, NVIDIA 2080 Ti ມີ 11 GB), ມັນຈະບໍ່ເປັນການຕັດສິນໃຈທີ່ສະຫລາດທີ່ສຸດທີ່ຈະສູນເສຍຈໍານວນດັ່ງກ່າວ.

ຕໍ່ມາຂ້ອຍຈະຂຽນເພີ່ມເຕີມກ່ຽວກັບການສ້າງຕາຕະລາງ hash ສໍາລັບບັດວີດີໂອທີ່ບໍ່ມີບັນຫາກັບຄວາມເລິກຂອງ probing, ເຊັ່ນດຽວກັນກັບວິທີການທີ່ຈະໃຊ້ຄືນໃຫມ່ຂອງຊ່ອງທີ່ຖືກລົບ.

ການວັດແທກຄວາມເລິກຂອງສຽງ

ເພື່ອກໍານົດຄວາມເລິກ probing ຂອງກະແຈ, ພວກເຮົາສາມາດສະກັດ hash ຂອງຄີ (ດັດຊະນີຕາຕະລາງທີ່ເຫມາະສົມຂອງມັນ) ຈາກດັດຊະນີຕາຕະລາງຕົວຈິງຂອງມັນ:

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

ເນື່ອງຈາກວ່າ magic ຂອງຕົວເລກສອງຂອງສອງທີ່ສົມບູນຂອງສອງແລະຄວາມຈິງທີ່ວ່າຄວາມສາມາດຂອງຕາຕະລາງ hash ແມ່ນສອງກັບພະລັງງານຂອງສອງ, ວິທີການນີ້ຈະເຮັດວຽກເຖິງແມ່ນວ່າໃນເວລາທີ່ດັດຊະນີທີ່ສໍາຄັນໄດ້ຖືກຍ້າຍໄປຈຸດເລີ່ມຕົ້ນຂອງຕາຕະລາງ. ໃຫ້ເອົາລະຫັດທີ່ hashed ເປັນ 1, ແຕ່ຖືກໃສ່ເຂົ້າໄປໃນຊ່ອງສຽບ 3. ຫຼັງຈາກນັ້ນສໍາລັບຕາຕະລາງທີ່ມີຄວາມຈຸ 4 ພວກເຮົາໄດ້ຮັບ. (3 — 1) & 3, ຊຶ່ງເທົ່າກັບ 2 .

ສະຫລຸບ

ຖ້າທ່ານມີຄໍາຖາມຫຼືຄໍາເຫັນ, ກະລຸນາສົ່ງອີເມວຫາຂ້ອຍທີ່ Twitter ຫຼືເປີດຫົວຂໍ້ໃຫມ່ໃນ ຄັງເກັບມ້ຽນ.

ລະຫັດນີ້ຖືກຂຽນພາຍໃຕ້ການດົນໃຈຈາກບົດຄວາມທີ່ດີເລີດ:

ໃນອະນາຄົດ, ຂ້າພະເຈົ້າຈະສືບຕໍ່ຂຽນກ່ຽວກັບການປະຕິບັດຕາຕະລາງ hash ສໍາລັບບັດວີດີໂອແລະວິເຄາະການປະຕິບັດຂອງພວກເຂົາ. ແຜນການຂອງຂ້ອຍປະກອບມີລະບົບຕ່ອງໂສ້, Robin Hood hashing, ແລະ cuckoo hashing ໂດຍໃຊ້ການປະຕິບັດປະລໍາມະນູໃນໂຄງສ້າງຂໍ້ມູນທີ່ເປັນມິດກັບ GPU.

ແຫຼ່ງຂໍ້ມູນ: www.habr.com

ຕາຕະລາງ hash ງ່າຍດາຍສໍາລັບ GPU