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

Question

我需要有关如何在 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 开销字符串):

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

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

感谢您阅读本文!

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

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

__global__ void kernel(...)
{
for(int diag_num = 0; diag_num < size; diag_num++)
{
    init_one_diag(...);
    syncronize_threads();
}
}

但这意味着我必须只在一个 cuda block 上启动这个内核？(因为据我所知，不同 block 之间没有同步)

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

dev_init_diag<<<(diag - 1)/block_size + 1, block_size>>>(...)

新方法是否有效？

Answer 1

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

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

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

Parallelizing the Smith-Waterman Local Alignment Algorithm using CUDA

但由于两个原因无效:

1. 大多数线程在大量计算(反对角线)中保持不活动状态；
2. 内存访问高度未合并。

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

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

其中展示了如何重新制定 Smith–Waterman(反对角线)矩阵填充算法，以便可以一次并行执行一行(或一列)计算。强调了行(或列)计算如何使 GPU 内存访问连续(合并)并因此变得快速。尽管没有明确提及，但我相信这种方法也可以减轻(或完全消除)上述非事件线程问题。

编辑

GPU Computing Gems 书籍， Jade 版，专门用两章介绍 Smith-Waterman 算法，即

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

和

第 13 章，模式匹配的 GPU- super 计算机加速

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