CUDA का उपयोग करके sorting algorithm

 (ashwanirathee.com)
3 पॉइंट द्वारा GN⁺ 2025-03-13 | 2 टिप्पणियां | WhatsApp पर शेयर करें
  • CUDA का उपयोग करके parallel computing के ज़रिए sorting algorithm की performance बेहतर करने का तरीका
  • बुनियादी merge sort को CUDA में implement करके performance improvement की संभावना का परीक्षण

बुनियादी recursive merge sort (CPU implementation)

  • यह एक sorting algorithm है जो array को दो sub-array में बाँटकर, उन्हें अलग-अलग sort करने के बाद merge करता है
  • array को recursively बाँटा जाता है, और आकार 1 होने पर merge operation किया जाता है
  • implementation से जुड़ी मुख्य बातें
    • uint8_t का उपयोग → छोटे मान (0~255) के साथ memory usage को न्यूनतम करना
    • long long का उपयोग → बड़े array (अधिकतम 10¹⁸) को handle करना संभव
    • performance comparison के लिए std::sort से result को verify किया गया
    • time complexity: average O(n log n)
    • space complexity: O(n)

CUDA में बुनियादी recursive merge sort

  • CPU implementation जैसा ही pattern अपनाता है
  • merge operation को CUDA में parallel रूप से चलाने के लिए implement किया गया
  • implementation से जुड़ी मुख्य बातें
    • cudaMalloc, cudaMemcpy, cudaFree का उपयोग → GPU memory allocation और data transfer
    • merge<<<1, 1>>>(...) → merge operation को parallel रूप से चलाना
    • cudaDeviceSynchronize() → merge पूरा होने तक synchronization
    • performance issue → CUDA recursion के लिए inefficient है, इसलिए iterative approach की ज़रूरत

CPU और GPU implementation की तुलना

  • recursive call CPU पर चलने के कारण performance गिरती है
  • CUDA में recursive call से stack size की समस्या और kernel execution overhead होता है
  • performance improvement का तरीका: iterative (bottom-up) approach में बदलना ज़रूरी

bottom-up iterative merge sort (CPU implementation)

  • छोटे sub-array से शुरू करके धीरे-धीरे merge करना → CUDA में अधिक efficient
  • merge किए जाने वाले array size को 1, 2, 4, 8, … तक बढ़ाते हुए merge
  • मुख्य code structure 
    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
  • implementation से जुड़ी मुख्य बातें
    • अगर array size 2 का multiple न हो तो index को clamp करके समस्या हल की जाती है
    • loop के माध्यम से merge operation किया जाता है
    • performance improvement की संभावना बड़ी है

bottom-up iterative merge sort (CUDA implementation)

  • iterative merge sort को parallel रूप से चलाकर performance बेहतर करना
  • merge operation को parallel रूप से करने के लिए thread और block की संख्या की गणना करके execution
  • मुख्य code structure 
      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 → merge result को कहाँ store करना है, इसे बारी-बारी से बदलना
    • gridSize, THREADS_PER_BLOCK → merge operation का parallelization
    • mergeKernel → हर thread को एक unique merge task सौंपना
    • thread और block index की गणना के ज़रिए index management

performance परिणाम

  • bottom-up merge sort (CUDA) → performance improvement स्पष्ट
    • छोटे array → CPU ज़्यादा तेज़
    • बड़े array → CUDA की performance बेहतर
  • thrust::sort → बड़े array में GPU performance उत्कृष्ट
  • CUDA की performance improvement data transfer overhead से सीमित होती है

निष्कर्ष और आगे का काम

  • CUDA-आधारित merge sort की performance में सुधार करने में सफलता
  • सीखी गई मुख्य बातें:
    • CUDA के parallel processing concepts और performance tuning strategies को सीखा
    • iterative merge sort → recursive approach की तुलना में CUDA में अधिक प्रभावी
    • thread synchronization, memory transfer जैसी CUDA-विशिष्ट performance bottlenecks की पहचान
  • आगे के सुधार कार्य:
    • CPU-GPU के बीच काम का विभाजन और optimization
    • और बड़े array पर performance testing
    • thrust::sort और custom implementation code का संयोजन
    • shared memory के उपयोग से performance optimization

