ฉันโพสต์ไว้บน Github โปรเจ็กต์ใหม่ ตารางแฮช 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 ต้องใช้คำสั่งด่วนเพียงคำสั่งเดียว แต่ตัวดำเนินการโมดูลัสจะช้ากว่ามาก นี่เป็นสิ่งสำคัญในกรณีของการตรวจวัดเชิงเส้น เนื่องจากในการค้นหาตารางเชิงเส้น ดัชนีช่องจะต้องถูกล้อมไว้ในแต่ละช่อง และเป็นผลให้ต้นทุนการดำเนินการถูกเพิ่มแบบโมดูโลในแต่ละช่อง

ตารางจะจัดเก็บเฉพาะคีย์และค่าสำหรับแต่ละองค์ประกอบ ไม่ใช่แฮชของคีย์ เนื่องจากตารางเก็บเฉพาะคีย์ 32 บิต แฮชจึงถูกคำนวณอย่างรวดเร็ว โค้ดด้านบนใช้แฮช Murmur3 ซึ่งดำเนินการกะ XOR และการคูณเพียงไม่กี่ครั้งเท่านั้น

ตารางแฮชใช้เทคนิคการป้องกันการล็อคที่ไม่ขึ้นกับลำดับหน่วยความจำ แม้ว่าการดำเนินการเขียนบางอย่างจะขัดขวางลำดับของการดำเนินการอื่นๆ ตารางแฮชจะยังคงรักษาสถานะที่ถูกต้องไว้ เราจะพูดถึงเรื่องนี้ด้านล่าง เทคนิคนี้ใช้งานได้ดีกับการ์ดแสดงผลที่ทำงานหลายพันเธรดพร้อมกัน

คีย์และค่าในตารางแฮชถูกเตรียมใช้งานให้ว่างเปล่า

โค้ดสามารถปรับเปลี่ยนเพื่อรองรับคีย์และค่า 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 การดำเนินการใดๆ ที่เรียกเคอร์เนลจะเสร็จสิ้นก่อนการเรียกเคอร์เนลครั้งถัดไปเสมอ (อย่างน้อยภายในหนึ่งเธรด ในเธรดที่ต่างกัน เคอร์เนลจะถูกดำเนินการแบบขนาน) ในตัวอย่างข้างต้น หากคุณเรียกเคอร์เนลตัวใดตัวหนึ่งด้วย 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 มิลลิวินาทีในการแทรกและลบองค์ประกอบ จากนั้นจึงปล่อยองค์ประกอบเหล่านั้นออก unordered_map (การกำจัดองค์ประกอบนับล้านต้องใช้เวลามากเพราะภายใน unordered_map มีการจัดสรรหน่วยความจำหลายรายการ) พูดตามตรงว่า std:unordered_map ข้อจำกัดที่แตกต่างอย่างสิ้นเชิง เป็นเธรดการประมวลผล CPU เดี่ยว รองรับคีย์-ค่าทุกขนาด ทำงานได้ดีที่อัตราการใช้งานที่สูง และแสดงประสิทธิภาพที่เสถียรหลังจากการลบหลายครั้ง

ระยะเวลาของตารางแฮชสำหรับ GPU และการสื่อสารระหว่างโปรแกรมคือ 984 ms ซึ่งรวมถึงเวลาที่ใช้ในการวางตารางในหน่วยความจำและการลบออก (การจัดสรรหน่วยความจำ 1 GB หนึ่งครั้งซึ่งใช้เวลาพอสมควรใน CUDA) การแทรกและการลบองค์ประกอบ และการวนซ้ำองค์ประกอบเหล่านั้น สำเนาทั้งหมดเข้าและออกจากหน่วยความจำการ์ดแสดงผลจะถูกนำมาพิจารณาด้วย

ตารางแฮชใช้เวลาดำเนินการ 271 มิลลิวินาที ซึ่งรวมถึงเวลาที่ใช้โดยการ์ดแสดงผลในการแทรกและการลบองค์ประกอบ และไม่คำนึงถึงเวลาที่ใช้ในการคัดลอกลงในหน่วยความจำและการวนซ้ำตารางผลลัพธ์ หากตาราง GPU ใช้งานได้นานหรือหากตารางแฮชนั้นอยู่ในหน่วยความจำของการ์ดแสดงผลทั้งหมด (ตัวอย่างเช่นเพื่อสร้างตารางแฮชที่จะใช้โดยโค้ด GPU อื่นและไม่ใช่โปรเซสเซอร์กลาง) จากนั้น ผลการทดสอบมีความเกี่ยวข้อง

ตารางแฮชสำหรับการ์ดแสดงผลแสดงให้เห็นถึงประสิทธิภาพสูงเนื่องจากมีปริมาณงานสูงและการทำงานแบบขนานที่ใช้งานอยู่

ข้อ จำกัด

สถาปัตยกรรมตารางแฮชมีปัญหาบางประการที่ต้องระวัง:

การตรวจสอบเชิงเส้นถูกขัดขวางโดยการจัดกลุ่ม ซึ่งทำให้คีย์ในตารางวางได้ไม่สมบูรณ์
ปุ่มจะไม่ถูกลบออกโดยใช้ฟังก์ชัน delete และเมื่อเวลาผ่านไปพวกเขาก็เกะกะโต๊ะ

เป็นผลให้ประสิทธิภาพของตารางแฮชสามารถค่อยๆ ลดลง โดยเฉพาะอย่างยิ่งหากมีอยู่เป็นเวลานานและมีส่วนแทรกและการลบจำนวนมาก วิธีหนึ่งในการลดข้อเสียเหล่านี้คือการแฮชใหม่ลงในตารางใหม่ที่มีอัตราการใช้งานที่ค่อนข้างต่ำ และกรองคีย์ที่ถูกลบออกในระหว่างการแฮชใหม่

เพื่ออธิบายปัญหาที่อธิบายไว้ ฉันจะใช้โค้ดด้านบนเพื่อสร้างตารางที่มีองค์ประกอบ 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 บิต สิ่งนี้สามารถพิสูจน์ได้ หากในตัวอย่างข้างต้น ในตารางที่มีองค์ประกอบ 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 แต่ถูกแทรกลงในช่อง XNUMX จากนั้นสำหรับตารางที่มีความจุ XNUMX เราจะได้ (3 — 1) & 3ซึ่งเท่ากับ 2

ข้อสรุป

หากคุณมีคำถามหรือความคิดเห็นกรุณาส่งอีเมลถึงฉันได้ที่ Twitter หรือเปิดหัวข้อใหม่ใน ที่เก็บ.

รหัสนี้เขียนขึ้นภายใต้แรงบันดาลใจจากบทความที่ยอดเยี่ยม:

ในอนาคต ฉันจะเขียนเกี่ยวกับการใช้งานตารางแฮชสำหรับการ์ดวิดีโอต่อไปและวิเคราะห์ประสิทธิภาพของพวกเขา แผนของฉันรวมถึงการผูกมัด การแฮชของ Robin Hood และการแฮชของนกกาเหว่าโดยใช้การดำเนินการแบบอะตอมมิกในโครงสร้างข้อมูลที่เป็นมิตรกับ GPU

ที่มา: will.com

ตารางแฮชอย่างง่ายสำหรับ GPU