algorithm - 如何提高CUDA中needleman-wunsch算法的性能

标签 algorithm cuda gpu

我需要有关如何在 CUDA 中优化我的 Needleman-Wunsch 算法实现的建议。

我想优化我的代码以填充 CUDA 中的 DP 矩阵。由于矩阵元素之间的数据依赖性(每个下一个元素都依赖于其他元素 - 左到它,到它,再到它),我正在并行填充反对角矩阵元素,如下所示:

__global__ void alignment_kernel(int *T, char *A, char *B, int t_M, int t_N, int d) {
  int row = BLOCK_SIZE_Y * blockIdx.y + threadIdx.y;
  int col = BLOCK_SIZE_X * blockIdx.x + threadIdx.x;

  // Check if we are inside the table boundaries.
  if (!(row < t_M && col < t_N)) {
    return;
  }

  // Check if current thread is on the current diagonal
  if (row + col != d) {
    return;
  }

  int v1;
  int v2;
  int v3;
  int v4;
  v1 = v2 = v3 = v4 = INT_MIN;

  if (row > 0 && col > 0) {
    v1 = T[t_N * (row - 1) + (col - 1)] + score_matrix_read(A[row - 1], B[col - 1]);
  }
  if (row > 0 && col >= 0) {
    v2 = T[t_N * (row - 1) + col] + gap;
  }
  if (row >= 0 && col > 0) {
    v3 = T[t_N * row + (col - 1)] + gap;
  } 
  if (row == 0 && col == 0) {
    v4 = 0;
  }

  // Synchronize (ensure all the data is available)
  __syncthreads();

  T[t_N * row + col] = mmax(v1, v2, v3, v4);
}

然而,我的代码的一个明显问题是我进行了多次内核调用(代码如下)。到现在为止,我不知道如何使用线程来同步处理反对角线而不这样做。我认为这是达到更好性能的主要问题。

// Invoke kernel.
  for (int d = 0; d < t_M + t_N - 1; d++) {
    alignment_kernel<<< gridDim, blockDim >>>(d_T, d_A, d_B, t_M, t_N, d);
    //CHECK_FOR_CUDA_ERROR();
  }

我怎样才能并行处理反对角线,并可能使用共享内存来提高加速比?

除了这个问题,有没有办法并行执行 needleman-wunsch 算法的回溯步骤?

最佳答案

我目前也在研究 Needleman Wunsch 算法的并行实现(用于基因组作图器)。根据您要进行的对齐次数,每个线程进行一次对齐可能更有效。

然而,这里有一个publication并行执行单个对齐(在 GPU 上)。他们方法的新颖之处在于它不是按顺序生成矩阵,而是按对角线生成矩阵。他们不谈论他们如何在他们的出版物中回溯。他们在生成矩阵后将其发送回主机,然后他们使用 CPU 执行回溯。我认为 GPU 上的回溯会因为分支而变得非常低效。

关于algorithm - 如何提高CUDA中needleman-wunsch算法的性能,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/29829245/

相关文章:

algorithm - 每个数字都可以加减的子集求和问题

c++ - 如何从 2D 纹理中成功读取

CUDA fft - Cooley tukey,如何利用并行性?

c++ - 在传输到 CUDA GPU 时保持主机数据完好无损

python - 仅针对选定的GPU架构编译Cupy

algorithm - 从给定区域采样

mysql - 如何为数据库中的每篇文章制作一个简短的唯一 ID?

Java - 检查数组是否已排序后代

opengl - 使用 OpenGL 进行高效的 GPU 随机内存访问

linux - pci_lookup_name 怎么来的?