为什么cuFFT性能会受到重叠输入的影响？

Question

我正在尝试使用cuFFT的回调函数来快速执行输入格式转换（例如，计算8位整数输入数据的FFT，而无需首先将输入缓冲区显式转换为float）。在我的许多应用程序中，我需要计算重叠输入缓冲区上的FFT，as described in this previous SO question。通常，相邻的FFT可能会重叠1/4到1/8的FFT长度。

cuFFT及其类似FFTW的接口明确支持此via the idist parameter of the cufftPlanMany() function。具体来说，如果我想计算连续输入之间4096个采样重叠的32768大小的FFT，我会设置idist = 32768 - 4096。这确实工作正常，它会产生正确的输出。

但是，以这种方式使用cuFFT时，我看到了奇怪的性能下降。我已经发明了实现两种不同的方式这种格式转换和重叠的测试：

1. 明确告知CUFFT关于输入的重叠性质：设置idist = nfft - overlap如我如上所述。安装一个负载回调函数，只需根据提供给回调的缓冲区索引从int8_t到float进行转换即可。

2. 不要告诉cuFFT关于输入的重叠性质;谎言它dset idist = nfft。然后，让回调函数通过计算每个FFT输入应读取的正确索引来处理重叠。

A test program implementing both of these approaches with timing and equivalence tests is available in this GitHub gist。为了简洁起见，我没有在这里重现这一切。该程序计算一批1024个32768点的FFT，其重叠4096个样本;输入数据类型是8位整数。当我在我的机器上运行（与GeForce GTX的660 GPU，使用CUDA 8.0 RC在Ubuntu 16.04），我得到以下结果：

executing method 1...done in 32.523 msec 
executing method 2...done in 26.3281 msec 

方法2是相当快的，这是我不希望。看的回调函数的实现：

方法1：

template <typename T> 
__device__ cufftReal convert_callback(void * inbuf, size_t fft_index, 
    void *, void *) 
{ 
    return (cufftReal)(((const T *) inbuf)[fft_index]); 
} 

方法2：

template <typename T> 
__device__ cufftReal convert_and_overlap_callback(void *inbuf, 
    size_t fft_index, void *, void *) 
{ 
    // fft_index is the index of the sample that we need, not taking 
    // the overlap into account. Convert it to the appropriate sample 
    // index, considering the overlap structure. First, grab the FFT 
    // parameters from constant memory. 
    int nfft = overlap_params.nfft; 
    int overlap = overlap_params.overlap; 
    // Calculate which FFT in the batch that we're reading data for. This 
    // tells us how much overlap we need to account for. Just use integer 
    // arithmetic here for speed, knowing that this would cause a problem 
    // if we did a batch larger than 2Gsamples long. 
    int fft_index_int = fft_index; 
    int fft_batch_index = fft_index_int/nfft; 
    // For each transform past the first one, we need to slide "overlap" 
    // samples back in the input buffer when fetching the sample. 
    fft_index_int -= fft_batch_index * overlap; 
    // Cast the input pointer to the appropriate type and convert to a float. 
    return (cufftReal) (((const T *) inbuf)[fft_index_int]); 
} 

方法2具有显著更复杂的回调函数，一个甚至包含一个非编译时间值的整数除法！我希望这比方法1慢得多，但我看到的是相反的情况。对此有没有很好的解释？当输入重叠时，cuFFT可能会以不同的方式处理它的处理，从而导致性能下降？

看起来像我如果可以从回调中删除索引计算（但需要将重叠指定为cuFFT），那么应该能够实现比方法2快得多的性能。

编辑：在运行我的测试程序nvvp后，我可以看到cuFFT绝对似乎正在以不同方式构造其计算。很难使内核符号名称的意义，但内核调用打破这样的：

方法1：

1. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex14packR2C_kernelIjfEEvNS_19spRealComplexR2C_stIT_T0_EE：3.72毫秒
2. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK>：7.71毫秒
3. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK> ：12.75毫秒（是，它被调用两次）
4. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelTexIjfL9fftAxii_t1EEEvP7ComplexIT0_EjT_15coordDivisors_tIS6_E7coord_tIS6_ESA_S6_S3_：7.49毫秒

方法2：

1. spRadix0128C::kernel1MemCallback<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, L1, ALL, WRITEBACK>：5.15毫秒
2. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK>：12.88毫秒
3. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelTexIjfL9fftAxii_t1EEEvP7ComplexIT0_EjT_15coordDivisors_tIS6_E7coord_tIS6_ESA_S6_S3_：7.51毫秒

有趣的是，它看起来像CUFFT调用两个内核来实际使用方法1计算FFT（当cuFFT知道重叠时），但使用方法2（不知道t的方法）这些FFT是重叠的），它只用一个来完成这项工作。对于在这两种情况下使用的内核，似乎在方法1和方法2之间使用相同的网格参数。

我不明白为什么它应该在这里使用不同的实现，尤其是在输入步幅istride == 1。在变换输入端获取数据时，它应该只使用不同的基地址;我认为算法的其余部分应该完全相同。

编辑2：我看到一些更离奇的行为。我意外地意识到，如果我没有合适地销毁cuFFT手柄，我会发现测量性能存在差异。例如，我修改了测试程序以跳过cuFFT手柄的销毁，然后以不同的顺序执行测试：方法1，方法2，然后方法2和方法1。我得到了以下结果：

executing method 1...done in 31.5662 msec 
executing method 2...done in 17.6484 msec 
executing method 2...done in 17.7506 msec 
executing method 1...done in 20.2447 msec 

这样的表现似乎取决于是否有存在其他CUFFT计划创建测试用例计划时改变！使用Profiler，我发现内核启动的结构在两种情况下不会改变;内核似乎都执行得更快。我对这种效果也没有合理的解释。

Answer 1

在@llukas的建议下，我向NVIDIA提交了关于该问题的错误报告（https://partners.nvidia.com/bug/viewbug/1821802，如果您已注册为开发人员）。他们承认重叠计划的表现较差。他们实际上表示，在这两种情况下使用的内核配置都不是最理想的，他们计划最终改进它。没有给出ETA，但它可能会在而不是处于下一个版本（8.0刚刚在上周发布）。最后，他们表示，从CUDA 8.0开始，没有任何解决方法可以使cuFFT使用更有效的方法来处理输入。

Answer 2

如果您指定非标准步幅（批次/转换无关紧要）cuFFT在内部使用不同的路径。

ad edit 2： 这可能是GPU Boost调整GPU上的时钟。CUFFT计划不会对其他

方式影响一个获得更稳定的结果：

1. 运行热身内核（任何会充分GPU做工还是不错的），然后你的问题
2. 增加批量大小
3. 运行测试几次，并利用GPU的平均
4. 锁定时钟（在GeForce不是真的有可能 - 特斯拉能做到这一点）