CUDA 中 Smith-Waterman 算法的矩阵填充

Posted 2023-03-23

【中文标题】CUDA 中 Smith-Waterman 算法的矩阵填充

【英文标题】：Matrix filling for the Smith-Waterman algorithm in CUDA

【发布时间】：2014-05-28 16:40:24

【问题描述】：

我需要关于如何优化我在 CUDA 中实现 Smith-Waterman 算法的建议。

我要优化的部分是填充矩阵。由于矩阵元素之间的数据依赖性（每个下一个元素都依赖于其他元素 - 向左、向上和向左），我正在并行填充反对角矩阵元素，如图所示下图：

我的程序在循环中运行

int diag = 1;
for(int x = 0; x < size_b; x++)

    block_size = 1024;
    if(block_size > diag)
    
        block_size = diag;
    
    SAFE_KERNEL_CALL((dev_init_diag<<<(diag - 1)/block_size + 1, block_size>>>(H, size_a, size_b, x,
                      sequence_a, sequence_b, false, x_offset, y_offset, null_ind)));
    diag++;

如您所见，每个对角线都有一个内核调用。

因为我有相当大的矩阵（旁边有21000 元素），所以有很多内核调用。结果，我对CUDA内核调用的开销很大，浪费了大约一半的处理时间，这可以从Visual Profiler的截图中看出（看Profiler的开销字符串）：

所以，问题是如何摆脱多个内核调用并消除这种开销。

有一件重要的事情需要注意： 我为每个对角线调用一个新内核的原因是我需要在下一次调用之前同步线程和块，据我所知，同步 CUDA 块的唯一方法是完成内核并重新启动它。不过，对于这种算法，可能会有更好的解决方案。

感谢您阅读本文！

/////////////////////////////////////// //////////////

好的，谢谢您的回复！ 还有一个问题，更多关于 CUDA 的信息： 所以，我必须实现一个新的内核，大概是这样的：

__global__ void kernel(...)

for(int diag_num = 0; diag_num < size; diag_num++)

    init_one_diag(...);
    syncronize_threads();

但这意味着我只能在一个 cuda 块上启动这个内核？（因为我知道不同块之间没有同步）

在我以这种方式启动内核之前：

dev_init_diag>>(...)

新方法是否会同样有效？

【问题讨论】：

你可以试试动态并行？

【参考方案1】：

我建议您阅读现有文献，为 Smith-Waterman 算法实现矩阵填充问题的有效方法。

根据您的代码描述，您选择并行填充对角线，并且为每个对角线启动一个内核。正如您所提到的，由于多个内核启动，这是非常无效的。

一个简单的替代方法是构建一个负责计算所有反对角线的核函数。该内核应该以至少等于最长对角线的线程数启动。内核执行的迭代次数等于要计算的对角线的数量。对于比最长的对角线短的对角线，只有一部分线程保持活动状态。这种方法在

Parallelizing the Smith-Waterman Local Alignment Algorithm using CUDA

但由于两个原因无效：

对于大量计算（反对角线），大多数线程保持非活动状态； 内存访问高度未合并。

中的方法提供了对角矩阵填充的替代方法

Acceleration of the Smith–Waterman algorithm using single and multiple graphics processors

这里展示了如何重新制定 Smith-Waterman（对角线）矩阵填充算法，以便一次可以并行执行一行（或一列）计算。强调行（或列）计算如何允许 GPU 内存访问是连续的（合并的）并因此快速。虽然没有明确提及，但我相信这种方法也可以缓解（或完全消除）上述非活动线程的问题。

编辑

GPU Computing Gems 这本书，翡翠版，有两章专门介绍 Smith-Waterman 算法，即，

第 11 章，使用 Smith-Waterman 算法准确扫描序列数据库

和

第 13 章，GPU-超级计算机加速模式匹配

后者是第二个提到的方法的同一作者的一章。前者包含优化的 CUDA 代码的逐步推导，这可能对未来的用户有用。

【讨论】：

谢谢！然后，我还有一个问题，最后我通过编辑我的第一条消息发布了它。（因为这个cmets不支持代码粘贴）

@IlyaAfanasiev 关于你的第二个问题，我建议你设置一个在多个块上工作的代码。然后，您可以通过利用 CUDA 占用计算器或使用分析器（例如，Visual Profiler）来调整代码的性能，尤其是块大小。我忘了提到一个参考资料，您可以在其中找到史密斯-沃特曼算法的 CUDA 实现的逐步推导（请参阅我编辑的帖子）。虽然在将算法应用于许多序列的情况下它可能会发挥更大的作用，但它可能对你来说还是很有趣的。