使用 CUDA 的矩阵乘法答案

【问题标题】：Matrix Multiplication using CUDA使用 CUDA 的矩阵乘法
【发布时间】：2012-08-15 10:28:47
【问题描述】：

我对 CUDA 上的矩阵乘法很感兴趣。结果乘积矩阵始终为零。我已经阅读了一些示例代码，例如 matrix multiplication in cuda 来解决我的问题，但都是徒劳的。

除了 0 的不稳定结果之外，“宽度”（下面的代码）的最大尺寸甚至不是 512。我无法调试问题所在。也许我们可以在 StackOverflow 上讨论它。

我指的是“编程大规模并行处理器”

#include<cuda.h>
#include<stdio.h>

int main(void) {
    void MatrixMultiplication(float *, float *, float *, int);
    const int Width = 5;
    float M[Width*Width], N[Width*Width], P[Width*Width];
    for(int i = 0; i < (Width*Width) ; i++) {
        M[i] = 5;
        N[i] = 5;
        P[i] = 0;
    }
    MatrixMultiplication(M, N, P, Width);
    for(int i = 0; i < (Width*Width) ; i++) {
        printf("%d \n", P[i]);
    }
    int quit;
    scanf("%d",&quit);
    return 0;
}

//Matrix multiplication kernel - thread specification
__global__ void MatrixMulKernel(float *Md, float *Nd, float *Pd, int Width) {
    //2D Thread ID
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    //Pvalue stores the Pd element that is computed by the thread
    float Pvalue = 0;

    for(int k = 0; k < Width ; ++k) {
        float Mdelement = Md[ty*Width + k];
        float Ndelement = Nd[k*Width + tx];
        Pvalue += (Mdelement*Ndelement);
    }

    Pd[ty*Width + tx] = Pvalue;
}

void MatrixMultiplication(float *M, float *N, float *P, int Width) {
    int size = Width*Width*sizeof(float);
    float *Md, *Nd, *Pd;

    //Transfer M and N to device memory
    cudaMalloc((void**)&Md, size);
    cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice);
    cudaMalloc((void**)&Nd, size);
    cudaMemcpy(Nd,N,size,cudaMemcpyHostToDevice);

    //Allocate P on the device
    cudaMalloc((void**)&Pd,size);

    //Setup the execution configuration
    dim3 dimBlock(Width,Width);
    dim3 dimGrid(1,1);

    //Launch the device computation threads!
    MatrixMulKernel<<<dimGrid,dimBlock>>>(Md,Nd,Pd,Width);

    //Transfer P from device to host
    cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost);

    //Free device matrices
    cudaFree(Md);
    cudaFree(Nd);
    cudaFree(Pd);
}

【问题讨论】：

要获得正确的代码格式，您需要将所有代码缩进 4 个空格。您可以通过突出显示您的代码并按 Ctrl + K 轻松完成此操作。
谢谢杰夫！只是打算这样做
如果您不需要坚持自己的代码，CUDA C 编程指南有一个很棒的 matrix-mul 实现，可以处理具有除 2 的幂之外的其他维度的矩阵，并使用共享内存进行了优化。强烈推荐它用于现实世界的使用和学习。
@dave 水滴组成了浩瀚的海洋。要达到一个级别，您必须一步一步地遵循。例如如果我直接使用附录 A 中给出的代码，我可能永远不会知道什么是 __syncthreads() - 不考虑友好的语法 ;-)

标签： c cuda

【解决方案1】：

到目前为止，你都做得很好：

for(int i = 0; i < (Width*Width) ; i++) {
    printf("%d \n", P[i]);
}

我把它改成了 %f（因为它是一个浮点数），它们都打印得很好 :)

$ ./test.exe
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000
125.000000

【讨论】：

确实！尽管我没有阅读您的答案就想通了，只是打算发布它。

【解决方案2】：

我知道出了什么问题。我们来分析一下：

第 1 点：去除单调的“零值”

如上所述，您必须将 printf("%d \n", P[i]); 替换为 printf("%f \n", P[i]);

第 2 点：为什么程序会以 Width 512 的值失败？

实际上，即使是像 23 这样的小值也会失败。为什么？因为 23*23 > 512（截至今天，GPU 可以拥有的每个块的最大线程数！）

【讨论】：

【解决方案3】：

在您的 MatrixMulKernel 函数中，您的 for 循环就像

for(int k = 0; k < Width ; ++k) 
{
    //rest of code      
}

您必须使用Width*Width 而不是Width，因为您的数组大小为Width*Width。

【讨论】：

使用 CUDA 并行性的全部意义在于消除计算开销。在这种情况下，每个线程只负责乘积矩阵的 1 个结果。使用“宽度”迭代可以找到产品矩阵的一个结果（元素）。所以 Width*Width 在任何情况下都不起作用。