CUDAを活用したソートアルゴリズム

(ashwanirathee.com)

3 ポイント投稿者 GN⁺ 2025-03-13 | 2件のコメント | WhatsAppで共有

CUDAを使用し、並列コンピューティングによってソートアルゴリズムの性能を改善する方法
基本的なマージソート（merge sort）をCUDAで実装し、性能改善の可能性を探る

基本的な再帰マージソート（CPU実装）

配列を2つの部分配列に分割し、それぞれをソートした後にマージする方式のソートアルゴリズム
再帰的に配列を分割し、サイズが1になったらマージ処理を実行
実装に関する主なポイント
- uint8_tを使用 → 小さな値（0〜255）でメモリ使用量を最小化
- long longを使用 → 大きな配列（最大10¹⁸）を処理可能
- 性能比較のためstd::sortで結果を検証
- 計算量: 平均 O(n log n)
- 空間計算量: O(n)

CUDAでの基本的な再帰マージソート

CPU実装と同じパターンに従う
マージ処理をCUDA上で並列実行するよう実装
実装に関する主なポイント
- cudaMalloc、cudaMemcpy、cudaFreeを使用 → GPUメモリの確保とデータ転送
- merge<<<1, 1>>>(...) → マージ処理を並列実行
- cudaDeviceSynchronize() → マージ完了まで同期を実行
- 性能上の問題 → CUDAは再帰処理に非効率なため、反復的アプローチが必要

CPUとGPU実装の比較

再帰呼び出しがCPUで実行されるため性能低下が発生
CUDAでの再帰呼び出しはスタックサイズの問題とカーネル実行オーバーヘッドを生む
性能改善方法: 反復的な（bottom-up）アプローチへの切り替えが必要

ボトムアップ反復マージソート（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の倍数でない場合はインデックスをクランプして問題を解決
- ループによってマージ処理を実行
- 性能改善の余地が大きい

ボトムアップ反復マージソート（CUDA実装）

反復マージソートを並列実行して性能を改善
マージ処理を並列に行うため、スレッド数とブロック数を計算して実行

主なコード構造

  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 → 各スレッドに固有のマージ処理を割り当て
- スレッドおよびブロックのインデックス計算によるインデックス管理

性能結果

ボトムアップマージソート（CUDA） → 性能改善は明確
- 小さな配列 → CPUの方が高速
- 大きな配列 → CUDAが性能面で優位
thrust::sort → 大きな配列ではGPU性能が優秀
CUDAの性能改善はデータ転送オーバーヘッドによって制限される

結論と今後の作業

CUDAベースのマージソート性能改善に成功
学んだ主なポイント:
- CUDAの並列処理の概念と性能チューニング戦略を学習
- 反復マージソート → 再帰的アプローチよりCUDAで効果的
- スレッド同期、メモリ転送など、CUDA固有の性能ボトルネックを発見
今後の改善項目:
- CPU-GPU間の処理分担と最適化
- より大きな配列に対する性能テスト
- thrust::sortとユーザー実装コードの結合
- 共有メモリの使用による性能最適化

2件のコメント

kandk 2025-03-14

CUDAではradix sortで実装されていたようですが、参考資料を見て同じように実装した経験があります。

GN⁺ 2025-03-13

Hacker Newsの意見

GPUで高速にソートする方法ではない。CUDAで最速のアルゴリズムはOnesweepで、GPUの並列性を活用し、制約を乗り越えるために複雑な技術を使っている
- Linebenderは、こうしたアイデアをGPUでより移植性高く適用する取り組みを進めている
- 関連資料はLinebenderのWikiページで確認できる
他の意見と同様に、このアルゴリズムは適切ではない。Onesweepのようなアルゴリズムはすばらしいが、理解するのが難しい
- 中核アルゴリズムである基数ソートを見ると、より理解しやすい
- 基数ソートは並列化を非常に簡単に実装でき、美しくエレガントなアプローチである
面白いおもちゃ問題として扱うにはよい。スレッド調整オプションを活用すれば性能向上が見込めるかもしれない
- Nsightを使って性能に影響する要素を把握するのも面白い
- 他のソートアルゴリズムも検討する必要がある。bitonic sortのようなソーティングネットワークはより多くの作業が必要だが、並列ハードウェアではより高速に動作する可能性がある
- H100で10Mを約10msでソートする単純な実装を行った。さらに作業すればもっと高速化できる
Futhark言語は、こうしたアルゴリズムをGPU上でより便利に使えるようにしてくれる
- 非常に高水準な言語で、GPU命令にコンパイルされ、Pythonライブラリとしてアクセスできる
- Webサイトにはマージソート実装の例がある
大学時代にCUDAでbitonic sortを実装した小さなプロジェクトを思い出す
- bitonic sortの実装はGitHubで確認できる
GPUスレッドインデクシングの概念を説明したノートがよい
- ベクトル化されたソートの性能上の利点に関する個人ブログ記事を紹介している
高速な基数ソート実装が気に入っている
- CudaドライバAPIと簡単に連携でき、CUBと違ってランタイムAPIに限定されない
- Onesweepも含まれているが、動かすことはできなかった
Unityと一緒に使おうとしたが、データ転送のボトルネックを克服できなかった
- コンピュートシェーダー使用時にもオーバーヘッドはあるが、それほど大きくはない
GPUでソートする価値があるには大きな配列が必要だ
- RAMとGPU間のデータ転送には時間がかかる
時間節約のために要約すると、誰かがGPUでソートアルゴリズムを書いたが遅かった
- 最新技術ではなく、作者がGPUを効果的に使う方法を知っているのかは不明だ
- 個人的なGPUプログラミング実験にすぎない

CUDAを活用したソートアルゴリズム

基本的な再帰マージソート（CPU実装）

CUDAでの基本的な再帰マージソート

CPUとGPU実装の比較

ボトムアップ反復マージソート（CPU実装）

ボトムアップ反復マージソート（CUDA実装）

性能結果

結論と今後の作業

関連記事

2件のコメント

Hacker Newsの意見