CUDA编程练习（四） 矩阵相乘

求求各位狠狠地批评我吧！我太想进步了，同时还能爽到😭

1 想法

这是个经典问题，也是很值得探究的有些难度的问题，也不是一天两天能把方法写完的，所以该文章按照以下路径去做这件事情：

• 写一个dummy版本，就不考虑任何内存优化，就硬算，一个线程一个数字；（涉及到线程id到矩阵对应元素的转化，笨，但是也考基本功）
• 写一个马赛克法，方便扩展，同时考虑了内存优化；（几本书上都有）
• 待补充。。。

2 代码

2.1 CPU计算与常量设置

#define MAT_M 2000
#define MAT_K 2000
#define MAT_N 2000
#define SCALE_REDUCTION 10
#define BLOCK_SIZE 16
#define BLOCK_PER_ROW ((MAT_N+(BLOCK_SIZE-1))/BLOCK_SIZE)  // 能不用寄存器我干嘛要用！
#define BLOCK_PER_COL ((MAT_M+(BLOCK_SIZE-1))/BLOCK_SIZE)
#define BLOCK_ON_K ((MAT_K+BLOCK_SIZE-1)/BLOCK_SIZE)
#define PARABLOCK_ON_ROW 64
#define random_double(RAND_MAX) (rand() / double(RAND_MAX))


// main函数部分：
int i, j;
for (int l=0; l<MAT_M*MAT_N; l++) {
    dimtrans_1to2(l, MAT_N, i, j);
    float tmp=0;
    for (int k=0; k<MAT_K; k++) {
        tmp += a[i*MAT_K+k]*b[k*MAT_N+j];
    }
    cpu_out[l] = tmp;
}

CPU主打一个暴力干脆不伤脑，有助于身心健康，延年益寿

都是没什么技术含量的东西，就放一起

2.2 dummy

某历史问题：nvcc -arch=sm_89 example.cu -o example

这样貌似能解决不支持的原子操作？4090的运算能力：8.9

2.2.1 1D

// 以后除非特殊要求，一般都是直接多block+可扩展，单的有啥可写
__global__ void dummy_1D_matMul(float *out, float *a, float *b, int m, int k, int n) {  // 乘完后，是m行n列
    int bid=blockIdx.x * blockDim.x + threadIdx.x;
    int leap=blockDim.x * gridDim.x;
    double tmp=0;
    int resNum=m*n;

    for (int i=bid; i<resNum; i+=leap) {
        // 这个线程负责哪个元素？
        int tar_i=i/n;
        int tar_j=i%n;

        // 怎么负责？累加
        for (int j=0; j<k; j++) {
            tmp += a[tar_i*k+j] * b[j*n+tar_j];  // 一个一行k个，一个一行n个，不难理解吧？
        }
        out[i] = tmp;
        tmp=0;
    }
}

为什么不同步？

你看，这都是各自干各自的，哪里需要同步啊？

为什么不用二维线程？

别急，我也喜欢二维线程

2.2.2 2D

注意