cuFFTストリームの同時実行

Question

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

、いくつかのメモリ・コピー（黄色のバー）は、いくつかのカーネル演算（、紫褐色、ピンクバー）と並行していました。しかし、カーネル実行はまったく同時ではありませんでした。あなたが気づくごとに、各カーネルは互いに厳密に従っていました。以下は、デバイスへのメモリコピーとカーネルの起動に使用したコードです。そして、私がいた

：私は（同期）すべてのメモリコピーを終えて、一度にストリームにすべてのカーネルを送信し、私は、次のプロファイリング結果を得たように

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]); 

    } 

は、その後、私は私のコードを変更しましたカーネルが並行して実行されていないことを確認しました。

私はあなたが#includeしたり、あなたのコード内で前に設定するか渡す「-default-ストリームスレッドごとの」コマンドライン引数またはの#define CUDA_API_PER_THREAD_DEFAULT_STREAMでフル並行性を活用する方法を詳細に説明している1 linkを見ました。これはCUDA 7で導入された機能です。上記のリンクにあるMacBook Pro Retina 15 'のサンプルコードをGeForce GT750M（上記のリンクと同じマシン）で実行し、同時にカーネルを実行することができました。しかし、私はcuFFTカーネルを並行して走らせることができませんでした。

私はこれを見つけました。誰かがcuFFTカーネルがGPU全体を占有していると言って、2つのcuFFTカーネルは並行して走っていません。それから私は立ち往生した。私は、CUFFTが並行カーネルを有効にするかどうかを扱う正式な文書は見つけていないので、これは本当ですか？これを回避する方法はありますか？

Answer 1

あなたが表示したコードの前にcufftSetStream()と呼ばれ、各プランが別々のストリームに関連付けられるように、それぞれplanr2c[j]に適しているとします。あなたが投稿したコードには表示されません。カフのカーネルと他のカフのカーネルが重なり合うようにするには、のカーネルを起動して別々のストリームにする必要があります。したがって、画像0のためのcufft exec呼び出しは、例えば画像1のcufft exec呼び出しとは異なるストリームに起動されなければならない。ために

ための任意の二つのCUDA動作が重複する可能性を持っている、彼らは、異なるストリームに起動されなければなりません。

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

一次近似の128x128 FFTは〜15,000スレッドをスピンアップするので、もしスレッドブロックがそれぞれ〜500スレッドであれば、それは30スレッドブロックになり、GPUはかなり占有され、追加のカーネル用。 GT750m probably has 2 Kepler SMsをa maximum of 16 blocks per SMとすると、最大瞬時容量は32ブロックになるので、実際にはプロファイラ自体でカーネルのブロックとスレッドの総数を知ることができます。そして、この容量数は、共有メモリ使用量、レジスタ使用量、または他の要因のために特定のカーネルのために減少される可能性がある。

あなたが上で実行されているものは何でもGPUの瞬間容量（SMあたりの最大ブロックは*のSMの数）は、カーネルの重複（同時性）の可能性を決定します。カーネルを1回起動するだけでその容量を超えると、GPUが「満たされ」、カーネルの同時実行が妨げられることがあります。

CUFFTカーネルを同時に実行することは理論的に可能でなければなりません。しかし、カーネル並行処理のシナリオ（CUFFTなど）と同じように、カーネルのリソース使用率は、並行性を実際に見るにはかなり低くなければなりません。通常、リソース使用率が低い場合は、スレッド/スレッドブロックの数が比較的少ないカーネルを意味します。これらのカーネルは通常、実行に時間がかかりません。実際には同時実行性を目撃するのがさらに難しくなります（起動遅延やその他の遅延要因が発生する可能性があるため）。並行カーネルを目撃する最も簡単な方法は、異常に低いリソース要件と異常に長い実行時間を組み合わせたカーネルを持つことです。 CUFFTカーネルやその他のカーネルでは、これは一般的なシナリオではありません。コピーおよび計算の

オーバーラップはCUFFTとストリームのさらに便利な機能です。並行性の考え方は、マシンの容量とリソースの制約を理解する根拠がなく、それ自体はやや不合理です。カーネルの同時実行が（「私は2つのカーネルが同時に実行させることができるはず」）任意に達成した場合は、同時に実行している2つのカーネルを取得した後、例えば、容量やリソース仕様を考慮せずに、その後、次の論理的なステップは、になります同時に4,8,16のカーネルに行きます。しかし現実は、マシンが同時に多くの作業を処理することができないということです。単一のカーネルの起動時に十分な並列性（十分に「十分なスレッド」と解釈される）を覚えたら、追加のカーネル起動による追加の並列処理を公開すると、通常はマシンが速く実行されたり、作業が迅速に処理されたりすることはありません。