Algoritme Pengurutan Menggunakan CUDA

 (ashwanirathee.com)
3 poin oleh GN⁺ 2025-03-13 | 2 komentar | Bagikan ke WhatsApp
  • Cara meningkatkan performa algoritme pengurutan melalui komputasi paralel menggunakan CUDA
  • Mengeksplorasi kemungkinan peningkatan performa dengan mengimplementasikan merge sort dasar menggunakan CUDA

Merge Sort Rekursif Dasar (Implementasi CPU)

  • Algoritme pengurutan yang membagi array menjadi dua subarray, mengurutkan masing-masing, lalu menggabungkannya
  • Array dibagi secara rekursif, dan ketika ukurannya menjadi 1, proses penggabungan dilakukan
  • Hal-hal utama terkait implementasi
    • Menggunakan uint8_t → meminimalkan penggunaan memori untuk nilai kecil (0~255)
    • Menggunakan long long → memungkinkan pemrosesan array besar (hingga 10¹⁸)
    • Hasil diverifikasi dengan std::sort untuk perbandingan performa
    • Kompleksitas waktu: rata-rata O(n log n)
    • Kompleksitas ruang: O(n)

Merge Sort Rekursif Dasar di CUDA

  • Mengikuti pola yang sama dengan implementasi CPU
  • Diimplementasikan agar proses penggabungan berjalan paralel di CUDA
  • Hal-hal utama terkait implementasi
    • Menggunakan cudaMalloc, cudaMemcpy, cudaFree → alokasi memori GPU dan transfer data
    • merge<<<1, 1>>>(...) → menjalankan proses penggabungan secara paralel
    • cudaDeviceSynchronize() → melakukan sinkronisasi hingga penggabungan selesai
    • Masalah performa → CUDA tidak efisien untuk pemrosesan rekursif sehingga diperlukan pendekatan iteratif

Perbandingan Implementasi CPU dan GPU

  • Penurunan performa terjadi karena pemanggilan rekursif dijalankan di CPU
  • Pemanggilan rekursif di CUDA menimbulkan masalah ukuran stack dan overhead eksekusi kernel
  • Cara meningkatkan performa: perlu beralih ke pendekatan iteratif (bottom-up)

Merge Sort Iteratif Bottom-Up (Implementasi CPU)

  • Menggabungkan mulai dari subarray kecil secara bertahap → lebih efisien di CUDA
  • Ukuran array yang digabung ditingkatkan menjadi 1, 2, 4, 8, … sambil terus melakukan merge
  • Struktur kode utama 
    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
  • Hal-hal utama terkait implementasi
    • Jika ukuran array bukan kelipatan 2, masalah diatasi dengan melakukan clamping pada indeks
    • Proses merge dilakukan melalui loop
    • Potensi peningkatan performa sangat besar

Merge Sort Iteratif Bottom-Up (Implementasi CUDA)

  • Meningkatkan performa dengan menjalankan merge sort iteratif secara paralel
  • Untuk menjalankan proses merge secara paralel, jumlah thread dan block dihitung terlebih dahulu sebelum eksekusi
  • Struktur kode utama 
      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));  
  }
  • Poin-poin utama
    • flipflop → mengganti lokasi penyimpanan hasil merge
    • gridSize, THREADS_PER_BLOCK → menjalankan paralelisasi pekerjaan merge
    • mergeKernel → menetapkan tugas merge unik ke setiap thread
    • Pengelolaan indeks melalui perhitungan indeks thread dan block

Hasil Performa

  • Merge sort bottom-up (CUDA) → peningkatan performa terlihat jelas
    • Array kecil → CPU lebih cepat
    • Array besar → CUDA unggul dalam performa
  • thrust::sort → performa GPU unggul pada array besar
  • Peningkatan performa CUDA dibatasi oleh overhead transfer data

