将 3D 网格转换为 2D 数组索引

Question

假设我想将以下 C 例程翻译成 CUDA 内核。

而且，我想使用网格中的所有维度来运行内核。

如何计算矩阵的行和列的索引？

void OuterProduct(float* A, float* B, float** C, int N)
{
    for(int r=0 ; r<N ; r++)
    {
        for(int c=0 ; c<N ; c++)
        {
            for(int cc=0 ; cc<N ; cc++)
            {
                (*C)[r * N + c] += A[r * N + cc] * B[cc * N + c];
            }
        }
    }
}

以下是我的理解：

假设上面的例程是要将两个 3x3 矩阵相乘。因此，计算次数为 3x3x3 = 27。因此，我们需要 27 个线程来完成乘法运算。

假设我们将在每个块中使用一个线程。所以，我们需要 27 个区块。

dim3 threads_per_block(3, 3, 3);
dim3 blocks_per_grid(3, 3, 3);

typedef float I;  
__global__ void OuterProductKernel(I* A, I* B, I* C, int N)
{
    int dimx = N;
    int dimy = N;
    int dimz = N;

    int r = blockIdx.x * blockDim.x + threadIdx.x;
    int c = blockIdx.y * blockDim.y + threadIdx.y;
    int d = blockIdx.z * blockDim.z + threadIdx.z;

    if (r < N && c < N && d < N) 
    {
        int loc_c = d * dimx * dimy + c * dimx + r;
 
        for (int cc=0; cc<N; cc++) 
        {
            int loc_a = (cc * dimx * dimy) + (c * dimx) + r;
            int loc_b = (d * dimx * dimy) + (cc * dimx) + r;
                    C[loc_c] += A[loc_a]*B[loc_b];
        }
    }
}

这是正确的吗？我认为不是。

你能告诉我计算

loc_a

、

loc_b

和

loc_c

的正确理由吗？

Answer 1

CUDA（作为 C 和 C++）使用 行优先顺序，所以代码像

int loc_c = d * dimx * dimy + c * dimx + r;

应该改写为

int loc_c = d * dimx * dimy + r * dimx + c;

与其他“位置”相同：loc_a 和 loc_b。

还有：

1. 确保 C 数组归零，你永远不会在代码中这样做
2. 如果能看到调用代码就好了