5

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

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

实验（代码下载）

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

FFTW + OpenMP
以下代码在.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的 top-of-the-line 上执行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 快。
4

1 回答 1

4

问题可能已经过时，尽管这是一个可能的解释（对于 cuFFT 的缓慢性）。

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

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 的提升。

于 2016-04-26T16:47:14.577 回答