Kesimpulan dan Pekerjaan Lanjutan

  • Berhasil meningkatkan performa merge sort berbasis CUDA
  • Hal-hal utama yang dipelajari:
    • Mempelajari konsep pemrosesan paralel CUDA dan strategi tuning performa
    • Merge sort iteratif → lebih efektif di CUDA dibanding pendekatan rekursif
    • Menemukan bottleneck performa khas CUDA seperti sinkronisasi thread dan transfer memori
  • Pekerjaan peningkatan berikutnya:
    • Memisahkan dan mengoptimalkan pekerjaan antara CPU-GPU
    • Menguji performa pada array yang lebih besar
    • Menggabungkan thrust::sort dengan kode implementasi kustom
    • Optimasi performa melalui penggunaan shared memory

Bacaan terkait

2 komentar

 
kandk 2025-03-14

Sepertinya itu diimplementasikan dengan radix sort di CUDA; saya juga pernah punya pengalaman mengimplementasikan hal yang sama sebagai referensi.
 
GN⁺ 2025-03-13
Opini Hacker News

  • Bukan cara untuk melakukan pengurutan cepat di GPU. Algoritme tercepat di CUDA adalah Onesweep, yang menggunakan teknik kompleks untuk memanfaatkan paralelisme GPU dan mengatasi keterbatasannya

    • Linebender sedang mengerjakan penerapan ide-ide ini agar lebih portabel di GPU
    • Materi terkait dapat dilihat di halaman wiki Linebender

  • Seperti pendapat lain, algoritme ini tidak cocok. Algoritme seperti Onesweep memang keren tetapi sulit dipahami

    • Jika melihat radix sort sebagai algoritme intinya, ini bisa lebih mudah dipahami
    • Radix sort dapat diimplementasikan dengan sangat mudah untuk diparalelkan, dan merupakan pendekatan yang indah serta elegan

  • Menarik untuk dijadikan toy problem. Dengan memanfaatkan opsi pengaturan thread, mungkin ada peningkatan performa

    • Menyenangkan juga menggunakan Nsight untuk memahami faktor-faktor yang memengaruhi performa
    • Perlu mempertimbangkan algoritme pengurutan lain juga. Network sort seperti bitonic sort memerlukan lebih banyak pekerjaan, tetapi bisa berjalan lebih cepat pada perangkat keras paralel
    • Saya membuat implementasi sederhana yang mengurutkan 10M dalam sekitar 10ms di H100. Dengan lebih banyak usaha, ini bisa dibuat lebih cepat

  • Bahasa Futhark membuat algoritme semacam ini lebih nyaman digunakan di GPU

    • Ini adalah bahasa tingkat sangat tinggi yang dikompilasi menjadi instruksi GPU, dan dapat diakses sebagai library Python
    • Di situs webnya ada contoh implementasi merge sort

  • Ini mengingatkan pada proyek kecil saat kuliah ketika mengimplementasikan bitonic sort di CUDA

    • Implementasi bitonic sort dapat dilihat di GitHub

  • Catatan yang menjelaskan konsep pengindeksan thread GPU sangat bagus

    • Memperkenalkan posting blog pribadi tentang keuntungan performa dari vectorized sort

  • Saya menyukai implementasi radix sort yang cepat

    • Mudah bekerja dengan Cuda driver API, dan tidak seperti CUB, tidak terbatas pada runtime API
    • Onesweep juga disertakan, tetapi saya belum berhasil membuatnya berjalan

  • Pernah mencoba menggunakannya bersama Unity, tetapi tidak berhasil mengatasi bottleneck transfer data

    • Saat menggunakan compute shader juga ada overhead, tetapi tidak terlalu besar

  • Agar layak diurutkan di GPU, diperlukan array yang besar

    • Transfer data antara RAM dan GPU memerlukan waktu

  • Untuk menghemat waktu, ringkasannya: seseorang menulis algoritme pengurutan di GPU tetapi hasilnya lambat

    • Ini bukan teknik terbaru, dan tidak jelas apakah penulisnya tahu cara menggunakan GPU secara efektif
    • Ini hanya eksperimen pribadi dalam pemrograman GPU