CUDA循环遍历下三角矩阵。

Question

如果有一个矩阵，我只想访问矩阵的下三角部分。我正试图找到一个好的线程索引，但到目前为止我还没有成功。有什么想法吗？ 我需要一个索引来循环遍历下三角矩阵，假设这是我的矩阵：

1 2 3 4
5 6 7 8
9 0 1 2
3 5 6 7

索引应该用于

1 
5 6
9 0 1
3 5 6 7

在这个例子中，1D数组的位置0、4、5、8、9、10、12、13、14、15。

CPU循环如下：

for(i = 0; i < N; i++){
    for(j = 0; j <= i; j++){
             .......

N是行数。我在内核中尝试了一些东西：

 __global__ void Kernel(int N) {

        int row = blockIdx.x * blockDim.x + threadIdx.x;
        int col = blockIdx.y * blockDim.y + threadIdx.y;
        if((row < N) && (col<=row) )
           printf("%d\n", row+col);
      }

然后这样调用：

 dim3 Blocks(1,1);
 dim3 Threads(N,N);
 Kernel<<< Blocks, Threads>>>(N);

但是它完全不起作用。 我得到的结果：

0
1
2
2
3
4

- Manolete

---

我认为条件应该是 if ((row<N) && (col<=row))。另外，您应该考虑增加网格大小。N 的值是多少？ - sgarizvi

你部分地说对了，但我认为整个想法都是错误的。 - Manolete

这个想法是正确的。你要启动多少个线程？如果每个块的线程数超过了设备限制，那么内核启动将会失败。 - sgarizvi

@sgar91 请查看我的最后一次编辑，那正是我现在正在执行的，但我并没有接近目标，可能是我在这里做错了什么。按照4x4矩阵的例子，我使用N=4个线程在每个维度上启动内核。 - Manolete

一定还有其他的问题，可能只是一些很小的错误，因为程序逻辑是正确的，但它只打印了前3行。 - sgarizvi

Answer 1

您正在启动线程网格，然后禁用对角线上方的所有线程，即约50%的线程将无事可做，非常低效。

您代码的简单修复方法是修复索引：

__global__ void Kernel(int N)
{
  int row = blockIdx.x * blockDim.x + threadIdx.x;
  int col = blockIdx.y * blockDim.y + threadIdx.y;
  if((row < N) && (col<=row) )
    printf("%d\n", row * N + col);
}

也许更高效但更复杂的解决方案是启动正确数量的线程并转换索引。请查看此答案以获取起点...

- Tom

Answer 2

问题在于我们正在索引一个一维数组，因此为了进行映射，我们需要将行索引乘以列数。因此，按照以下示例操作：

__global__ void Kernel(int N) {
        int row = blockIdx.x * blockDim.x + threadIdx.x;
        int col = blockIdx.y * blockDim.y + threadIdx.y;
        if((row < N) && (col<=row) )
           printf("%d\n", row*N + col);
 }

- Manolete

---

@harrism：我认为Manolete是在和我差不多的时间独立写下这个的。 - Tom