CUDA：访问两个维度的任意长矩阵

Question

嘿， 目前我正在使用仅在一维索引的线程来访问矩阵的所有元素，如下所示：

// Thread-ID
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// Offset:
int offset = gridDim.x * blockDim.x;

while(idx < MATRIX_ROWS * MATRIX_COLS)
{   
    row = idx % MATRIX_ROWS;
    col = idx / MATRIX_ROWS;    

    matrix[ row ][ col ] = ...;
    idx += offset;
}

现在我想知道如何使用二维索引访问任意长矩阵。我希望一个块始终访问一行的单个元素。类似的东西（x-index 指的是矩阵的列，y-index 指的是矩阵的行）：

// Thread-IDs
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offset = gridDim.x * blockDim.x;

while(idx < MATRIX_COLS)
{   
    matrix[ idy ][ idx ] = ...;
    idx += offset;
}

现在让我们假设矩阵的行数比调用内核时开始块的行数多： 当启动 N 个块时，前 N 个块矩阵的行处理正确，但其他行呢？你会怎么做？ 谢谢！

编辑：我想出了一个想法，但我不知道这是否是某种“丑陋”的编码！？

// Thread-IDs
int idx0 = blockIdx.x * blockDim.x + threadIdx.x;
int idx = idx0;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offsetx = gridDim.x * blockDim.x;
int offsety = gridDim.y * blockDim.y;

while(idx < MATRIX_COLS && idy < MATRIX_ROWS)
{   
    matrix[ idy ][ idx ] = ...;

    idx += offsetx;
    if(idx > MATRIX_COLS)
    {
        // Jump to nex row and start from 'beginning' concerning columns
        idy += offsety;
        idx = idx0;
    }
}

Answer 1

也许这样的东西就是您正在寻找的？

// Thread-IDs
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offsetx = gridDim.x * blockDim.x;
int offsety = gridDim.y * blockDim.y;

for(row = idy; row < MATRIX_ROWS; i+=offsety) {
    float * pos = matrix + row;

#pragma unroll
    for(col = idx; col < MATRIX_COLS; col+=offsetx) {
        pos[col] = .....;
    }
}

如果 MATRIX_COLS 是预处理器定义或常量，编译器可能能够展开内部循环并提供更多性能。

编辑：代码的第一个版本是用列优先顺序编写的，这在我的脑海中留下了深刻的印象，因此请忽略此处的评论。应该以这种方式合并访问。