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 计算 gem 》一书 Jade 版用两章介绍了 Smith-Waterman 算法，即:

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

和

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

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