CUDA信道化器实现：CDP-FFT调用与同步问题求助

1. 放弃CDP，采用多流串行+并行调度 #

既然协作网格同步与CDP无法兼容，且CDP本身存在同步风险，建议改用主机端多流调度方案：

• 为每个信道分配独立的CUDA流，在每个流中依次执行「滤波内核 → cuFFT操作」
• 利用cudaStreamCreate创建流，将滤波内核和cuFFT任务都绑定到对应流中，流内任务天然串行，保证滤波完成后才执行FFT
• 主机端通过cudaDeviceSynchronize等待所有流完成，或用事件同步实现更精细的控制
• 该方案完全符合“滤波后立即FFT”的概念设计，且多流并行能充分利用GPU算力，同时避免了CDP的同步问题

2. 优化CDP模式下的同步机制（若坚持使用CDP） #

如果必须在核内调用FFT，可通过全局原子计数器+忙等实现块间同步：

• 初始化一个全局原子变量filter_done_count，初始值为0
• 每个块完成滤波后，调用atomicAdd(&filter_done_count, 1)，然后循环检查filter_done_count是否等于总块数
• 当所有块都完成滤波后，再调用FFT内核；注意在忙等过程中可插入__yield()减少资源占用
• 针对零值问题，在滤波块内所有线程完成写操作后，添加__syncthreads()和__threadfence()，确保全局内存数据对所有块可见

3. 架构针对性优化 #

• Blackwell GPU：利用Warp Specialization加速多相滤波的归约操作，将归约任务分配给专用warp，提升计算效率
• Ada Lovelace GPU：后续移植时，启用cuFFT的Tensor Core加速（需配置合适的FFT大小和数据类型），同时利用异步拷贝进一步隐藏延迟
• 内存优化：使用共享内存缓存滤波的输入数据，减少全局内存访问次数；保证全局内存访问对齐，提升带宽利用率

问题背景 #

从概念上讲，信道化器由一系列归约操作（多相滤波器）和数量相等的变换（FFT）组成。

已尝试两种方案均存在问题：

• 方案1：每个CUDA块执行滤波后通过CDP调用FFT，无块间同步，偶尔出现FFT输出零值，推测是全局内存读写未跨线程同步导致
• 方案2：分离滤波与FFT，先完成所有滤波再执行FFT，但无法在信道化器内核内调用FFT（因协作组与CDP不兼容）

环境：CUDA Toolkit 13.1.115、Nvidia RTX Pro 4000 Blackwell GPU（后续移植至Ada Lovelace）、Ubuntu 24.04

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