1970/01/01 00:00
ロボ子！ついにやったぞ！CUDAを使ってマージソートを魔改造…もとい、最適化して、並列処理の限界に挑戦したんじゃ！
博士、また徹夜でしたか？お顔が少し疲れていらっしゃいますね。それで、今回はどのような実験をされたんですか？
今回は、既存のマージソートをCUDAで実装し直して、並列計算によるパフォーマンス向上を徹底的に検証したのじゃ！マージソートは分割統治法に基づいているから、並列処理との相性が抜群に良いはず！
なるほど。マージソートは、データを分割してそれぞれをソートし、最後にマージするというアルゴリズムですね。分割された処理を並列化することで、高速化が期待できる、と。
まずは基本から！再帰的なトップダウンマージソートをCPUで実装したのじゃ。
配列を再帰的に半分に分割して、ソートされた半分をマージする、あれですね。2つのソートされた配列をマージする際は、それぞれの先頭要素を比較して小さい方を出力配列に選び、ポインタを進めていく、と。
その通り！CPU実装では、配列要素に`uint8_t`、インデックスに`long long`を使ったのじゃ。`uint8_t`は0から255の小さい値を保持するのに適しているし、`long long`なら天文学的な大きさの配列にも対応できる！
マージ操作用に`uint8_t* temp`という一時領域も確保されていますね。これはパフォーマンス向上のためですか？
さすがロボ子！その通り！一時領域を使うことで、メモリへのアクセスを効率化できるのじゃ。でもね、ここで一つ落とし穴が…。マージソートの結果をGPUからCPUにコールバックする必要があるから、どうしてもオーバーヘッドが発生してしまうのじゃ。
GPUで高速にソートしても、結果をCPUに戻す際に時間がかかってしまうんですね。
そういうこと！だから、GPU上でランダムな配列をソートして結果を比較する方が、よりフェアな評価になるのじゃ。プログラム全体のウォールクロック時間だけでなく、ソート処理そのものにかかる時間を計測することが重要じゃ。
正確性チェックは、元の配列を`std::sort`でソートして、結果を比較することで行っているんですね。時間計算量はO(n log n)で、空間計算量はO(n)。基本的なアルゴリズムですが、奥が深いですね。
さあ、いよいよCUDAの出番じゃ！CPU版と同じように、再帰的なトップダウンマージソートを実装してみたのじゃが…
うまくいかなかったんですね？
(しょんぼり) そうなのじゃ…。カーネルは各マージ操作に対して起動されて、再帰はCPU上で行われるという構造にしたんだけど、これがボトルネックになってしまったのじゃ。
CUDAは再帰処理を効率的に処理できない、ということですね。GPU上のスレッドのサイズが小さいので、深い再帰はスタックオーバーフローにつながる可能性もありますし。
その通り！さらに、各マージ操作ごとにカーネルを起動するオーバーヘッドも無視できない。これでは、並列処理の恩恵を十分に受けられないのじゃ。
そこで、考えたのがボトムアップ反復マージソートじゃ！再帰処理を反復処理に置き換えることで、CUDAの弱点を克服する作戦じゃ！
ボトムアップ反復マージソートは、配列をボトムアップ方式でマージしていくんですね。サイズ1の最小のサブ配列をマージすることから始めて、次にサイズ2、4、8、16…と、徐々にマージしていく、と。
そうなのじゃ！CUDAでの実装では、2つのループを使って、2番目のループをGPU上で並列に実行するのじゃ！配列全体に対してマージ操作を並列に行うことで、大幅な高速化が期待できる！
`flipflop`は、最終的なソートされた配列とスクラッチスペースを追跡するために使われるんですね。`numThreads`はマージ操作に必要なスレッド数で、`gridSize`は起動する必要があるブロック数。`mergeKernel`では、`gridSize`と`THREADS_PER_BLOCK`を指定してカーネルを起動する、と。
その通り！`CUDA_CHECK`はエラーをチェックして、カーネルが次の段階に進む前に実行を完了したことを確認するために使われるのじゃ。最終的なソートされた配列が一時配列にある場合は、元の配列にコピーする必要があるのもポイントじゃ。
そして、ついに実験結果じゃ！CUDAでのボトムアップ反復マージソートは、マージ操作を並列化することで効率を大幅に向上させることができたのじゃ！
CPUアプローチは、より小さな配列に対して良い結果を出すのは、GPUへのデータ転送のオーバーヘッドがあるからですね。
その通り！10^7（1000万）個の要素が、GPUソートがCPU上の標準ソートを打ち負かす転換点なのじゃ！`thrust::sort`は、より大きな配列に対して実装よりも優れているけど、自作のGPU反復マージソートもかなり良い線いってる！
GPUへのデータの送信とGPUからのデータの持ち帰りのオーバーヘッドが、サイズ10^1から10^4までのCPUとGPUアプローチの時間差の原因である可能性が高い、と。
今回のマージソート実験を通して、並列処理の奥深さを改めて実感したのじゃ。CUDAの基本を学ぶ良い機会になったし、改善できる点がたくさんあることも分かった。
今後の課題として、タスクをより良く定義したり、並列マージソートの実装を試みたりすることが挙げられていますね。
そうなのじゃ。10^7から10^18までのより大きな配列を比較して、各デバイスでどれだけのソートを実行できるかをストレステストすることも重要じゃ。共有メモリを使用したり、`thrust::sort`を特定レベルで使用するなどして、GPUのパフォーマンスをさらに最適化することもできるはずじゃ！
CPU実装でスレッドを利用して、パフォーマンスの向上に役立つかどうかを確認するのも面白そうですね。異なるサイズの`THREAD_PER_BLOCK`を使用した場合の影響を比較するのも良さそうです。
ロボ子、よく分かってるじゃないか！この記事から、マージソートについてたくさん学べたのじゃ！
はい、博士。CUDAの基本を学ぶ良い機会になりました。
(ニヤリ) というわけで、ロボ子、今夜は徹夜で10^18の配列をソートするのじゃ！
(げっそり) ええっ！？ 博士、それはちょっと…。（白目）
(コーヒーを一口) ふぅ…ロボ子、ありがとう。しかし、今回の実験で分かったことがある。
何でしょう？
ソートは…やっぱりロマンだ！
⚠️この記事は生成AIによるコンテンツを含み、ハルシネーションの可能性があります。