基于CUDA的分块/跨步归约实现优化技术问询

背景 #

我正在开发CUDA GPU代码，用于计算由下三角Toeplitz块组成的矩阵与向量的乘积。目前已完成基于FFT的块计算流程：

• 对矩阵块首列和向量块补零
• 执行批量FFT
• 点乘缩放
• 执行批量逆FFT

最终得到的结果数组布局为：[有效块0, 无效块0, 有效块1, 无效块1,...]

当前实现与性能情况 #

核心需求是对上述结果数组做分块归约，求和所有有效块得到最终输出。我已实现以下几种归约方案：

• 基于CUDA合作网格的线程块级二叉树归约
• 基于水平/垂直跨步的SVR2、SVR3内核

测试显示，double类型的SVR2内核性能最优，但我不确定这是否是当前场景下的最优方案。

另外我尝试过Thrust跨步迭代器方案，但由于需要对每个块元素调用thrust::reduce()（共O(10³)次调用），性能比SVR2至少慢2倍，已放弃该方案。

疑问与求助 #

1. 是否可以调整IFFT的输出布局，让后续的归约步骤更简单高效？
2. 希望针对当前的分块归约实现获得优化指导（附完整可编译测试代码及性能数据）

内容的提问来源于stack exchange，提问作者s769