cuda - 2 つの GPU で同時に cudaMalloc を呼び出すとパフォーマンスが低下する

Question

ユーザーのシステム上の GPU 間で処理負荷を分割するアプリケーションがあります。基本的に、メイン アプリケーション スレッドによって定期的にトリガー されると、 GPU 処理間隔を開始する GPU ごとに CPU スレッドがあります。

GPU 処理間隔の例として、次の画像 (NVIDIA の CUDA プロファイラー ツールを使用して生成) を検討してください。ここでは、アプリケーションは単一の GPU を使用しています。

ご覧のとおり、GPU 処理時間の大部分は 2 つの並べ替え操作によって消費されます。私はこれに Thrust ライブラリ (thrust::sort_by_key) を使用しています。また、実際の並べ替えを開始する前に、thrust::sort_by_key がフードの下でいくつかの cudaMallocs を呼び出しているようです。

次に、アプリケーションが処理負荷を 2 つの GPU に分散させた同じ処理間隔を考えます。

完璧な世界では、2 つの GPU の処理間隔が 1 つの GPU の処理間隔の正確に半分になることが予想されます (各 GPU が半分の作業を行っているため)。ご覧のように、cudaMalloc が同時に呼び出されると、ある種の競合の問題により時間がかかるように見えるため (2 ～ 3 倍長くなる場合もあります)、部分的にはそうではありません。2 つの GPU のメモリ割り当てスペースは完全に独立しているため、cudaMalloc でシステム全体のロックを行うべきではないため、これが必要な理由がわかりません。GPU ごとのロックの方が合理的です。

cudaMalloc の同時呼び出しに問題があるという私の仮説を証明するために、2 つの CPU スレッド (GPU ごとに) がそれぞれ cudaMalloc を数回呼び出す非常に単純なプログラムを作成しました。最初にこのプログラムを実行して、別々のスレッドが cudaMalloc を同時に呼び出さないようにしました。

割り当てごとに最大 175 マイクロ秒かかることがわかります。次に、cudaMalloc を同時に呼び出すスレッドでプログラムを実行しました。

ここでは、各呼び出しに約 538 マイクロ秒、つまり前のケースの 3 倍の時間がかかりました! 言うまでもなく、これはアプリケーションの速度を大幅に低下させており、2 つ以上の GPU で問題が悪化するのは当然のことです。

Linux と Windows でこの動作に気付きました。Linux では Nvidia ドライバー バージョン 319.60 を使用し、Windows では 327.23 バージョンを使用しています。CUDA ツールキット 5.5 を使用しています。

考えられる理由: これらのテストでは GTX 690 を使用しています。このカードは基本的に、同じユニットに収容された 2 つの 680 のような GPU です。これは私が実行した唯一の「マルチ GPU」設定なので、cudaMalloc の問題は 690 の 2 つの GPU 間のハードウェア依存関係と関係があるのでしょうか?

Answer 1

私は NVIDIA ドライバーの内部に精通していないので、これはやや憶測です。

表示されている速度低下は、デバイスの malloc を同時に呼び出す複数のスレッドからの競合によって引き起こされるドライバー レベルの競合です。デバイス メモリの割り当てには、ドライバー レベルのコンテキスト切り替えと同様に、多数の OS システム コールが必要です。どちらの操作にも、かなりの量のレイテンシーがあります。2 つのスレッドが同時にメモリを割り当てようとするときに表示される余分な時間は、両方のデバイスにメモリを割り当てるために必要なシステム コールのシーケンス全体で、あるデバイスから別のデバイスに切り替えることによる追加のドライバー レイテンシが原因である可能性があります。

これを軽減できるいくつかの方法を考えることができます。

• 初期化中に割り当てられたメモリのスラブで動作するデバイス用の独自のカスタム スラスト メモリ アロケータを作成することで、スラスト メモリ割り当てのシステム コール オーバーヘッドをゼロに減らすことができます。これにより、各 内のシステム コールのオーバーヘッドがすべて取り除かれますsort_by_keyが、独自のユーザー メモリ マネージャーを作成する作業は簡単ではありません。一方、スラスト コードの残りの部分はそのまま残ります。
• 別のソートライブラリに切り替えて、一時メモリの割り当てを自分で管理することができます。初期化フェーズですべての割り当てを行う場合、1 回限りのメモリ割り当てのコストは、各スレッドの存続期間にわたってほぼゼロになるまで償却できます。

私が書いたマルチ GPU CUBLAS ベースの線形代数コードでは、両方のアイデアを組み合わせて、一度だけ割り当てられたデバイス メモリ プールで動作するスタンドアロンのユーザー空間デバイス メモリ マネージャーを作成しました。中間デバイス メモリ割り当てのオーバーヘッド コストをすべて削除すると、速度が大幅に向上することがわかりました。あなたのユースケースは、同様の戦略から利益を得るかもしれません。