CUDA Kullanarak Sıralama Algoritması

 (ashwanirathee.com)
3 puan yazan GN⁺ 2025-03-13 | 2 yorum | WhatsApp'ta paylaş
  • CUDA kullanarak paralel hesaplama üzerinden sıralama algoritmasının performansını iyileştirmenin yöntemi
  • Temel bir merge sort'u CUDA ile uygulayarak performans iyileştirme olasılığını inceleme

Temel özyinelemeli merge sort (CPU uygulaması)

  • Diziyi iki alt diziye bölüp her birini sıraladıktan sonra birleştiren bir sıralama algoritması
  • Dizi özyinelemeli olarak bölünür, boyut 1 olduğunda birleştirme işlemi gerçekleştirilir
  • Uygulamayla ilgili başlıca noktalar
    • uint8_t kullanımı → küçük değerlerle (0~255) bellek kullanımını en aza indirme
    • long long kullanımı → büyük dizileri (en fazla 10¹⁸) işleyebilme
    • Performans karşılaştırması için sonuçlar std::sort ile doğrulanır
    • Zaman karmaşıklığı: ortalama O(n log n)
    • Alan karmaşıklığı: O(n)

CUDA'da temel özyinelemeli merge sort

  • CPU uygulamasıyla aynı deseni izler
  • Birleştirme işlemi CUDA'da paralel çalışacak şekilde uygulanır
  • Uygulamayla ilgili başlıca noktalar
    • cudaMalloc, cudaMemcpy, cudaFree kullanımı → GPU belleği ayırma ve veri aktarımı
    • merge<<<1, 1>>>(...) → birleştirme işlemini paralel çalıştırma
    • cudaDeviceSynchronize() → birleştirme tamamlanana kadar senkronizasyon
    • Performans sorunu → CUDA özyinelemeyi verimsiz işlediği için yinelemeli bir yaklaşım gerekir

CPU ve GPU uygulamalarının karşılaştırması

  • Özyinelemeli çağrılar CPU'da çalıştığı için performans düşüşü oluşur
  • CUDA'da özyinelemeli çağrılar yığın boyutu sorunları ve kernel çalıştırma ek yükü oluşturur
  • Performans iyileştirme yöntemi: yinelemeli (bottom-up) yaklaşıma geçiş gerekli

Bottom-up yinelemeli merge sort (CPU uygulaması)

  • Küçük alt dizilerden başlayarak kademeli birleştirme → CUDA'da daha verimli
  • Birleştirilen dizi boyutu 1, 2, 4, 8, … şeklinde artırılarak birleştirme yapılır
  • Başlıca kod yapısı 
    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
  • Uygulamayla ilgili başlıca noktalar
    • Dizi boyutu 2'nin katı olmadığında indeksleri clamp ederek sorun çözülür
    • Döngüler üzerinden birleştirme işlemi gerçekleştirilir
    • Performans iyileştirme potansiyeli yüksektir

Bottom-up yinelemeli merge sort (CUDA uygulaması)

  • Yinelemeli merge sort paralel çalıştırılarak performans iyileştirilir
  • Birleştirme işini paralel yürütmek için thread ve block sayıları hesaplanıp çalıştırılır
  • Başlıca kod yapısı 
      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));  
  }
  • Başlıca noktalar
    • flipflop → birleştirme sonucunun kaydedildiği konumu değiştirme
    • gridSize, THREADS_PER_BLOCK → birleştirme işini paralelleştirme
    • mergeKernel → her thread'e benzersiz bir birleştirme işi atama
    • Thread ve block indeks hesaplamalarıyla indeks yönetimi

Performans sonuçları

  • Bottom-up merge sort (CUDA) → performans iyileşmesi nettir
    • Küçük diziler → CPU daha hızlı
    • Büyük diziler → CUDA performans avantajı sağlar
  • thrust::sort → büyük dizilerde GPU performansı güçlü
  • CUDA'nın performans iyileştirmesi veri aktarım ek yüküyle sınırlanır

