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

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);
}

也许更有效但更复杂的解决方案是启动正确数量的线程并转换索引.看看这个答案的起点......

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);
 }