Cuda矩阵乘法给出错误答案

Question

更新！

我当前的代码不检查越界内存访问。当我运行 cuda memcheck 时，它说即使对于只有 2 x 2 的矩阵，内存访问也很糟糕！我正在访问我不应该以某种方式访问​​的内存，这就是问题所在！

要检查内存访问是否越界，请运行 cuda-memcheck ./(在此处插入可执行文件)

下面显示的是我的矩阵乘法代码：

dim3 block(32,32);
dim3 grid( (n+31)/32, (n+31)/32 );
matrixMul<<<grid,block>>>(d_C, d_A, d_B, n, k);

kA 和 kB 是其中包含值的矩阵（为了更容易，它们都是 2）。

m, n, k 对于我的方阵来说都是同一个数

kC 是存储答案的矩阵。

#ifndef _MATRIXMUL_KERNEL_H_
#define _MATRIXMUL_KERNEL_H_

#include <stdio.h>

__global__ void matrixMul(float *kC, float *kA, float *kB, int n, int k)
{

    int tx = blockIdx.x * 32 + threadIdx.x;
    int ty = blockIdx.y * 32 + threadIdx.y;
    float value = 0;

    for (int i=0;i<n;i++)
    {
        float elementA=kA[ty*n+i];
        float elementB=kB[i*k+tx];
        value += elementA*elementB;
    }

    kC[ty*n+tx] = value;
}

#endif // #ifndef _MATRIXMUL_KERNEL_H_

Answer 1

根据您定义线程网格的方式，您应该像这样在内核代码中添加线程检查：

#ifndef _MATRIXMUL_KERNEL_H_
#define _MATRIXMUL_KERNEL_H_

#include <stdio.h>

__global__ void matrixMul(float *kC, float *kA, float *kB, int n, int k)
{

    int tx = blockIdx.x * 32 + threadIdx.x;
    int ty = blockIdx.y * 32 + threadIdx.y;

    if ((ty < n) && (tx < n)) { // add this line
      float value = 0;

      for (int i=0;i<n;i++)
      {
        float elementA=kA[ty*n+i];
        float elementB=kB[i*k+tx];
        value += elementA*elementB;
      }

      kC[ty*n+tx] = value;
    }  //  add this line
}

#endif // #ifndef _MATRIXMUL_KERNEL_H_

否则有效数组数组之外的线程会破坏您的结果。事情适用于 32x32 的倍数，因为没有无效线程。在这种情况下，您将启动所需数量的线程。但在其他情况下，您正在启动额外的线程。如果允许这些额外的线程计算无效的矩阵位置，则会破坏结果。