Intel CPU で FFTW と OpenMP を使用して、多くの FFT を計算するコンピューター ビジョン アプリケーションを高速化することを望んでいます。ただし、さまざまな FFT の問題サイズについて、cuFFT は OpenMP を使用した FFTW よりも遅いことがわかりました。
以下の実験と議論では、cuFFT はバッチ化された 2D FFT の FFTWよりも遅いことがわかりました。cuFFT が非常に遅いのはなぜですか? また、cuFFT を高速化するためにできることはありますか?
実験 (コードのダウンロード)
私たちのコンピューター ビジョン アプリケーションでは、サイズ 256x256 の小さな平面の束で前方 FFT が必要です。深さ 32 のHOG機能でFFT を実行しているので、バッチ モードを使用して、関数呼び出しごとに 32 の FFT を実行します。通常、バッチ サイズ 32 でサイズ 256x256 の約 8 つの FFT 関数呼び出しを行います。
FFTW + OpenMP
次のコードは、 .NET で16.0msで実行されIntel i7-2600 8-core CPUます。
int depth = 32; int nRows = 256; int nCols = 256; int nIter = 8;
int n[2] = {nRows, nCols};
//if nCols is even, cols_padded = (nCols+2). if nCols is odd, cols_padded = (nCols+1)
int cols_padded = 2*(nCols/2 + 1); //allocate this width, but tell FFTW that it's nCols width
int inembed[2] = {nRows, 2*(nCols/2 + 1)};
int onembed[2] = {nRows, (nCols/2 + 1)}; //default -- equivalent ot onembed=NULL
float* h_in = (float*)malloc(sizeof(float)*nRows*cols_padded*depth);
memset(h_in, 0, sizeof(float)*nRows*cols_padded*depth);
fftwf_complex* h_freq = reinterpret_cast<fftwf_complex*>(h_in); //in-place version
fftwf_plan forwardPlan = fftwf_plan_many_dft_r2c(2, //rank
n, //dims -- this doesn't include zero-padding
depth, //howmany
h_in, //in
inembed, //inembed
depth, //istride
1, //idist
h_freq, //out
onembed, //onembed
depth, //ostride
1, //odist
FFTW_PATIENT /*flags*/);
double start = read_timer();
#pragma omp parallel for
for(int i=0; i<nIter; i++){
fftwf_execute_dft_r2c(forwardPlan, h_in, h_freq);
}
double responseTime = read_timer() - start;
printf("did %d FFT calls in %f ms \n", nIter, responseTime);
cuFFT
次のコードは、最上位で21.7msNVIDIA K20 GPUで実行されます。ストリームを使用しても、cuFFT は複数の FFT を同時に実行しないことに注意してください。
int depth = 32; int nRows = 256; int nCols = 256; int nIter = 8;
int n[2] = {nRows, nCols};
int cols_padded = 2*(nCols/2 + 1); //allocate this width, but tell FFTW that it's nCols width
int inembed[2] = {nRows, 2*(nCols/2 + 1)};
int onembed[2] = {nRows, (nCols/2 + 1)}; //default -- equivalent ot onembed=NULL in FFTW
cufftHandle forwardPlan;
float* d_in; cufftComplex* d_freq;
CHECK_CUFFT(cufftPlanMany(&forwardPlan,
2, //rank
n, //dimensions = {nRows, nCols}
inembed, //inembed
depth, //istride
1, //idist
onembed, //onembed
depth, //ostride
1, //odist
CUFFT_R2C, //cufftType
depth /*batch*/));
CHECK_CUDART(cudaMalloc(&d_in, sizeof(float)*nRows*cols_padded*depth));
d_freq = reinterpret_cast<cufftComplex*>(d_in);
double start = read_timer();
for(int i=0; i<nIter; i++){
CHECK_CUFFT(cufftExecR2C(forwardPlan, d_in, d_freq));
}
CHECK_CUDART(cudaDeviceSynchronize());
double responseTime = read_timer() - start;
printf("did %d FFT calls in %f ms \n", nIter, responseTime);
その他の注意事項
- GPU 版で
cudaMemcpyは、CPU と GPU 間の s は計算時間に含まれません。 - ここに示すパフォーマンスの数値は、いくつかの実験の平均であり、各実験には 8 つの FFT 関数呼び出しがあります (合計 10 回の実験、つまり 80 個の FFT 関数呼び出し)。
- 多くの問題サイズ (例: 128x128、256x256、512x512、1024x1024) をすべて深さ = 32 で試しました。プロファイラーに基づくと、
nvvp1024x1024 などの一部のサイズは GPU を完全に飽和させることができます。しかし、これらすべてのサイズで、CPU FFTW+OpenMP は cuFFT より高速です。