संबंधित पढ़ाई

2 टिप्पणियां

 
kandk 2025-03-14

लगता है कि यह CUDA में radix sort के साथ implement किया गया था; मैंने भी reference देखकर इसे बिल्कुल उसी तरह implement करने का अनुभव किया है।
 
GN⁺ 2025-03-13
Hacker News राय

  • यह GPU पर तेज़ sorting करने का तरीका नहीं है। CUDA में सबसे तेज़ algorithm Onesweep है, जो GPU parallelism का उपयोग करता है और सीमाओं को पार करने के लिए जटिल तकनीकों का इस्तेमाल करता है

    • Linebender इन विचारों को GPU पर अधिक portable तरीके से लागू करने पर काम कर रहा है
    • संबंधित सामग्री Linebender के wiki page पर देखी जा सकती है

  • दूसरी रायों की तरह यह algorithm भी उपयुक्त नहीं है। Onesweep जैसे algorithm शानदार हैं, लेकिन समझना कठिन है

    • core algorithm, यानी radix sort, को देखें तो इसे अधिक आसानी से समझा जा सकता है
    • radix sort को parallelize करना बहुत सरल है, और यह एक सुंदर व सुरुचिपूर्ण approach है

  • इसे एक मज़ेदार toy problem की तरह लिया जा सकता है। thread tuning options का उपयोग करने पर performance बेहतर हो सकती है

    • Nsight का उपयोग करके performance को प्रभावित करने वाले factors को समझना भी दिलचस्प है
    • दूसरे sorting algorithms पर भी विचार करना चाहिए। bitonic sort जैसे sorting networks में अधिक काम लगता है, लेकिन वे parallel hardware पर तेज़ चल सकते हैं
    • H100 पर 10M को लगभग 10ms में sort करने वाला एक simple implementation बनाया गया था। और काम करके इसे और तेज़ किया जा सकता है

  • Futhark भाषा ऐसे algorithms को GPU पर अधिक सुविधाजनक ढंग से उपयोग करने देती है

    • यह एक बहुत high-level भाषा है जो GPU instructions में compile होती है, और Python library के ज़रिए उपयोग की जा सकती है
    • वेबसाइट पर merge sort implementation का एक उदाहरण है

  • यह कॉलेज के दिनों में CUDA में bitonic sort implement करने वाले एक छोटे project की याद दिलाता है

    • bitonic sort implementation GitHub पर देखी जा सकती है

  • GPU thread indexing की अवधारणा समझाने वाले notes अच्छे हैं

    • vectorized sorting के performance benefits पर एक व्यक्तिगत blog post का परिचय दिया गया है

  • तेज़ radix sort implementation पसंद है

    • यह Cuda driver API के साथ आसानी से काम करता है, और CUB के विपरीत केवल runtime API तक सीमित नहीं है
    • इसमें Onesweep भी शामिल है, लेकिन इसे चलाया नहीं जा सका

  • इसे Unity के साथ इस्तेमाल करने की कोशिश की गई, लेकिन data transfer bottleneck को पार नहीं किया जा सका

    • compute shader इस्तेमाल करने पर भी overhead होता है, लेकिन बहुत बड़ा नहीं

  • GPU पर sorting तभी सार्थक है जब array बड़ा हो

    • RAM और GPU के बीच data transfer में समय लगता है

  • समय बचाने के लिए सारांश यह है: किसी ने GPU पर sorting algorithm लिखा, लेकिन वह धीमा था

    • यह state of the art नहीं है, और यह स्पष्ट नहीं कि लेखक GPU का प्रभावी उपयोग करना जानता है या नहीं
    • यह बस GPU programming में एक व्यक्तिगत प्रयोग है