อัลกอริทึมการเรียงลำดับด้วย CUDA

• วิธีปรับปรุงประสิทธิภาพของอัลกอริทึมการเรียงลำดับด้วยการประมวลผลแบบขนานโดยใช้ CUDA
• สำรวจความเป็นไปได้ในการเพิ่มประสิทธิภาพด้วยการนำ merge sort พื้นฐานไปใช้งานด้วย CUDA

merge sort แบบเรียกซ้ำพื้นฐาน (การใช้งานบน CPU)

• เป็นอัลกอริทึมการเรียงลำดับที่แบ่งอาร์เรย์ออกเป็น 2 อาร์เรย์ย่อย จากนั้นเรียงลำดับแต่ละส่วนแล้วจึงรวมเข้าด้วยกัน
• แบ่งอาร์เรย์แบบเรียกซ้ำ และเมื่อขนาดเหลือ 1 จะทำขั้นตอนการรวม
• ประเด็นสำคัญในการใช้งาน
  • ใช้ uint8_t → ลดการใช้หน่วยความจำด้วยค่าขนาดเล็ก (0~255)
  • ใช้ long long → รองรับอาร์เรย์ขนาดใหญ่ (สูงสุด 10¹⁸)
  • ตรวจสอบผลลัพธ์ด้วย std::sort เพื่อเปรียบเทียบประสิทธิภาพ
  • ความซับซ้อนด้านเวลา: O(n log n)
  • ความซับซ้อนด้านพื้นที่: O(n)

merge sort แบบเรียกซ้ำพื้นฐานใน CUDA

• ใช้รูปแบบเดียวกับการใช้งานบน CPU
• ออกแบบให้ขั้นตอนการรวมทำงานแบบขนานบน CUDA
• ประเด็นสำคัญในการใช้งาน
  • ใช้ cudaMalloc, cudaMemcpy, cudaFree → จัดสรรหน่วยความจำบน GPU และโอนย้ายข้อมูล
  • merge<<<1, 1>>>(...) → รันขั้นตอนการรวมแบบขนาน
  • cudaDeviceSynchronize() → ซิงโครไนซ์จนกว่าการรวมจะเสร็จสิ้น
  • ปัญหาด้านประสิทธิภาพ → CUDA ไม่มีประสิทธิภาพกับการเรียกซ้ำ จึงจำเป็นต้องใช้แนวทางแบบวนซ้ำ

เปรียบเทียบการใช้งานบน CPU และ GPU

• ประสิทธิภาพลดลงเพราะการเรียกซ้ำทำงานบน CPU
• การเรียกซ้ำใน CUDA มีปัญหาเรื่องขนาดสแตกและ kernel execution overhead
• วิธีปรับปรุงประสิทธิภาพ: จำเป็นต้องเปลี่ยนเป็นแนวทางแบบวนซ้ำ (bottom-up)

merge sort แบบวนซ้ำจากล่างขึ้นบน (การใช้งานบน CPU)

• รวมอาร์เรย์ย่อยขนาดเล็กทีละขั้น → มีประสิทธิภาพมากกว่าใน CUDA
• เพิ่มขนาดของอาร์เรย์ที่นำมารวมเป็น 1, 2, 4, 8, … แล้วรวมต่อเนื่อง
• โครงสร้างโค้ดหลัก

  MERGE_SORT(arr, temp, start, end)  
    FOR sub_size ← 1 TO end STEP 2 × sub_size DO  
        FOR left ← 0 TO end STEP 2 × sub_size DO  
            mid ← MIN(left + sub_size - 1, end)  
            right ← MIN(left + 2 * sub_size - 1, end)  
            MERGE(arr, temp, left, mid, right)  
        ENDFOR  
    ENDFOR  
  END MERGE_SORT  
  

• ประเด็นสำคัญในการใช้งาน
  • หากขนาดอาร์เรย์ไม่เป็นจำนวนเท่าของ 2 จะแก้ปัญหาโดย clamp ดัชนี
  • ทำขั้นตอนการรวมผ่านลูป
  • มีศักยภาพในการเพิ่มประสิทธิภาพสูง

