为什么cuFFT这么慢？

Question

我希望加速计算机视觉应用程序,在Intel CPU上使用FFTW和OpenMP计算许多FFT.但是,对于各种FFT问题大小,我发现cuFFT比使用OpenMP的FFTW慢.

在下面的实验和讨论中,我发现对于批量2D FFT ,cuFFT 比FFTW 慢.为什么cuFFT这么慢,有什么办法让cuFFT运行得更快？

---

实验(代码下载)

我们的计算机视觉应用需要在一堆256x256的小型平面上进行正向FFT.我在HOG功能上运行FFT,深度为32,因此我使用批处理模式为每个函数调用执行32次FFT.通常,我做大约8个大小为256x256的FFT函数调用,批量大小为32.

FFTW + OpenMP的
下面的代码执行在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.7ms内执行NVIDIA 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版本中,cudaMemcpyCPU和GPU之间的s 不包括在我的计算时间内.
• 这里给出的性能数是几个实验的平均值,其中每个实验有8个FFT函数调用(总共10个实验,所以80个FFT函数调用).
• 我尝试了很多问题大小(例如128x128,256x256,512x512,1024x1024),所有问题都是深度= 32.基于nvvp分析器,1024x1024等一些尺寸能够完全饱和GPU.但是,对于所有这些尺寸,CPU FFTW + OpenMP比cuFFT快.

Answer 1

问题可能已经过时,但这是一个可能的解释(因为cuFFT的缓慢).

在构建数据时cufftPlanMany,GPU的数据排列不是很好.实际上,使用32的istride和ostride意味着没有数据读取被合并.有关读取模式的详细信息,请参见此处

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.