【参加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;
    }
}