Sonuç ve gelecek çalışmalar

  • CUDA tabanlı merge sort'ta performans iyileştirmesi başarıyla sağlandı
  • Öğrenilen başlıca noktalar:
    • CUDA'nın paralel işleme kavramları ve performans tuning stratejileri öğrenildi
    • Yinelemeli merge sort → CUDA'da özyinelemeli yaklaşımdan daha etkilidir
    • Thread senkronizasyonu, bellek aktarımı gibi CUDA'ya özgü performans darboğazları keşfedildi
  • Gelecekteki iyileştirme çalışmaları:
    • CPU-GPU arasında iş bölümünün ayrılması ve optimize edilmesi
    • Daha büyük diziler için performans testleri
    • thrust::sort ile kullanıcı uygulaması kodunun birleştirilmesi
    • Paylaşılan bellek kullanımıyla performans optimizasyonu

İlgili okumalar

2 yorum

 
kandk 2025-03-14

CUDA'da radix sort ile uygulanmış gibi görünüyor; ben de referanstakini birebir uygulama deneyimi yaşamıştım.
 
GN⁺ 2025-03-13
Hacker News görüşleri

  • Bu, GPU'da hızlı sıralama yapmanın yolu değil. CUDA'daki en hızlı algoritma Onesweep; GPU paralelliğinden yararlanmak ve sınırlamaların üstesinden gelmek için karmaşık teknikler kullanıyor

    • Linebender, bu fikirleri GPU'ya daha taşınabilir şekilde uygulamak için çalışmalar yürütüyor
    • İlgili materyaller Linebender'ın wiki sayfasında bulunabilir

  • Diğer görüşlerde de belirtildiği gibi bu algoritma uygun değil. Onesweep gibi algoritmalar etkileyici, ancak anlaşılması zor

    • Temel algoritma olan radix sort'a bakmak daha kolay anlaşılmasını sağlayabilir
    • Radix sort, paralelleştirilmesi son derece kolay bir şekilde uygulanabilir ve güzel, zarif bir yaklaşım

  • Eğlenceli bir oyuncak problem olarak ele almak için iyi. Thread ayarlama seçenekleri kullanılırsa performans artışı olabilir

    • Performansı etkileyen unsurları anlamak için Nsight kullanmak da eğlenceli
    • Diğer sıralama algoritmalarını da düşünmek gerekiyor. Bitonic sort gibi sorting network'ler daha fazla iş gerektirir, ancak paralel donanımda daha hızlı çalışabilir
    • H100'de 10M öğeyi yaklaşık 10ms'de sıralayan basit bir implementasyon yaptım. Daha fazla çalışmayla daha da hızlandırılabilir

  • Futhark dili, bu tür algoritmaların GPU'da daha rahat kullanılmasını sağlıyor

    • GPU komutlarına derlenen çok yüksek seviyeli bir dil ve Python kütüphanesi üzerinden erişilebiliyor
    • Web sitesinde merge sort implementasyonu örneği var

  • Üniversite yıllarında CUDA'da bitonic sort gerçekleştirdiğim küçük bir proje aklıma geldi

    • Bitonic sort implementasyonu GitHub'da görülebilir

  • GPU thread indexing kavramını açıklayan notlar iyi

    • Vektörleştirilmiş sıralamanın performans avantajlarına dair kişisel bir blog yazısını tanıtıyor

  • Hızlı bir radix sort implementasyonunu beğendim

    • Cuda driver API ile kolay çalışıyor ve CUB'dan farklı olarak runtime API ile sınırlı değil
    • Onesweep de dahil edilmiş, ancak çalıştırmayı başaramadım

  • Unity ile birlikte kullanmayı denedim ama veri aktarımı darboğazını aşamadım

    • Compute shader kullanıldığında da ek yük var, ama çok büyük değil

  • GPU'da sıralamanın değerli olması için büyük diziler gerekiyor

    • RAM ile GPU arasındaki veri aktarımı zaman alıyor

  • Zaman kazandırması için özetle: birisi GPU'da bir sıralama algoritması yazmış ama yavaştı

    • Son teknoloji değil ve yazarın GPU'yu etkili kullanmayı bilip bilmediği belirsiz
    • Sadece kişisel bir GPU programlama deneyi