【参加CUDA线上训练营】--储存单元及矩阵乘法

1. GPU的存储单元

GPU的存储单元分为两大类： 板子上芯片周围的显存颗粒(on board)，读取速度相对慢，如下图中的local memory，global memory，constant memory，texture memory。 在GPU芯片内部(on chip)，读取速度相对快，如下图中 下图中的箭头双向表示可以读写，单向表示只能读。这些memory可以进一步细分： R/W可读可写memory： registers、local memory：线程私有memory，每个线程私有访问。 shared memory：一个block内的线程都可以访问，可以数据共享通信。 global memory：每个线程都可以读写。 R 只读memory：constant memory，texture memory是每个线程都可以读。 global memory，constant memory，texture memory和主机之间都可以通信读写，通常显卡说明书写的显存大小，就是global memory

在这里插入图片描述

2. GPU存储单元的分配与释放

1.申请GPU存储单元 当我们要为一个方阵M(m * m)申请GPU的存储单元时，使用下面的函数： cudaMalloc((void**) &d_m，sizeof(int) * m * m)，参数含义如下： 1）d_m：指向存储在Device端数据的地址的指针 2）sizeof(int) * m * m：存储在Device端空间的大小

2.释放GPU申请的存储单元函数 cudaFree(d_m); d_m：指向存储在Device端数据的地址的指针，从CPU内存传输到GPU存储单元 cudaMemcpy(d_m，h_m，sizeof(int）* m * m，cudaMemcpyHostToDevice），各参数的设置为： d_m：传输的目的地，GPU存储单元 h_m：数据的源地址，CPU存储单元 sizeof(int）* m * m：数据传输的大小 cudaMemcpyHostToDevice：数据传输的方向，CPU到GPU

在这里插入图片描述

2.矩阵乘法

CPU处理

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) 
{
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

GPU算法分析

在这里插入图片描述

在这里插入图片描述

算法实现：

__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
}