2

そのため、CUDA ストリーム機能と組み合わせて cuFFT を使用しています。私が抱えている問題は、cuFFT カーネルを完全な同時実行で実行できないように見えることです。以下は、nvvp から得た結果です。各ストリームは、サイズ 128x128 の 128 個の画像に対して 2D バッチ FFT のカーネルを実行しています。3 つのストリームをセットアップして、3 つの独立した FFT バッチ プランを実行します。

ここに画像の説明を入力

図からわかるように、一部のメモリ コピー (黄色のバー) が一部のカーネル計算 (紫、茶色、ピンクのバー) と並行していました。しかし、カーネルの実行はまったく並行していませんでした。お気づきのように、各カーネルは厳密に相互に追従していました。以下は、デバイスへのメモリ コピーとカーネルの起動に使用したコードです。

    for (unsigned int j = 0; j < NUM_IMAGES; j++ ) {
        gpuErrchk( cudaMemcpyAsync( dev_pointers_in[j],
                                image_vector[j],
                                NX*NY*NZ*sizeof(SimPixelType),
                                cudaMemcpyHostToDevice,
                                streams_fft[j]) );
        gpuErrchk( cudaMemcpyAsync( dev_pointers_out[j],
                                out,
                                NX*NY*NZ*sizeof(cufftDoubleComplex),
                                cudaMemcpyHostToDevice,
                                streams_fft[j] ) );
        cufftExecD2Z( planr2c[j],
                  (SimPixelType*)dev_pointers_in[j],
                  (cufftDoubleComplex*)dev_pointers_out[j]);

    }

次に、すべてのメモリ コピー (同期) を完了し、すべてのカーネルを一度にストリームに送信するようにコードを変更したところ、次のプロファイリング結果が得られました。

ここに画像の説明を入力

その後、カーネルが同時に実行されていないことが確認されました。

「–default-stream per-thread」コマンドライン引数を渡すか、#include またはコードに #define CUDA_API_PER_THREAD_DEFAULT_STREAM を渡すことにより、完全な同時実行性を利用するようにセットアップする方法を詳細に説明する1 つのリンクを見ました。これは CUDA 7 で導入された機能です。GeForce GT750M を搭載した MacBook Pro Retina 15' (上記のリンクと同じマシン) で上記のリンクのサンプル コードを実行したところ、カーネルを同時に実行できました。しかし、cuFFT カーネルを並行して実行することはできませんでした。

次に、cuFFTカーネルがGPU全体を占有するため、2つのcuFFTカーネルが並行して実行されることはないと誰かが言っているこのリンクを見つけました。それから私は立ち往生しました。CUFFT が並行カーネルを有効にするかどうかを説明する正式なドキュメントが見つからなかったためです。これは本当ですか？これを回避する方法はありますか？

4

1 に答える 1

2

各プランが個別のストリームに関連付けられるcufftSetStream()ように、表示したコードの前に、それぞれに適切に電話したと思います。planr2c[j]あなたが投稿したコードには表示されません。実際にカフト カーネルを他のカフト カーネルとオーバーラップさせたい場合は、それらのカーネルを別のストリームに起動する必要があります。したがって、たとえば、イメージ 0 の cufft exec 呼び出しは、イメージ 1 の cufft exec 呼び出しとは異なるストリームに起動する必要があります。

2 つの CUDA 操作がオーバーラップできるようにするには、それらを異なるストリームで起動する必要があります。

そうは言っても、カーネル実行による同時メモリ コピーは、妥当なサイズの FFT に期待されるものです。

128x128 FFT を一次近似にすると、最大 15,000 スレッドがスピンアップするため、スレッド ブロックがそれぞれ最大 500 スレッドの場合、それは 30 スレッド ブロックになり、GPU がかなり占有されたままになり、カーネルを追加するための「余地」があまり残りません。(プロファイラー自体でカーネルの合計ブロックとスレッドを実際に検出できます。) GT750mにはおそらく 2 つの Kepler SMがあり、SM ごとに最大 16 ブロックであるため、最大瞬間容量は 32 ブロックです。また、この容量の数値は、共有メモリの使用、レジスタの使用、またはその他の要因により、特定のカーネルで減少する可能性があります。

実行している GPU の瞬間的な容量 (SM あたりの最大ブロック * SM の数) によって、カーネルのオーバーラップ (同時実行) の可能性が決まります。1 回のカーネル起動でその容量を超えると、GPU が「いっぱい」になり、しばらくの間カーネルの同時実行が妨げられます。

理論的には、CUFFT カーネルを同時に実行することは可能です。ただし、CUFFT などのカーネルの同時実行シナリオと同様に、実際に同時実行を確認するには、これらのカーネルのリソース使用量をかなり低くする必要があります。通常、リソースの使用率が低い場合は、スレッド/スレッドブロックの数が比較的少ないカーネルを意味します。これらのカーネルは通常、実行に時間がかからないため、同時実行を実際に確認することはさらに困難になります (起動の遅延やその他の遅延要因が邪魔になる可能性があるため)。並行カーネルを確認する最も簡単な方法は、リソース要件が異常に低く、ランタイムが異常に長いカーネルを使用することです。これは通常、CUFFT カーネルまたはその他のカーネルの典型的なシナリオではありません。

コピーと計算のオーバーラップは、CUFFT を使用したスト​​リームの便利な機能です。また、マシンの容量とリソースの制約を理解する根拠がない同時実行の考え方は、それ自体がいくぶん不合理です。たとえば、カーネルの同時実行性が任意に達成可能 (「任意の 2 つのカーネルを同時に実行できるはず」) であり、容量やリソースの詳細を考慮しない場合、2 つのカーネルを同時に実行した後、次の論理的なステップは次のようになります。 4、8、16 カーネルに同時に移動します。しかし現実には、機械はそれほど多くの作業を同時に処理することはできません。1 回のカーネル起動で十分な並列処理 (「十分なスレッド」と大まかに翻訳) を公開したら、

于 2016-04-16T01:59:02.650 に答える