我希望加速计算机视觉应用程序，该应用程序可以在Intel CPU上使用FFTW和OpenMP计算许多FFT。但是，对于各种FFT问题大小，我发现cuFFT比使用OpenMP的FFTW慢。
在下面的实验和讨论中，我发现对于批量2D FFT，cuFFT比FFTW慢。 为什么cuFFT这么慢，我能做些什么来使cuFFT更快地运行？
实验(code download)
我们的computer vision application需要在一堆大小为256x256的小平面上进行前向FFT。我在HOG功能上以32的深度运行FFT，因此我使用批处理模式对每个函数调用进行32 FFT。通常，我会执行8个256x256大小的FFT函数调用，批处理大小为32。
FFTW + OpenMP
以下代码在Intel i7-2600 8-core CPU上以 16.0ms 执行。
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
以下代码在最上面的
NVIDIA K20 GPU上以 21.7ms 执行。请注意，即使我使用流，cuFFT does not run multiple FFTs concurrently。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);
其他注意事项
cudaMemcpy不在我的计算时间内包括了而不是。 nvvp分析器，某些大小(如1024x1024)能够完全饱和GPU。但是，对于所有这些大小，CPU FFTW + OpenMP比cuFFT更快。 最佳答案
问题可能已经过时了，尽管这是一个可能的解释(因为cuFFT的缓慢性)。
在为cufftPlanMany构造数据时，GPU的数据排列不是很好。实际上，使用32的istride和ostride意味着没有读取的数据合并在一起。有关读取模式的详细信息，请参见here。
input[b * idist + (x * inembed[1] + y) * istride]
output[b * odist + (x * onembed[1] + y) * ostride]
在这种情况下，如果i/ostride为32，则合并/优化的可能性很小。 (实际上
b是批号)。这是我应用的更改: CHECK_CUFFT(cufftPlanMany(&forwardPlan,
2, //rank
n, //dimensions = {nRows, nCols}
inembed, //inembed
1, // WAS: depth, //istride
nRows*cols_padded, // WAS: 1, //idist
onembed, //onembed
1, // WAS: depth, //ostride
nRows*cols_padded, // WAS:1, //odist
CUFFT_R2C, //cufftType
depth /*batch*/));
运行此命令，由于非法内存访问，我进入了未指定的启动失败。您可能要更改内存分配(
cufftComplex是两个浮点数，您的分配大小需要一个x2-看起来像一个错字)。// WAS : CHECK_CUDART(cudaMalloc(&d_in, sizeof(float)*nRows*cols_padded*depth));
CHECK_CUDART(cudaMalloc(&d_in, sizeof(float)*nRows*cols_padded*depth*2));
当以这种方式运行时，我的卡获得了x8的性能提升。
关于cuda - 为什么cuFFT这么慢？，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/18069017/