merge sort แบบวนซ้ำจากล่างขึ้นบน (การใช้งานบน CUDA)

• ปรับปรุงประสิทธิภาพด้วยการรัน merge sort แบบวนซ้ำให้ทำงานแบบขนาน
• เพื่อให้ขั้นตอนการรวมทำงานแบบขนาน จะคำนวณจำนวน thread และ block ก่อนแล้วจึงรัน
• โครงสร้างโค้ดหลัก

    void mergeSort(uint8_t* arr, uint8_t* temp, long long n) {  
        bool flipflop = true;  
        long long size;  
        for (size = 1; size < n; size *= 2) {  
            numThreads = max(n / (2 * size), (long long)1);  
            gridSize = (numThreads + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;  
            mergeKernel<<<gridSize, THREADS_PER_BLOCK>>>(flipflop ? arr : temp, flipflop ? temp : arr, size, n);  
            CUDA_CHECK(cudaGetLastError());  
            CUDA_CHECK(cudaDeviceSynchronize());  
            flipflop = !flipflop;  
        }  
        if (!flipflop) CUDA_CHECK(cudaMemcpy(arr, temp, n * sizeof(uint8_t), cudaMemcpyDeviceToDevice));  
    }  
  

• ประเด็นสำคัญ
  • flipflop → สลับตำแหน่งจัดเก็บผลลัพธ์ของการรวม
  • gridSize, THREADS_PER_BLOCK → ใช้สำหรับทำให้ขั้นตอนการรวมทำงานแบบขนาน
  • mergeKernel → กำหนดงานรวมที่ไม่ซ้ำกันให้แต่ละ thread
  • จัดการดัชนีผ่านการคำนวณดัชนีของ thread และ block

ผลลัพธ์ด้านประสิทธิภาพ

• merge sort แบบ bottom-up (CUDA) → ปรับปรุงประสิทธิภาพได้ชัดเจน
  • อาร์เรย์ขนาดเล็ก → CPU เร็วกว่า
  • อาร์เรย์ขนาดใหญ่ → CUDA เหนือกว่าด้านประสิทธิภาพ
• thrust::sort → GPU มีประสิทธิภาพยอดเยี่ยมกับอาร์เรย์ขนาดใหญ่
• การปรับปรุงประสิทธิภาพของ CUDA ถูก จำกัดด้วย overhead จากการโอนย้ายข้อมูล

บทสรุปและงานต่อไป

• ประสบความสำเร็จในการปรับปรุงประสิทธิภาพของ merge sort บนพื้นฐาน CUDA
• สิ่งสำคัญที่ได้เรียนรู้:
  • เรียนรู้แนวคิดการประมวลผลแบบขนานของ CUDA และกลยุทธ์การจูนประสิทธิภาพ
  • merge sort แบบวนซ้ำ → มีประสิทธิภาพบน CUDA มากกว่าแนวทางแบบเรียกซ้ำ
  • พบคอขวดด้านประสิทธิภาพเฉพาะของ CUDA เช่น การซิงโครไนซ์ thread และการโอนย้ายหน่วยความจำ
• งานปรับปรุงในอนาคต:
  • แยกและปรับแต่งงานระหว่าง CPU-GPU
  • ทดสอบประสิทธิภาพกับอาร์เรย์ที่ใหญ่ขึ้น
  • ผสาน thrust::sort เข้ากับโค้ดที่พัฒนาขึ้นเอง
  • ปรับแต่งประสิทธิภาพด้วยการใช้ shared memory

[kandk]

ดูเหมือนว่าจะมีการทำเป็น radix sort ไว้ใน CUDA และผมก็เคยมีประสบการณ์ลองทำตามตัวอย่างอ้างอิงแบบเดียวกันเลยครับ

[GN⁺]

ความคิดเห็นจาก Hacker News

• ไม่ใช่วิธีจัดเรียงบน GPU ที่เร็วที่สุด อัลกอริทึมที่เร็วที่สุดบน CUDA คือ Onesweep ซึ่งใช้เทคนิคซับซ้อนเพื่อดึงความขนานของ GPU มาใช้และก้าวข้ามข้อจำกัดต่าง ๆ

  • Linebender กำลังทำงานเพื่อนำแนวคิดเหล่านี้ไปใช้กับ GPU ในแบบที่พกพาได้มากขึ้น
  • สามารถดูข้อมูลที่เกี่ยวข้องได้จากหน้าวิกิของ Linebender

• เช่นเดียวกับความเห็นอื่น ๆ อัลกอริทึมนี้ไม่เหมาะสม แม้อัลกอริทึมอย่าง Onesweep จะน่าทึ่ง แต่ก็เข้าใจได้ยาก

  • หากดูที่อัลกอริทึมหลักอย่าง radix sort จะเข้าใจได้ง่ายกว่า
  • radix sort สามารถทำให้ขนานได้ง่ายมาก และเป็นแนวทางที่สวยงามและเรียบง่าย

• เป็นปัญหาแนวของเล่นที่หยิบมาลองเล่นได้สนุก หากใช้ตัวเลือกการปรับจูน thread ก็อาจช่วยเพิ่มประสิทธิภาพได้

  • การใช้ Nsight เพื่อดูว่าอะไรส่งผลต่อประสิทธิภาพก็สนุกดีเช่นกัน
  • ควรพิจารณาอัลกอริทึมจัดเรียงแบบอื่นด้วย network sort อย่าง bitonic sort ต้องทำงานมากกว่า แต่สามารถรันบนฮาร์ดแวร์แบบขนานได้เร็วกว่า
  • เคยทำ implementation แบบง่ายที่จัดเรียง 10M บน H100 ได้ในเวลาประมาณ 10ms และยังทำให้เร็วกว่านี้ได้อีกหากลงงานเพิ่ม

• ภาษา Futhark ทำให้อัลกอริทึมลักษณะนี้ใช้งานบน GPU ได้สะดวกขึ้น

  • เป็นภาษาระดับสูงมากที่คอมไพล์เป็นคำสั่ง GPU และเข้าถึงได้ผ่านไลบรารี Python
  • บนเว็บไซต์มีตัวอย่าง implementation ของ merge sort

• ทำให้นึกถึงโปรเจ็กต์เล็ก ๆ ตอนเรียนมหาวิทยาลัยที่เคยเขียน bitonic sort บน CUDA

  • สามารถดู implementation ของ bitonic sort ได้บน GitHub

• โน้ตที่อธิบายแนวคิดการทำ indexing ของ thread บน GPU นั้นดี

  • มีการแนะนำบล็อกโพสต์ส่วนตัวเกี่ยวกับข้อได้เปรียบด้านประสิทธิภาพของการจัดเรียงแบบเวกเตอร์

• ชอบ implementation ของ radix sort ที่เร็ว

  • ทำงานร่วมกับ Cuda driver API ได้ง่าย และไม่ถูกจำกัดอยู่แค่ runtime API แบบ CUB
  • มี Onesweep รวมมาด้วย แต่ยังทำให้มันทำงานไม่ได้

• เคยพยายามใช้ร่วมกับ Unity แต่ไม่สามารถเอาชนะคอขวดจากการส่งข้อมูลได้

  • แม้ตอนใช้ compute shader ก็ยังมี overhead อยู่ แต่ไม่มากนัก

• ถ้าจะให้คุ้มกับการจัดเรียงบน GPU ก็ต้องเป็นอาร์เรย์ขนาดใหญ่

  • การส่งข้อมูลระหว่าง RAM กับ GPU ใช้เวลา

• สรุปแบบประหยัดเวลา: มีคนเขียนอัลกอริทึมจัดเรียงบน GPU ขึ้นมา แต่ช้า

  • ไม่ใช่เทคโนโลยีล่าสุด และยังไม่ชัดเจนว่าผู้เขียนรู้วิธีใช้ GPU อย่างมีประสิทธิภาพหรือไม่
  • เป็นเพียงการทดลองเขียนโปรแกรม GPU ส่วนตัวเท่านั้น