使用cuda C完成矩阵相乘算法详解

写在前面：

• 本笔记为 NVIDIA CUDA初级教程视频的笔记
• 视频连接：https://www.bilibili.com/video/BV1kx411m7Fk?p=10
• P9和P10讲解了矩阵相乘的cuda 实现
  

使用cuda C完成矩阵相乘算法详解

---

矩阵相乘大家应该都不陌生。

设有两个矩阵M和N，假设M和N都是方阵，维度均为width × width

如果M和N均为1000 × 1000的矩阵，总共要进行1000000次点乘。其中，每次点乘有1000次乘法和1000次加法。

Matirx Multiply：CPU实现

---

先来看看使用普通的c代码在CPU端如何实现

void MatrixMulOnHost(float* M,float* N,float* P,int width)
{
   
    for(int i=0;i<width;++i)
        for(int j=0;i<width;j++)
        {
   
            //sum对应每一次点乘（M的某一行×N的某一列）的结果
            float sum = 0;
            for(int k=0;k<width;k++)
            {
   
                float a = M[i*width+k];
                float b = N[k*width+j];
                sum+=a*b;
            }
            P[i*width+j]=sum;//乘累加的结果放到对应位置上
        }
}

可以看到循环计算结果P矩阵里的每一个元素。计算过程非常清晰。

从这里可以看到，这个计算存在非常大的并行性，即结果矩阵P里的每一个元素结果的计算与P中其他元素是不相关的，没有依赖性。

所以我们可以在GPU端上实现矩阵相乘。

Matirx Multiply：GPU实现

---

可以看到总共有3步：

1. 管理内存（在GPU上分配空间，将CPU端数据拷贝到GPU端）
2. GPU上并行处理（启动kernel函数）
3. 将结果拷贝回到CPU端

第1步：在算法框架中添加 CUDA memory transfers

第2步：CUDA C编程实现kernel

可以看到这里有2个问题

1. 使用线程的索引代替了双重循环，并行去做就可以
2. 不需要锁或同步，如果数据之间有依赖，存在同步问题，但这里每一个结果矩阵的元素是独立的，与别的元素无关，所以不需要锁存。

第3步 CUDA C编程调用kernel

源代码

#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdlib.h>

#define WIDTH 16

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int width)
{
   
	int tx = threadIdx.x;
	int ty = threadIdx.y;

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

int main(void)
{
   
	float M[16][16], N[16][16], P[16][16];
	int Width = 16;
	int NUM = 192;
	//初始化示例数据
	for (